cccl
CCCL 2.6.1
| Guia do colaborador | Recipientes de dev | Discórdia | Godbolt | Projeto Github | Documentação |
|---|
Bem -vindo às bibliotecas de computação do CUDA Core (CCCL), onde nossa missão é tornar o CUDA mais agradável.
Este repositório unifica três bibliotecas essenciais do CUDA C ++ em um repositório único e conveniente:
O objetivo do CCCL é fornecer aos desenvolvedores CUDA C ++ os blocos de construção que facilitam a redação de código seguro e eficiente. Reunindo essas bibliotecas simplifica seu processo de desenvolvimento e amplia sua capacidade de alavancar o poder do CUDA C ++. Para obter mais informações sobre a decisão de unificar esses projetos, consulte o anúncio aqui.
O conceito para as bibliotecas de computação do CUDA Core (CCCL) cresceu organicamente a partir de projetos que foram desenvolvidos independentemente ao longo dos anos com um objetivo semelhante: fornecer abstrações C ++ de alta qualidade, de alto desempenho e fáceis de usar para desenvolvedores de CUDA. Naturalmente, houve muita sobreposição entre os três projetos, e ficou claro que a comunidade seria melhor servida, unificando -os em um único repositório.
O impulso é a biblioteca de algoritmos paralela C ++ que inspirou a introdução de algoritmos paralelos na biblioteca padrão C ++. A interface de alto nível da Thrust aprimora bastante a produtividade do programador, permitindo a portabilidade do desempenho entre as GPUs e as CPUs multicore por meio de backnds configuráveis que permitem o uso de várias estruturas de programação paralelas (como CUDA, TBB e OpenMP).
O Cub é uma biblioteca específica de CUDA de nível inferior, projetada para algoritmos paralelos de velocidade de luz em todas as arquiteturas da GPU. Além dos algoritmos em todo o dispositivo, ele fornece algoritmos cooperativos , como redução em todo o bloco e varredura em toda a urdidura, fornecendo aos desenvolvedores do kernel da CUDA blocos de construção para criar kernels personalizados e de velocidade.
Libcudacxx é a biblioteca padrão CUDA C ++. Ele fornece uma implementação da biblioteca padrão C ++ que funciona no código do host e do dispositivo. Além disso, ele fornece abstrações para recursos de hardware específico do CUDA, como primitivas de sincronização, controle de cache, atômicos e muito mais.
O principal objetivo do CCCL é preencher uma função semelhante que a biblioteca C ++ padrão preencha para as ferramentas padrão de C ++: fornecer ferramentas de velocidade de uso geral para desenvolvedores de CUDA C ++, permitindo que eles se concentrem na solução dos problemas que importam. Unificar esses projetos é o primeiro passo para realizar esse objetivo.
Este é um exemplo simples demonstrando o uso da funcionalidade CCCL da impulso, Cub e Libcudacxx.
Ele mostra como usar o impulso/cub/libcudacxx para implementar um kernel de redução paralela simples. Cada bloco de thread calcula a soma de um subconjunto da matriz usando cub::BlockReduce . A soma de cada bloco é então reduzida a um único valor usando um atômico adicionado via cuda::atomic_ref de libcudacxx.
Em seguida, mostra como a mesma redução pode ser feita usando o algoritmo reduce do Thrust e compara os resultados.
Experimente ao vivo em 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 ;
}Tudo no CCCL é apenas para cabeçalho. Portanto, os usuários precisam se preocupar apenas com a forma como obtêm os arquivos de cabeçalho e como os incorporam ao sistema de construção.
A maneira mais fácil de começar a usar o CCCL é através do CUDA Toolkit, que inclui os cabeçalhos do CCCL. Quando você compila com nvcc , ele adiciona automaticamente os cabeçalhos CCCL ao seu caminho de inclusão para que você possa simplesmente #include qualquer cabeçalho do CCCL no seu código, sem nenhuma configuração adicional necessária.
Se compilar com outro compilador, você precisará atualizar o caminho de pesquisa do seu sistema de compilação para apontar para os cabeçalhos do CCCL na instalação do CTK (por exemplo, /usr/local/cuda/include ).
# include < thrust/device_vector.h >
# include < cub/cub.cuh >
# include < cuda/std/atomic > Os usuários que desejam permanecer na vanguarda do desenvolvimento do CCCL são incentivados a usar o CCCL do GitHub. Usando uma versão mais recente do CCCL com uma versão mais antiga do kit de ferramentas CUDA é suportada, mas não o contrário. Para obter informações completas sobre a compatibilidade entre o CCCL e o CUDA Toolkit, consulte o suporte à nossa plataforma.
Tudo no CCCL é somente para o cabeçalho, então a clonagem e incluí-lo em um projeto simples é tão fácil quanto o seguinte:
git clone https://github.com/NVIDIA/cccl.git
nvcc -Icccl/thrust -Icccl/libcudacxx/include -Icccl/cub main.cu -o mainNota Use
-Ie não-isystempara evitar colisões com os cabeçalhos do CCCL implicitamente incluídos pelonvccno kit de ferramentas CUDA. Todos os cabeçalhos do CCCL usam#pragma system_headerpara garantir que os avisos ainda sejam silenciados como se estivessem usando-isystem, consulte #527 para obter mais informações.
Uma compilação mínima que gera apenas regras de instalação pode ser configurada usando a predefinição install do 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 experimentais na instalação, use o diretório de predefinição e construção de install-unstable .
Para instalar apenas as bibliotecas experimentais, use o diretório install-unstable-only de instalação e construção.
O CCCL também fornece pacotes de CONDA de cada lançamento através do canal conda-forge :
conda config --add channels conda-forge
conda install cccl Isso instalará o mais recente CCCL com o $CONDA_PREFIX/include/ $CONDA_PREFIX/lib/cmake/ DIRETORITOS. É descoberta pelo CMake via find_package(CCCL) e pode ser usado por quaisquer compiladores no ambiente do CONDA. Para mais informações, consulte esta introdução a Conde-Forge.
Se você deseja usar a mesma versão CCCL que enviou com um kit de ferramentas CUDA, por exemplo, CUDA 12.4, você pode instalar o CCCL com:
conda config --add channels conda-forge
conda install cuda-cccl cuda-version=12.4 O metapackage cuda-cccl instala a versão cccl enviada com o kit de ferramentas CUDA correspondente à cuda-version . Se você deseja atualizar para o cccl mais recente depois de instalar cuda-cccl , desinstale cuda-cccl antes de atualizar cccl :
conda uninstall cuda-cccl
conda install -c conda-forge ccclObserve que também existem pacotes do CONDA com nomes como
cuda-cccl_linux-64. Esses pacotes contêm as versões CCCL enviadas como parte do kit de ferramentas CUDA, mas são projetadas para uso interno pelo kit de ferramentas CUDA. Instalecccloucuda-cccl, para compatibilidade com compiladores do CONDA. Para obter mais informações, consulte a receita do CCCL CONDA-FORGE.
O CCCL usa o CMake para toda a infraestrutura de compilação e instalação, incluindo testes, bem como alvos para vincular em outros projetos de CMake. Portanto, o Cmake é a maneira recomendada de integrar o CCCL em outro projeto.
Para um exemplo completo de como fazer isso usando o CMake Package Manager, consulte nosso projeto de exemplo básico.
Outros sistemas de construção devem funcionar, mas apenas o CMake é testado. As contribuições para simplificar a integração do CCCL em outros sistemas de construção são bem -vindas.
Interessado em contribuir para melhorar o CCCL? Confira nosso guia contribuinte para obter uma visão geral abrangente de tudo o que você precisa saber para configurar seu ambiente de desenvolvimento, fazer alterações, executar testes e enviar um PR.
Objetivo: Esta seção descreve onde os usuários podem esperar que o CCCL compile e seja executado com sucesso.
Em geral, o CCCL deve trabalhar em todos os lugares em que o kit de ferramentas CUDA é suportado, no entanto, o diabo está nos detalhes. As seções abaixo descrevem os detalhes de suporte e teste para diferentes versões do kit de ferramentas CUDA, compiladores de host e dialetos C ++.
Resumo:
Os usuários da CCCL são incentivados a capitalizar os mais recentes aprimoramentos e "Live at Head", sempre usando a versão mais recente do CCCL. Para uma experiência perfeita, você pode atualizar o CCCL independentemente de todo o kit de ferramentas do CUDA. Isso é possível porque o CCCL mantém a compatibilidade com a versão mais recente do patch de cada liberação menor do CTK da série de versão principal atual e anterior. Em alguns casos excepcionais, a versão menor mínima suportada pelo lançamento do CUDA Toolkit pode precisar ser mais recente do que a versão mais antiga em sua principal série de versão. Por exemplo, o CCCL requer uma versão mínima suportada de 11.1 da série 11.x devido a uma questão inevitável do compilador presente no CTK 11.0.
Quando um novo CTK importante é lançado, lançamos suporte para a versão principal mais antiga suportada.
| Versão CCCL | Suporta a versão CUDA Toolkit |
|---|---|
| 2.x | 11.1 - 11.8, 12.x (apenas as mais recentes lançamentos de patches) |
| 3.x (futuro) | 12.x, 13.x (apenas as mais recentes lançamentos de patches) |
O código bem-comportado usando o CCCL mais recente deve compilar e executar com sucesso com qualquer versão CTK suportada. Exceções podem ocorrer para novos recursos que dependem dos novos recursos do CTK, para que esses recursos não funcionassem em versões mais antigas do CTK. Por exemplo, o suporte C ++ 20 não foi adicionado ao nvcc até o CUDA 12.0, portanto, os recursos do CCCL que dependem do C ++ 20 não funcionariam com o CTK 11.X.
Os usuários podem integrar uma versão mais recente do CCCL em um CTK mais antigo, mas não o contrário. Isso significa que uma versão mais antiga do CCCL não é compatível com um CTK mais recente. Em outras palavras, o CCCL nunca é compatível com o kit de ferramentas CUDA.
A tabela abaixo resume a compatibilidade do CTK e CCCL:
| Versão ctk | Versão CCCL incluída | CCCL desejado | Suportado? | Notas |
|---|---|---|---|---|
CTK XY | CCCL MAJOR.MINOR | Cccl MAJOR.MINOR+n | ✅ | Alguns novos recursos podem não funcionar |
CTK XY | CCCL MAJOR.MINOR | CCCL MAJOR+1.MINOR | ✅ | Possíveis quebras; Alguns novos recursos podem não estar disponíveis |
CTK XY | CCCL MAJOR.MINOR | CCCL MAJOR+2.MINOR | CCCL suporta apenas duas versões principais do CTK | |
CTK XY | CCCL MAJOR.MINOR | Cccl MAJOR.MINOR-n | CCCL não é compatível com a frente | |
CTK XY | CCCL MAJOR.MINOR | CCCL MAJOR-n.MINOR | CCCL não é compatível com a frente |
Para obter mais informações sobre a versão do CCCL, a compatibilidade da API/ABI e as mudanças de quebra, consulte a seção de versão abaixo.
Salvo indicação em contrário, o CCCL suporta todos os mesmos sistemas operacionais que o CUDA Toolkit, que estão documentados aqui:
Salvo indicação em contrário, o CCCL suporta os mesmos compiladores de host que o mais recente kit de ferramentas CUDA, que estão documentados aqui:
Ao usar kits de ferramentas CUDA mais antigos, também suportamos apenas os compiladores host do mais recente kit de ferramentas do CUDA, mas pelo menos o compilador host mais recente de qualquer kit de ferramentas CUDA mais recente.
Podemos manter o apoio de compiladores adicionais e aceitaremos patches correspondentes da comunidade com correções razoáveis. Mas não investiremos um tempo significativo na trincagem ou na fixação de problemas para compiladores mais antigos.
No espírito de "Você só apóia o que testar", consulte nossa visão geral do CI para obter mais informações sobre exatamente o que testamos.
Salvo indicação em contrário, o CCCL suporta todos os mesmos recursos de arquiteturas/computação da GPU que o CUDA Toolkit, que estão documentados aqui: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#compute-capability
Observe que alguns recursos podem suportar apenas determinadas arquiteturas/recursos de computação.
A estratégia de teste da CCCL atinge um equilíbrio entre testar o maior número possível de configurações e a manutenção de tempos de IC razoáveis.
Para versões CUDA Toolkit, o teste é feito contra as versões mais antigas e mais recentes. Por exemplo, se a versão mais recente do kit de ferramentas CUDA for 12.3, os testes serão realizados contra 11.1 e 12.3. Para cada versão do CUDA, as compilações são concluídas contra todos os compiladores de host suportados com todos os dialetos C ++ suportados.
A estratégia de teste e a matriz estão em constante evolução. A matriz definida no arquivo ci/matrix.yaml é a fonte definitiva de verdade. Para obter mais informações sobre o nosso pipeline de IC, consulte aqui.
Objetivo: Esta seção descreve como o CCCL é versão em versão, garantias de estabilidade da API/ABI e diretrizes de compatibilidade para minimizar as dores de cabeça de atualização.
Resumo
cub:: ou thrust:: namespacescuda:: podem acontecer a qualquer momento, mas serão refletidas incrementando a versão ABI que é incorporada em um espaço de nome em linha para todos os símbolos cuda:: . Várias versões da ABI podem ser suportadas simultaneamente.Nota: Antes de mesclar a fusão, Cub e Libcudacxx nesse repositório, cada biblioteca foi de forma independente de acordo com a versão semântica. Começando com a versão 2.1, todas as três bibliotecas sincronizaram suas versões de liberação em seus repositórios separados. Avançando, o CCCL continuará sendo lançado sob uma única versão semântica, com 2.2.0 sendo a primeira versão do repositório NVIDIA/CCCL.
Uma mudança de ruptura é uma alteração para funcionalidade explicitamente suportada entre versões lançadas que exigiriam que um usuário trabalhe para atualizar para a versão mais recente.
No limite, qualquer mudança tem o potencial de quebrar alguém em algum lugar. Como resultado, nem todas as alterações possíveis de quebra de fonte são consideradas mudanças na API pública que justificam a maior versão semântica.
As seções abaixo descrevem os detalhes de interromper as alterações na API e ABI da CCCL.
A API pública da CCCL é a totalidade da funcionalidade intencionalmente exposta para fornecer a utilidade da biblioteca.
Em outras palavras, a API pública da CCCL vai além de apenas assinaturas de funções e inclui (mas não se limita a):
Além disso, a API pública da CCCL não inclui nenhum dos seguintes:
_ ou __detail , incluindo o detail:: namespace ou uma macrodetail/ diretório ou subdiretórioEm geral, o objetivo é evitar quebrar qualquer coisa na API pública. Tais alterações são feitas apenas se oferecer aos usuários um melhor desempenho, APIs mais fáceis de entender e/ou APIs mais consistentes.
Qualquer mudança de ruptura na API pública exigirá o número de versão principal da CCCL. De acordo com a compatibilidade da versão menor do CUDA, as mudanças de quebra de API e os principais solavancos da versão CCCL só ocorrerão coincidindo com uma nova versão principal da versão do kit de ferramentas CUDA.
Qualquer coisa que não parte da API pública possa mudar a qualquer momento sem aviso prévio.
A API pública de todos os componentes da CCCL compartilham uma versão semântica unificada do MAJOR.MINOR.PATCH .
Somente a versão mais recentemente lançada é suportada. Como regra, os recursos e as correções de bugs não são retrucados para a versão ou ramificações lançadas anteriormente.
O método preferido para consultar a versão é usar CCCL_[MAJOR/MINOR/PATCH_]VERSION conforme descrito abaixo. Para compatibilidade com versões anteriores, as definições da versão de impulso/cub/libcudacxxx estão disponíveis e sempre serão consistentes com CCCL_VERSION . Observe que o impulso/cub usa um esquema MMMmmmpp , enquanto o cccl e libcudacxx usam MMMmmmppp .
| Cccl | libcudacxx | Impulso | FILHOTE | |
|---|---|---|---|---|
| Cabeçalho | <cuda/version> | <cuda/std/version> | <thrust/version.h> | <cub/version.h> |
| Versão principal | CCCL_MAJOR_VERSION | _LIBCUDACXX_CUDA_API_VERSION_MAJOR | THRUST_MAJOR_VERSION | CUB_MAJOR_VERSION |
| Versão menor | CCCL_MINOR_VERSION | _LIBCUDACXX_CUDA_API_VERSION_MINOR | THRUST_MINOR_VERSION | CUB_MINOR_VERSION |
| Versão de patch/subminador | CCCL_PATCH_VERSION | _LIBCUDACXX_CUDA_API_VERSION_PATCH | THRUST_SUBMINOR_VERSION | CUB_SUBMINOR_VERSION |
| Versão concatenada | CCCL_VERSION (MMMmmmppp) | _LIBCUDACXX_CUDA_API_VERSION (MMMmmmppp) | THRUST_VERSION (MMMmmmpp) | CUB_VERSION (MMMmmmpp) |
A interface binária do aplicativo (ABI) é um conjunto de regras para:
O ABI de uma biblioteca inclui, mas não está limitado a:
Uma mudança de quebra de ABI é qualquer alteração que resulte em uma alteração no ABI de uma função ou tipo na API pública. Por exemplo, adicionar um novo membro de dados a uma estrutura é uma mudança de quebra de ABI, pois altera o tamanho do tipo.
No CCCL, as garantias sobre a ABI são as seguintes:
thrust:: e cub:: namespaces podem quebrar a ABI a qualquer momento sem aviso prévio.thrust:: e cub:: símbolos inclui as arquiteturas CUDA usadas para compilação. Portanto, um thrust:: ou cub:: símbolo pode ter um ABI diferente se:-x cu ) vs C ++ Source ( -x cpp )cuda:: também podem quebrar a ABI a qualquer momento. No entanto, os símbolos cuda:: incorporam um número de versão da ABI que é incrementado sempre que ocorre uma quebra de ABI. Várias versões da ABI podem ser suportadas simultaneamente e, portanto, os usuários têm a opção de reverter para uma versão anterior da ABI. Para mais informações, consulte aqui.Quem deve se preocupar com a ABI?
Em geral, os usuários do CCCL só precisam se preocupar com os problemas da ABI ao criar ou usar um artefato binário (como uma biblioteca compartilhada) cuja API inclui direta ou indiretamente os tipos fornecidos pelo CCCL.
Por exemplo, considere se libA.so foi construído usando o CCCL Versão X e sua API pública inclui uma função como:
void foo (cuda::std::optional< int >); Se outra biblioteca, libB.so , for compilada usando o CCCL Versão Y e usa foo do libA.so , isso pode falhar se houve uma quebra de ABI entre a versão X e Y Ao contrário das mudanças de quebra de API, as quebras de ABI geralmente não requerem alterações de código e exigem apenas recompilar tudo para usar a mesma versão da ABI.
Para saber mais sobre a ABI e por que é importante, veja o que é ABI e o que o C ++ deve fazer sobre isso?.
Como mencionado acima, nem todas as mudanças possíveis de quebra de fonte constituem uma mudança de ruptura que exigiria o número de versão principal da API da CCCL.
Os usuários são incentivados a aderir às seguintes diretrizes, a fim de minimizar o risco de interrupções de acidentalmente, dependendo de partes do CCCL que não fazem parte da API pública:
thrust:: , cub:: , nv:: ou cuda:: namespaces, a menos que uma exceção seja observada para um símbolo específico, por exemplo, especializar cuda::std::iterator_traitsthrust:: , cub:: , cuda:: ou nv:: namespaces.thrust:: , cub:: , cuda:: , ou nv:: namespaces._ , __ ou com detail em qualquer lugar em seu nome, incluindo um detail:: namespace ou macro#include o arquivo de cabeçalho que declara esse símbolo. Em outras palavras, não confie nos cabeçalhos implicitamente incluídos por outros cabeçalhos.Partes desta seção foram inspiradas nas diretrizes de compatibilidade de Abseil.
Faremos o possível para notificar os usuários antes de fazer alterações de quebra na API pública, ABI ou modificar as plataformas e compiladores suportados.
Conforme apropriado, as deprecações virão na forma de avisos programáticos que podem ser desativados.
O período de depreciação dependerá do impacto da mudança, mas geralmente dura pelo menos 2 liberações de versão menor.
Em breve!
Para uma visão geral detalhada do pipeline do CI, consulte Ci-Overview.md.
Projetos relacionados à missão da CCCL de tornar o CUDA mais agradável:
Seu projeto usa o CCCL? Abra um PR para adicionar seu projeto a esta lista!
