cccl
CCCL 2.6.1
| Guía de contribuyentes | Contenedores de desarrollo | Discordia | Mochila | Proyecto Github | Documentación |
|---|
Bienvenido a las Bibliotecas CUDA Core Compute (CCCL) donde nuestra misión es hacer que CUDA sea más encantador.
Este repositorio unifica tres bibliotecas CUDA C ++ esenciales en un solo repositorio conveniente:
El objetivo de CCCL es proporcionar a los desarrolladores CUDA C ++ a los bloques de construcción que faciliten la escritura de código seguro y eficiente. Reunir a estas bibliotecas optimiza su proceso de desarrollo y amplía su capacidad para aprovechar el poder de CUDA C ++. Para obtener más información sobre la decisión de unificar estos proyectos, consulte el anuncio aquí.
El concepto para las Bibliotecas CUDA Core Compute (CCCL) creció orgánicamente de los proyectos de empuje, CUB y Libcudacxx que se desarrollaron de forma independiente a lo largo de los años con un objetivo similar: proporcionar abstracciones de C ++ de alta calidad, alta rendimiento y fácil de usar para los desarrolladores de CUDA. Naturalmente, hubo mucha superposición entre los tres proyectos, y quedó claro que la comunidad estaría mejor atendiéndolos en un solo repositorio.
El empuje es la biblioteca de algoritmos paralelos C ++ que inspiró la introducción de algoritmos paralelos a la biblioteca estándar C ++. La interfaz de alto nivel de Thrust mejora enormemente la productividad del programador al tiempo que permite la portabilidad de rendimiento entre las GPU y las CPU multinúcleo a través de backends configurables que permiten usar múltiples marcos de programación paralelos (como CUDA, TBB y OpenMP).
CUB es una biblioteca de nivel inferior específica de CUDA diseñada para algoritmos paralelos de velocidad de luz en todas las arquitecturas de GPU. Además de los algoritmos de todo el dispositivo, proporciona algoritmos cooperativos como la reducción de todo el bloque y el escaneo de toda la urdimbre, proporcionando a los desarrolladores de núcleos CUDA con bloques de construcción para crear núcleos personalizados de velocidad de luz.
LibCudacxx es la biblioteca estándar CUDA C ++. Proporciona una implementación de la biblioteca estándar C ++ que funciona tanto en código de host como de dispositivo. Además, proporciona abstracciones para características de hardware específicas de CUDA como primitivas de sincronización, control de caché, atómica y más.
El objetivo principal de CCCL es cumplir un papel similar que la biblioteca estándar de C ++ se llena para C ++ estándar: proporcionar herramientas de velocidad de uso general y de velocidad a los desarrolladores CUDA C ++, lo que les permite concentrarse en resolver los problemas que importan. Unificar estos proyectos es el primer paso para darse cuenta de ese objetivo.
Este es un ejemplo simple que demuestra el uso de la funcionalidad CCCL a partir de empuje, Cub y Libcudacxx.
Muestra cómo usar empuje/cub/libcudacxx para implementar un núcleo de reducción paralelo simple. Cada bloque de subprocesos calcula la suma de un subconjunto de la matriz usando cub::BlockReduce . La suma de cada bloque se reduce a un solo valor utilizando una adición atómica a través de cuda::atomic_ref de libcudacxx.
Luego muestra cómo se puede hacer la misma reducción utilizando el algoritmo reduce de Thrust y compara los resultados.
¡Pruébalo en vivo en 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 ;
}Todo en CCCL solo es encabezado. Por lo tanto, los usuarios solo necesitan preocuparse por cómo obtienen los archivos de encabezado y cómo los incorporan a su sistema de compilación.
La forma más fácil de comenzar a usar CCCL es a través del kit de herramientas CUDA que incluye los encabezados CCCL. Cuando se compila con nvcc , agrega automáticamente los encabezados CCCL a su ruta de inclusión para que simplemente pueda #include cualquier encabezado CCCL en su código sin necesidad de configuración adicional.
Si se compila con otro compilador, deberá actualizar la ruta de búsqueda de su sistema de compilación para apuntar a los encabezados CCCL en su instalación CTK (p. Ej., /usr/local/cuda/include ).
# include < thrust/device_vector.h >
# include < cub/cub.cuh >
# include < cuda/std/atomic > Se alienta a los usuarios que desean permanecer a la vanguardia del desarrollo de CCCL a usar CCCL de GitHub. Se admite usar una versión más nueva de CCCL con una versión anterior del Kit CUDA Toolkit, pero no al revés. Para obtener información completa sobre la compatibilidad entre CCCL y el CUDA Toolkit, consulte el soporte de nuestra plataforma.
Todo en CCCL es solo de encabezado, por lo que clonarse e incluirlo en un proyecto simple es tan fácil como el siguiente:
git clone https://github.com/NVIDIA/cccl.git
nvcc -Icccl/thrust -Icccl/libcudacxx/include -Icccl/cub main.cu -o mainNota Use
-Iy no-isystempara evitar colisiones con los encabezados CCCL incluidos implícitamente pornvccdel CUDA Toolkit. Todos los encabezados CCCL usan#pragma system_headerpara garantizar que las advertencias aún se silenciarán como si usen-isystem, consulte #527 para obtener más información.
Una compilación mínima que solo genera reglas de instalación se puede configurar utilizando el preajuste install CMAKE:
git clone https://github.com/NVIDIA/cccl.git
cd cccl
cmake --preset install -DCMAKE_INSTALL_PREFIX=/usr/local/
cd build/install
ninja install Para incluir bibliotecas experimentales en la instalación, use el directorio de compilación y el directorio de compilación install-unstable .
Para instalar solo las bibliotecas experimentales, use el directorio de compilación y el directorio de compilación install-unstable-only .
CCCL también proporciona paquetes de conda de cada lanzamiento a través del canal conda-forge :
conda config --add channels conda-forge
conda install cccl Esto instalará el último CCCL en el $CONDA_PREFIX/include/ y $CONDA_PREFIX/lib/cmake/ directorios. CMake lo descubre a través de find_package(CCCL) y puede ser utilizado por cualquier compilador en el entorno de conda. Para obtener más información, consulte esta introducción a Conda-Forge.
Si desea usar la misma versión CCCL que se envió con un kit de herramientas CUDA particular, por ejemplo, CUDA 12.4, puede instalar CCCL con:
conda config --add channels conda-forge
conda install cuda-cccl cuda-version=12.4 El metapackage cuda-cccl instala la versión cccl que se envió con el kit de herramientas CUDA correspondiente a cuda-version . Si desea actualizar el último cccl después de instalar cuda-cccl , desinstale cuda-cccl antes de actualizar cccl :
conda uninstall cuda-cccl
conda install -c conda-forge ccclTenga en cuenta que también hay paquetes de conda con nombres como
cuda-cccl_linux-64. Esos paquetes contienen las versiones CCCL enviadas como parte del kit de herramientas CUDA, pero están diseñadas para uso interno por el kit de herramientas CUDA. Instaleccclocuda-ccclen su lugar, para la compatibilidad con los compiladores de conda. Para obtener más información, consulte la receta CCCL Conda-Forge.
CCCL utiliza CMake para toda la infraestructura de construcción e instalación, incluidas las pruebas y objetivos para vincular en otros proyectos CMake. Por lo tanto, CMake es la forma recomendada de integrar CCCL en otro proyecto.
Para obtener un ejemplo completo de cómo hacer esto utilizando el Administrador de paquetes CMake, consulte nuestro proyecto de ejemplo básico.
Otros sistemas de compilación deberían funcionar, pero solo se prueba Cmake. Las contribuciones para simplificar la integración de CCCL en otros sistemas de compilación son bienvenidas.
¿Interesado en contribuir a mejorar CCCL? Consulte nuestra guía de contribución para obtener una descripción completa de todo lo que necesita saber para configurar su entorno de desarrollo, realizar cambios, ejecutar pruebas y enviar un PR.
Objetivo: Esta sección describe dónde los usuarios pueden esperar que CCCL se compile y se ejecute con éxito.
En general, CCCL debería trabajar en todas partes, el kit de herramientas CUDA es compatible, sin embargo, el diablo está en los detalles. Las secciones a continuación describen los detalles del soporte y las pruebas de diferentes versiones del kit de herramientas CUDA, compiladores host y dialectos C ++.
Resumen:
Se alienta a los usuarios de CCCL a capitalizar las últimas mejoras y "vivir en la cabeza" al usar siempre la versión más reciente de CCCL. Para una experiencia perfecta, puede actualizar CCCL independientemente de todo el kit de herramientas CUDA. Esto es posible porque CCCL mantiene la compatibilidad hacia atrás con la última versión de parche de cada lanzamiento de CTK menor de la serie de versiones principales actuales y anteriores. En algunos casos excepcionales, la versión menor mínima compatible del lanzamiento de Cuda Toolkit puede deberse ser más nuevo que la versión más antigua dentro de su serie de versiones principales. Por ejemplo, CCCL requiere una versión mínima compatible de 11.1 de la serie 11.X debido a un problema de compilador inevitable presente en CTK 11.0.
Cuando se lanza un nuevo CTK importante, eliminamos el soporte para la versión principal más antigua compatible.
| Versión CCCL | Admite la versión CUDA Toolkit |
|---|---|
| 2.x | 11.1 - 11.8, 12.x (solo las últimas comunicadas de parche) |
| 3.x (futuro) | 12.x, 13.x (solo las últimas versiones de parche) |
El código de bienestar con el último CCCL debería compilar y ejecutarse con éxito con cualquier versión CTK compatible. Pueden ocurrir excepciones para nuevas características que dependen de nuevas características de CTK, por lo que esas características no funcionarían en versiones anteriores del CTK. Por ejemplo, el soporte de C ++ 20 no se agregó a nvcc hasta CUDA 12.0, por lo que las características de CCCL que dependen de C ++ 20 no funcionarían con CTK 11.x.
Los usuarios pueden integrar una versión más nueva de CCCL en un CTK más antiguo, pero no al revés. Esto significa que una versión anterior de CCCL no es compatible con un CTK más nuevo. En otras palabras, CCCL nunca es compatible con el kit de herramientas CUDA.
La siguiente tabla resume la compatibilidad de CTK y CCCL:
| Versión CTK | Incluida la versión CCCL | CCCL deseado | Compatible? | Notas |
|---|---|---|---|---|
CTK XY | CCCL MAJOR.MINOR | CCCL MAJOR.MINOR+n | ✅ | Algunas características nuevas pueden no funcionar |
CTK XY | CCCL MAJOR.MINOR | CCCL MAJOR+1.MINOR | ✅ | Posibles descansos; Es posible que algunas características nuevas no estén disponibles |
CTK XY | CCCL MAJOR.MINOR | CCCL MAJOR+2.MINOR | CCCL admite solo dos versiones principales de CTK | |
CTK XY | CCCL MAJOR.MINOR | CCCL MAJOR.MINOR-n | CCCL no es compatible con el avance | |
CTK XY | CCCL MAJOR.MINOR | CCCL MAJOR-n.MINOR | CCCL no es compatible con el avance |
Para obtener más información sobre el versiones de CCCL, la compatibilidad API/ABI y los cambios de ruptura, consulte la sección de versiones a continuación.
A menos que se especifique lo contrario, CCCL admite los mismos sistemas operativos que el kit de herramientas CUDA, que se documentan aquí:
A menos que se especifique lo contrario, CCCL admite los mismos compiladores host que el último Kit de herramientas CUDA, que se documentan aquí:
Cuando utilizamos kits de herramientas CUDA más antiguos, también solo admitemos los compiladores host del último kit de herramientas CUDA, pero al menos el compilador host más reciente de cualquier kit de herramientas CUDA más antiguo compatible.
Podemos retener el soporte de compiladores adicionales y aceptaremos parches correspondientes de la comunidad con soluciones razonables. Pero no invertiremos un tiempo significativo en el triado o solucionando problemas para compiladores más antiguos.
En el espíritu de "Usted solo apoya lo que prueba", vea nuestra descripción general del CI para obtener más información sobre exactamente lo que probamos.
A menos que se especifique lo contrario, CCCL admite las mismas capacidades de arquitecturas de GPU/compute que el kit de herramientas CUDA, que se documentan aquí: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#compute-capability
Tenga en cuenta que algunas características solo pueden admitir ciertas arquitecturas/capacidades de cálculo.
La estrategia de prueba de CCCL entaca un equilibrio entre las pruebas tantas configuraciones como sea posible y mantener tiempos de CI razonables.
Para las versiones CUDA Toolkit, las pruebas se realizan con las versiones más antiguas y más recientes. Por ejemplo, si la última versión del kit de herramientas CUDA es 12.3, las pruebas se realizan contra 11.1 y 12.3. Para cada versión CUDA, las compilaciones se completan en todos los compiladores host compatibles con todos los dialectos C ++ compatibles.
La estrategia de prueba y la matriz están evolucionando constantemente. La matriz definida en el archivo ci/matrix.yaml es la fuente definitiva de la verdad. Para obtener más información sobre nuestra tubería CI, consulte aquí.
OBJETIVO: Esta sección describe cómo se verifica CCCL, garantías de estabilidad API/ABI y pautas de compatibilidad para minimizar los dolores de cabeza de actualización.
Resumen
cub:: o thrust:: espacios de nombrescuda:: pueden ocurrir en cualquier momento, pero se reflejará incrementando la versión ABI que está integrada en un espacio de nombres en línea para todos los símbolos cuda:: . Múltiples versiones ABI pueden ser compatibles simultáneamente.Nota: Antes de fusionar el empuje, Cub y Libcudacxx en este repositorio, cada biblioteca fue versión independientemente de acuerdo con las versiones semánticas. Comenzando con la versión 2.1, las tres bibliotecas sincronizaron sus versiones de lanzamiento en sus repositorios separados. En el futuro, CCCL continuará lanzándose bajo una sola versión semántica, siendo 2.2.0 la primera versión del repositorio NVIDIA/CCCL.
Un cambio de ruptura es un cambio a la funcionalidad compatible explícitamente entre las versiones liberadas que requerirían que un usuario funcione para actualizar a la versión más nueva.
En el límite, cualquier cambio tiene el potencial de romper a alguien en algún lugar. Como resultado, no todos los cambios posibles de ruptura de la fuente se consideran que rompen cambios en la API pública que justifican aumentar la versión semántica principal.
Las secciones a continuación describen los detalles de los cambios de ruptura en la API y ABI de CCCL.
La API pública de CCCL es la totalidad de la funcionalidad intencionalmente expuesta para proporcionar la utilidad de la biblioteca.
En otras palabras, la API pública de CCCL va más allá de las firmas de la función justa e incluye (pero no se limita a):
Además, la API pública de CCCL no incluye nada de los siguientes:
_ o __detail , incluido el detail:: Namespace o una macrodetail/ directorio o subdirectorio del mismoEn general, el objetivo es evitar romper cualquier cosa en la API pública. Tales cambios se realizan solo si ofrecen a los usuarios un mejor rendimiento, las API más fáciles de entender y/o API más consistentes.
Cualquier cambio de ruptura en la API pública requerirá que el número de versión principal de CCCL. De acuerdo con la compatibilidad de la versión menor de CUDA, los cambios de ruptura de API y los golpes principales de la versión CCCL solo ocurrirán coincidiendo con una nueva versión de versión principal del CUDA Toolkit.
Cualquier cosa que no sea parte de la API pública puede cambiar en cualquier momento sin previo aviso.
La API pública de todos los componentes de CCCL comparte una versión semántica unificada de MAJOR.MINOR.PATCH .
Solo la versión lanzada más recientemente es compatible. Como regla general, las características y las correcciones de errores no están respaldadas a la versión o ramas lanzadas previamente.
El método preferido para consultar la versión es usar CCCL_[MAJOR/MINOR/PATCH_]VERSION como se describe a continuación. Para la compatibilidad hacia atrás, las definiciones de versión de empuje/cub/libcudacxxx están disponibles y siempre serán consistentes con CCCL_VERSION . Tenga en cuenta que el empuje/cachorro usa un esquema MMMmmmpp mientras que el CCCL y LibCudacxx usan MMMmmmppp .
| CCCL | libcudacxx | Empuje | CACHORRO | |
|---|---|---|---|---|
| Encabezamiento | <cuda/version> | <cuda/std/version> | <thrust/version.h> | <cub/version.h> |
| Versión principal | CCCL_MAJOR_VERSION | _LIBCUDACXX_CUDA_API_VERSION_MAJOR | THRUST_MAJOR_VERSION | CUB_MAJOR_VERSION |
| Versión menor | CCCL_MINOR_VERSION | _LIBCUDACXX_CUDA_API_VERSION_MINOR | THRUST_MINOR_VERSION | CUB_MINOR_VERSION |
| Versión de parche/subminor | CCCL_PATCH_VERSION | _LIBCUDACXX_CUDA_API_VERSION_PATCH | THRUST_SUBMINOR_VERSION | CUB_SUBMINOR_VERSION |
| Versión concatenada | CCCL_VERSION (MMMmmmppp) | _LIBCUDACXX_CUDA_API_VERSION (MMMmmmppp) | THRUST_VERSION (MMMmmmpp) | CUB_VERSION (MMMmmmpp) |
La interfaz binaria de la aplicación (ABI) es un conjunto de reglas para:
El ABI de una biblioteca incluye, pero no se limita a:
Un cambio de ruptura de ABI es cualquier cambio que resulte en un cambio en el ABI de una función o tipo en la API pública. Por ejemplo, agregar un nuevo miembro de datos a una estructura es un cambio de ruptura ABI, ya que cambia el tamaño del tipo.
En CCCL, las garantías sobre ABI son las siguientes:
thrust:: y los espacios de nombres cub:: pueden romper ABI en cualquier momento sin previo aviso.thrust:: y cub:: Symbols incluyen las arquitecturas CUDA utilizadas para la compilación. Por lo tanto, un símbolo thrust:: o cub:: puede tener un ABI diferente si:-x cu ) vs C ++ fuente ( -x cpp )cuda:: también pueden romper ABI en cualquier momento. Sin embargo, los símbolos cuda:: incrustan un número de versión ABI que se incrementa cada vez que se produce un descanso ABI. Se pueden admitir múltiples versiones ABI simultáneamente y, por lo tanto, los usuarios tienen la opción de volver a una versión ABI anterior. Para obtener más información, ver aquí.¿A quién debería preocuparse por Abi?
En general, los usuarios de CCCL solo deben preocuparse por los problemas de ABI al construir o usar un artefacto binario (como una biblioteca compartida) cuya API incluye directa o indirectamente los tipos proporcionados por CCCL.
Por ejemplo, considere si libA.so se construyó utilizando CCCL versión X y su API pública incluye una función como:
void foo (cuda::std::optional< int >); Si otra biblioteca, libB.so , se compila utilizando CCCL versión Y y usa foo de libA.so , entonces esto puede fallar si hubo un descanso ABI entre la Versión X e Y A diferencia de los cambios de ruptura de la API, los descansos ABI generalmente no requieren cambios en el código y solo requieren recompensar todo para usar la misma versión ABI.
Para obtener más información sobre ABI y por qué es importante, vea qué es ABI y ¿qué debe hacer C ++ al respecto?
Como se mencionó anteriormente, no todos los cambios posibles de ruptura de la fuente constituyen un cambio de ruptura que requeriría incrementar el número de versión principal de API de CCCL.
Se alienta a los usuarios que se adhieran a las siguientes pautas para minimizar el riesgo de interrupciones de accidentalmente dependiendo de partes de CCCL que no formen parte de la API pública:
thrust:: , cub:: , nv:: o cuda:: espacios de nombres a menos que se anote una excepción para un símbolo específico, por ejemplo, especialización cuda::std::iterator_traitsthrust:: , cub:: , cuda:: o nv:: Namespaces.thrust:: , cub:: , cuda:: o nv:: Spaces._ , __ o con detail en ningún lugar de su nombre, incluyendo un detail:: espacio de nombres o macro#include el archivo de encabezado que declara ese símbolo. En otras palabras, no confíe en encabezados incluidos implícitamente por otros encabezados.Las partes de esta sección se inspiraron en las pautas de compatibilidad de Abseil.
Haremos todo lo posible para notificar a los usuarios antes de hacer cambios de ruptura en la API pública, ABI o modificar las plataformas y compiladores compatibles.
Según corresponda, las deprecaciones vendrán en forma de advertencias programáticas que se pueden deshabilitar.
El período de deprecación dependerá del impacto del cambio, pero generalmente durará al menos 2 lanzamientos de versión menor.
¡Muy pronto!
Para obtener una descripción detallada de la tubería CI, consulte CI-Overview.md.
Proyectos relacionados con la misión de CCCL de hacer que CUDA sea más encantador:
¿Su proyecto usa CCCL? ¡Abra un PR para agregar su proyecto a esta lista!
