cccl
CCCL 2.6.1
| Guide des contributeurs | Conteneurs de développement | Discorde | Godbolt | Projet GitHub | Documentation |
|---|
Bienvenue aux bibliothèques Cuda Core Calcul (CCCL) où notre mission est de rendre Cuda plus délicieuse.
Ce référentiel unifie trois bibliothèques CUDA C ++ essentielles dans un seul référentiel pratique:
L'objectif de CCCL est de fournir aux développeurs CUDA C ++ des blocs de construction qui facilitent l'écriture de code sûr et efficace. Le rassemblement de ces bibliothèques rationalise votre processus de développement et élargit votre capacité à tirer parti de la puissance de CUDA C ++. Pour plus d'informations sur la décision d'unifier ces projets, consultez l'annonce ici.
Le concept des bibliothèques Cuda Core Calcul (CCCL) est née de manière organique des projets Thrust, Cub et LibcudAcxx qui ont été développés indépendamment au fil des ans avec un objectif similaire: fournir des abstractions C ++ de haute qualité, de haute performance et facile à utiliser pour les développeurs CUDA. Naturellement, il y a eu beaucoup de chevauchement parmi les trois projets, et il est devenu clair que la communauté serait mieux servie en les unifiant dans un seul référentiel.
La poussée est la bibliothèque d'algorithmes parallèles C ++ qui a inspiré l'introduction d'algorithmes parallèles à la bibliothèque standard C ++. L'interface de haut niveau de Thrust améliore considérablement la productivité des programmeurs tout en permettant la portabilité des performances entre les GPU et les CPU multicore via des backends configurables qui permettent d'utiliser plusieurs cadres de programmation parallèle (tels que CUDA, TBB et OpenMP).
Cub est une bibliothèque spécifique à CUDA de niveau inférieur conçu pour les algorithmes parallèles de vitesse de vitesse sur toutes les architectures GPU. En plus des algorithmes à l'échelle de l'appareil, il fournit des algorithmes coopératifs comme la réduction à l'échelle du bloc et le scan à l'échelle de la chaîne, fournissant aux développeurs du noyau CUDA avec des blocs de construction pour créer des noyaux personnalisés de vitesse et de lumière.
LibcudAcxx est la bibliothèque standard CUDA C ++. Il fournit une implémentation de la bibliothèque standard C ++ qui fonctionne à la fois dans l'hôte et le code de périphérique. De plus, il fournit des abstractions pour les fonctionnalités matérielles spécifiques à CUDA comme les primitives de synchronisation, le contrôle du cache, l'atomique, etc.
L'objectif principal de CCCL est de remplir un rôle similaire que la bibliothèque C ++ standard remplit pour les outils standard C ++: fournir des outils généraux à usage général pour les développeurs CUDA C ++, ce qui leur permet de se concentrer sur la résolution des problèmes qui comptent. L'unification de ces projets est la première étape vers la réalisation de cet objectif.
Ceci est un exemple simple démontrant l'utilisation de la fonctionnalité CCCL de Thrust, Cub et LibcudAcxx.
Il montre comment utiliser la poussée / cub / libcudacxx pour implémenter un simple noyau de réduction parallèle. Chaque bloc de filetage calcule la somme d'un sous-ensemble du tableau à l'aide de cub::BlockReduce . La somme de chaque bloc est ensuite réduite à une seule valeur à l'aide d'un ADD atomique via cuda::atomic_ref de LibcudAcxx.
Il montre ensuite comment la même réduction peut être effectuée en utilisant l'algorithme reduce de Thrust et compare les résultats.
Essayez-le en direct sur Godbolt!
# include < thrust/execution_policy.h >
# include < thrust/device_vector.h >
# include < cub/block/block_reduce.cuh >
# include < cuda/atomic >
# include < cuda/cmath >
# include < cuda/std/span >
# include < cstdio >
template < int block_size>
__global__ void reduce (cuda::std::span< int const > data, cuda::std::span< int > result) {
using BlockReduce = cub::BlockReduce< int , block_size>;
__shared__ typename BlockReduce::TempStorage temp_storage;
int const index = threadIdx. x + blockIdx. x * blockDim. x ;
int sum = 0 ;
if ( index < data. size ()) {
sum += data[ index ];
}
sum = BlockReduce (temp_storage). Sum (sum);
if (threadIdx. x == 0 ) {
cuda::atomic_ref< int , cuda::thread_scope_device> atomic_result (result. front ());
atomic_result. fetch_add (sum, cuda::memory_order_relaxed);
}
}
int main () {
// Allocate and initialize input data
int const N = 1000 ;
thrust::device_vector< int > data (N);
thrust::fill (data. begin (), data. end (), 1 );
// Allocate output data
thrust::device_vector< int > kernel_result ( 1 );
// Compute the sum reduction of `data` using a custom kernel
constexpr int block_size = 256 ;
int const num_blocks = cuda::ceil_div (N, block_size);
reduce<block_size><<<num_blocks, block_size>>>(cuda::std::span< int const >( thrust::raw_pointer_cast (data. data ()), data. size ()),
cuda::std::span< int >( thrust::raw_pointer_cast (kernel_result. data ()), 1 ));
auto const err = cudaDeviceSynchronize ();
if (err != cudaSuccess) {
std::cout << " Error: " << cudaGetErrorString (err) << std::endl;
return - 1 ;
}
int const custom_result = kernel_result[ 0 ];
// Compute the same sum reduction using Thrust
int const thrust_result = thrust::reduce (thrust::device, data. begin (), data. end (), 0 );
// Ensure the two solutions are identical
std::printf ( " Custom kernel sum: %d n " , custom_result);
std::printf ( " Thrust reduce sum: %d n " , thrust_result);
assert (kernel_result[ 0 ] == thrust_result);
return 0 ;
}Tout dans CCCL est uniquement en tête. Par conséquent, les utilisateurs ont seulement besoin de se préoccuper de la façon dont ils obtiennent les fichiers d'en-tête et de la façon dont ils les intègrent dans leur système de construction.
Le moyen le plus simple de commencer à utiliser CCCL est via la boîte à outils CUDA qui comprend les en-têtes CCCL. Lorsque vous compilez avec nvcc , il ajoute automatiquement des en-têtes CCCL à votre chemin d'inclusion afin que vous puissiez simplement #include n'importe quel en-tête CCCL dans votre code sans aucune configuration supplémentaire requise.
Si la compilation avec un autre compilateur, vous devrez mettre à jour le chemin de recherche de vos systèmes de build inclut pour pointer vers les en-têtes CCCL dans votre installation CTK (par exemple, /usr/local/cuda/include ).
# include < thrust/device_vector.h >
# include < cub/cub.cuh >
# include < cuda/std/atomic > Les utilisateurs qui souhaitent rester à la pointe du développement CCCL sont encouragés à utiliser CCCL à partir de GitHub. L'utilisation d'une version plus récente de CCCL avec une ancienne version de la boîte à outils CUDA est prise en charge, mais pas l'inverse. Pour des informations complètes sur la compatibilité entre CCCL et la boîte à outils CUDA, consultez notre prise en charge de la plate-forme.
Tout dans CCCL est uniquement en tête, donc le clonage et l'inclure dans un projet simple est aussi simple que le suivant:
git clone https://github.com/NVIDIA/cccl.git
nvcc -Icccl/thrust -Icccl/libcudacxx/include -Icccl/cub main.cu -o mainRemarque Utilisez
-Iet non-isystempour éviter les collisions avec les en-têtes CCCL implicitement incluses parnvccà partir de la boîte à outils CUDA. Tous les en-têtes CCCL utilisent#pragma system_headerpour s'assurer que les avertissements seront toujours réduits au silence comme si vous utilisez-isystem, voir # 527 pour plus d'informations.
Une génération minimale qui ne génère que des règles d'installation peut être configurée à l'aide de l' install CMake Preset:
git clone https://github.com/NVIDIA/cccl.git
cd cccl
cmake --preset install -DCMAKE_INSTALL_PREFIX=/usr/local/
cd build/install
ninja install Pour inclure des bibliothèques expérimentales dans l'installation, utilisez le répertoire préréglé et build d' install-unstable .
Pour installer uniquement les bibliothèques expérimentales, utilisez le répertoire préréglé et build et de construction install-unstable-only .
CCCL fournit également des packages conda de chaque version via la chaîne conda-forge :
conda config --add channels conda-forge
conda install cccl Cela installera le dernier CCCL sur $CONDA_PREFIX/include/ $CONDA_PREFIX/lib/cmake/ $. Il est découvrable par CMake via find_package(CCCL) et peut être utilisé par tous les compilateurs de l'environnement Conda. Pour plus d'informations, consultez cette introduction à Conda-Forge.
Si vous souhaitez utiliser la même version CCCL qui expédiée avec une boîte à outils CUDA particulière, par exemple CUDA 12.4, vous pouvez installer CCCL avec:
conda config --add channels conda-forge
conda install cuda-cccl cuda-version=12.4 Le MetaPackage cuda-cccl installe la version cccl qui a été expédiée avec la boîte à outils CUDA correspondant à cuda-version . Si vous souhaitez mettre à jour vers le dernier cccl après l'installation cuda-cccl , désinstallez cuda-cccl avant de mettre à jour cccl :
conda uninstall cuda-cccl
conda install -c conda-forge ccclRemarque Il existe également des packages conda avec des noms comme
cuda-cccl_linux-64. Ces packages contiennent les versions CCCL expédiées dans le cadre de la boîte à outils CUDA, mais sont conçues pour une utilisation interne par la boîte à outils CUDA. Installez à la placecccloucuda-cccl, pour une compatibilité avec les compilateurs conda. Pour plus d'informations, consultez la recette CCCL Conda-Forge.
CCCL utilise CMake pour toutes les infrastructures de build et d'installation, y compris les tests ainsi que des cibles à relier dans d'autres projets CMake. Par conséquent, CMake est le moyen recommandé d'intégrer CCCL dans un autre projet.
Pour un exemple complet de la façon de procéder à l'aide de CMake Package Manager, consultez notre exemple de base.
D'autres systèmes de construction devraient fonctionner, mais seul le CMake est testé. Les contributions pour simplifier l'intégration de CCCL dans d'autres systèmes de construction sont les bienvenues.
Vous souhaitez contribuer à améliorer CCCL? Consultez notre guide de contribution pour un aperçu complet de tout ce que vous devez savoir pour configurer votre environnement de développement, apporter des modifications, exécuter des tests et soumettre un RP.
Objectif: Cette section décrit où les utilisateurs peuvent s'attendre à ce que CCCL se compile et s'exécute avec succès.
En général, CCCL devrait travailler partout où la boîte à outils CUDA est prise en charge, cependant, le diable est dans les détails. Les sections ci-dessous décrivent les détails de la prise en charge et des tests pour différentes versions de la boîte à outils CUDA, des compilateurs hôtes et des dialectes C ++.
Résumé:
Les utilisateurs de CCCL sont encouragés à capitaliser sur les dernières améliorations et "Live at Head" en utilisant toujours la dernière version de CCCL. Pour une expérience transparente, vous pouvez mettre à niveau CCCL indépendamment de l'ensemble de la boîte à outils CUDA. Cela est possible car CCCL maintient la compatibilité vers l'arrière avec la dernière version de patch de chaque version mineure CTK de la série actuelle et précédente de la série de versions majeures. Dans certains cas exceptionnels, la version mineure minimale prise en charge de la version de la boîte à outils CUDA peut devoir être plus récente que la version la plus ancienne de sa série de versions principales. Par exemple, CCCL nécessite une version minimale prise en charge de 11.1 de la série 11.x en raison d'un problème de compilateur inévitable présent dans CTK 11.0.
Lorsqu'un nouveau CTK majeur est publié, nous supprimons la prise en charge de la plus ancienne version majeure prise en charge.
| Version CCCL | Prend en charge la version de la boîte à outils CUDA |
|---|---|
| 2.x | 11.1 - 11.8, 12.x (seules les dernières versions de correctifs) |
| 3.x (futur) | 12.x, 13.x (seules les dernières versions de correctifs) |
Le code bien élevé à l'aide du dernier CCCL devrait compiler et exécuter avec succès avec toute version CTK prise en charge. Des exceptions peuvent se produire pour de nouvelles fonctionnalités qui dépendent de nouvelles fonctionnalités CTK, de sorte que ces fonctionnalités ne fonctionneraient pas sur les anciennes versions du CTK. Par exemple, le support C ++ 20 n'a pas été ajouté à nvcc jusqu'à CUDA 12.0, donc les fonctionnalités CCCL qui dépendent de C ++ 20 ne fonctionneraient pas avec CTK 11.x.
Les utilisateurs peuvent intégrer une version plus récente de CCCL dans un CTK plus ancien, mais pas l'inverse. Cela signifie qu'une ancienne version de CCCL n'est pas compatible avec un CTK plus récent. En d'autres termes, CCCL n'est jamais en avant compatible avec la boîte à outils CUDA.
Le tableau ci-dessous résume la compatibilité du CTK et du CCCL:
| Version CTK | Version CCCL incluse | CCCL souhaité | Soutenu? | Notes |
|---|---|---|---|---|
CTK XY | CCCL MAJOR.MINOR | Cccl MAJOR.MINOR+n | ✅ | Certaines nouvelles fonctionnalités peuvent ne pas fonctionner |
CTK XY | CCCL MAJOR.MINOR | CCCL MAJOR+1.MINOR | ✅ | Pauses possibles; Certaines nouvelles fonctionnalités pourraient ne pas être disponibles |
CTK XY | CCCL MAJOR.MINOR | CCCL MAJOR+2.MINOR | CCCL ne prend en charge que deux versions majeures CTK | |
CTK XY | CCCL MAJOR.MINOR | Cccl MAJOR.MINOR-n | CCCL n'est pas compatible en avant | |
CTK XY | CCCL MAJOR.MINOR | CCCL MAJOR-n.MINOR | CCCL n'est pas compatible en avant |
Pour plus d'informations sur le versioning CCCL, la compatibilité API / ABI et les modifications de rupture, consultez la section de versioning ci-dessous.
Sauf indication contraire, CCCL prend en charge tous les mêmes systèmes d'exploitation que la boîte à outils CUDA, qui est documentée ici:
Sauf indication contraire, CCCL prend en charge les mêmes compilateurs hôtes que la dernière boîte à outils CUDA, qui est documentée ici:
Lorsque nous utilisons des boîtes à outils CUDA plus anciennes, nous ne prenons également en charge que les compilateurs hôtes de la dernière boîte à outils CUDA, mais au moins le compilateur hôte le plus récent de toute boîte à outils CUDA plus ancien prise en charge.
Nous pouvons conserver le soutien de compilateurs supplémentaires et accepterons les correctifs correspondants de la communauté avec des correctifs raisonnables. Mais nous n'investirons pas beaucoup de temps dans le triage ou la résolution des problèmes pour les compilateurs plus âgés.
Dans l'esprit de "vous ne soutient que ce que vous testez", consultez notre aperçu CI pour plus d'informations sur ce que nous testons exactement.
Sauf indication contraire, CCCL prend en charge toutes les mêmes architectures / capacités de calcul GPU que la boîte à outils CUDA, qui est documentée ici: https://docs.nvidia.com/cuda/cuda-c-programming-uide/index.html#compute-capabilité
Notez que certaines fonctionnalités ne peuvent prendre en charge que certaines architectures / capacités de calcul.
La stratégie de test de CCCL établit un équilibre entre le test autant de configurations que possible et le maintien de temps CI raisonnables.
Pour les versions CUDA Toolkit, les tests se font contre les versions les plus anciennes et les plus récentes. Par exemple, si la dernière version de la boîte à outils CUDA est de 12,3, des tests sont effectués contre 11.1 et 12.3. Pour chaque version CUDA, les builds sont complétés par rapport à tous les compilateurs hôtes pris en charge avec tous les dialectes C ++ pris en charge.
La stratégie de test et la matrice évoluent constamment. La matrice définie dans le fichier ci/matrix.yaml est la source définitive de vérité. Pour plus d'informations sur notre pipeline CI, voir ici.
Objectif: Cette section décrit comment CCCL est versé, les garanties de stabilité API / ABI et les directives de compatibilité pour minimiser les maux de tête de mise à niveau.
Résumé
cub:: ou thrust:: Espaces de nomscuda:: peut se produire à tout moment, mais sera reflété par l'incrément de la version ABI qui est intégrée dans un espace de noms en ligne pour tous les symboles cuda:: . Plusieurs versions ABI peuvent être prises en charge simultanément.Remarque: Avant de fusionner Thrust, Cub et LibCudAcxx dans ce référentiel, chaque bibliothèque a été versée indépendamment en fonction du versioning sémantique. À partir de la version 2.1, les trois bibliothèques ont synchronisé leurs versions de version dans leurs référentiels séparés. À l'avenir, CCCL continuera d'être publié sous une seule version sémantique, 2.2.0 étant la première version du référentiel NVIDIA / CCCL.
Un changement de rupture est un changement pour les fonctionnalités explicitement prises en charge entre les versions publiées qui obligeraient un utilisateur à fonctionner pour passer à la version plus récente.
Dans la limite, tout changement a le potentiel de briser quelqu'un quelque part. En conséquence, toutes les modifications possibles de rupture des sources ne sont pas considérées comme des modifications de rupture de l'API publique qui justifient de frapper la version sémantique majeure.
Les sections ci-dessous décrivent les détails de la rupture des modifications de l'API et ABI de CCCL.
L'API publique de CCCL est l'intégralité de la fonctionnalité intentionnellement exposée à fournir l'utilité de la bibliothèque.
En d'autres termes, l'API publique de CCCL va au-delà des seules signatures de fonction et comprend (mais sans s'y limiter):
De plus, l'API publique de CCCL n'inclut aucun des éléments suivants:
_ ou __detail , y compris le detail:: espace de noms ou macrodetail/ répertoire ou sous-répertoire de celui-ciEn général, l'objectif est d'éviter de rompre quoi que ce soit dans l'API publique. Ces modifications ne sont apportées que s'ils offrent aux utilisateurs de meilleures performances, des API plus faciles à comprendre et / ou des API plus cohérentes.
Tout changement de rupture de l'API public nécessitera de cogner le principal numéro de version de CCCL. Conformément à la compatibilité de la version mineure de CUDA, les modifications de rupture d'API et les bosses de version majeure CCCL ne feront qu'accéder à une nouvelle version majeure de la boîte à outils CUDA.
Tout ce qui ne fait pas partie de l'API publique peut changer à tout moment sans avertissement.
L'API publique de tous les composants de CCCL partage une version sémantique unifiée de MAJOR.MINOR.PATCH .
Seule la dernière version publiée est prise en charge. En règle générale, les fonctionnalités et les corrections de bogues ne sont pas recouvertes de version ou de branches publiées précédemment.
La méthode préférée pour interroger la version consiste à utiliser CCCL_[MAJOR/MINOR/PATCH_]VERSION comme décrit ci-dessous. Pour la compatibilité descendante, les définitions de version Thrust / Cub / LibCudAcxxx sont disponibles et seront toujours cohérentes avec CCCL_VERSION . Notez que Thrust / Cub utilise un schéma MMMmmmpp tandis que le CCCL et le libcudacxx utilisent MMMmmmppp .
| Cccl | libcudacxx | Poussée | Petit | |
|---|---|---|---|---|
| Tête | <cuda/version> | <cuda/std/version> | <thrust/version.h> | <cub/version.h> |
| Version majeure | CCCL_MAJOR_VERSION | _LIBCUDACXX_CUDA_API_VERSION_MAJOR | THRUST_MAJOR_VERSION | CUB_MAJOR_VERSION |
| Version mineure | CCCL_MINOR_VERSION | _LIBCUDACXX_CUDA_API_VERSION_MINOR | THRUST_MINOR_VERSION | CUB_MINOR_VERSION |
| Version patch / subminor | CCCL_PATCH_VERSION | _LIBCUDACXX_CUDA_API_VERSION_PATCH | THRUST_SUBMINOR_VERSION | CUB_SUBMINOR_VERSION |
| Version concaténée | CCCL_VERSION (MMMmmmppp) | _LIBCUDACXX_CUDA_API_VERSION (MMMmmmppp) | THRUST_VERSION (MMMmmmpp) | CUB_VERSION (MMMmmmpp) |
L'interface binaire d'application (ABI) est un ensemble de règles pour:
ABI d'une bibliothèque comprend, mais sans s'y limiter:
Un changement de rupture ABI est tout changement qui se traduit par un changement à l'ABI d'une fonction ou de type dans l'API publique. Par exemple, l'ajout d'un nouveau membre de données à une structure est un changement de rupture ABI car il modifie la taille du type.
Dans CCCL, les garanties concernant ABI sont les suivantes:
thrust:: et cub:: Les espaces de noms peuvent casser Abi à tout moment sans avertissement.thrust:: et cub:: Symboles comprend les architectures CUDA utilisées pour la compilation. Par conséquent, un symbole thrust:: ou cub:: peut avoir un ABI différent si:-x cu ) vs C ++ Source ( -x cpp )cuda:: peuvent également casser Abi à tout moment. Cependant, les symboles cuda:: intégrent un numéro de version ABI qui est incrémenté chaque fois qu'une rupture ABI se produit. Plusieurs versions ABI peuvent être prises en charge simultanément et, par conséquent, les utilisateurs ont la possibilité de revenir à une version ABI antérieure. Pour plus d'informations, consultez ici.Qui devrait se soucier d'Abi?
En général, les utilisateurs de CCCL n'ont qu'à se soucier des problèmes d'ABI lors de la construction ou de l'utilisation d'un artefact binaire (comme une bibliothèque partagée) dont l'API comprend directement ou indirectement les types fournis par CCCL.
Par exemple, considérez si libA.so a été construit à l'aide de CCCL Version X et son API publique comprend une fonction comme:
void foo (cuda::std::optional< int >); Y une autre bibliothèque, libB.so , est compilée à l'aide de CCCL X Y et utilise foo à partir de libA.so . Contrairement aux modifications de rupture de l'API, les ruptures ABI ne nécessitent généralement pas de modifications de code et nécessitent que tout recompilation pour utiliser la même version ABI.
Pour en savoir plus sur ABI et pourquoi il est important, voyez ce qu'est ABI et que devrait faire C ++ à ce sujet ?.
Comme mentionné ci-dessus, toutes les modifications de rupture des sources possibles ne constituent pas un changement de rupture qui nécessiterait l'intégration du numéro de version majeure de l'API de CCCL.
Les utilisateurs sont encouragés à respecter les directives suivantes afin de minimiser le risque de perturbations de la part accidentellement en fonction des parties de CCCL qui ne font pas partie de l'API publique:
thrust:: , cub:: , nv:: ou cuda:: Namespaces sauf si une exception est notée pour un symbole spécifique, par exemple, spécialisé cuda::std::iterator_traitsthrust:: , cub:: , cuda:: , ou nv:: Espaces.thrust:: , cub:: , cuda:: , ou nv:: Espaces de noms._ , __ ou avec detail n'importe où dans son nom, y compris un detail:: espace de noms ou macro#include le fichier d'en-tête qui déclare ce symbole. En d'autres termes, ne comptez pas sur les en-têtes implicitement inclus par d'autres en-têtes.Des parties de cette section ont été inspirées par les directives de compatibilité d'Abseil.
Nous ferons de notre mieux pour informer les utilisateurs avant d'apporter des modifications de rupture à l'API publique, ABI ou à la modification des plateformes et des compilateurs pris en charge.
Le cas échéant, les dépréciations se présentent sous la forme d'avertissements programmatiques qui peuvent être désactivés.
La période de dépréciation dépendra de l'impact du changement, mais durera généralement au moins 2 versions de versions mineures.
À venir!
Pour un aperçu détaillé du pipeline CI, voir Ci-Overview.md.
Des projets liés à la mission de CCCL de rendre Cuda plus délicieux:
Votre projet utilise-t-il CCCL? Ouvrez un PR pour ajouter votre projet à cette liste!
