cccl
CCCL 2.6.1
| Руководство участника | Разработчик контейнеров | Раздор | Godbolt | GitHub Project | Документация |
|---|
Добро пожаловать в библиотеки Cuda Core Compute (CCCL), где наша миссия - сделать CUDA более восхитительным.
Этот репозиторий объединяет три важные библиотеки CUDA C ++ в единый, удобный репозиторий:
Цель CCCL - предоставить разработчикам CUDA C ++ строительные блоки, которые облегчают запись безопасного и эффективного кода. Объединение этих библиотек оптимирует процесс разработки и расширяет вашу способность использовать силу CUDA C ++. Для получения дополнительной информации о решении объединить эти проекты, см. Объявление здесь.
Концепция для CUDA Core Compute Libraries (CCCL) органически выросла из проектов тяги, Cub и LibcudAcxx, которые были разработаны независимо на протяжении многих лет с аналогичной целью: обеспечить высококачественные, высокопроизводительные и простые в использовании абстракции C ++ для разработчиков CUDA. Естественно, было много совпадений между тремя проектами, и стало ясно, что сообщество будет лучше обслуживать их в единый репозиторий.
Уклон - это библиотека параллельных алгоритмов C ++, которая вдохновила внедрение параллельных алгоритмов в стандартную библиотеку C ++. Высокий интерфейс Thrust значительно повышает производительность программиста, одновременно обеспечивая переносимость производительности между графическими процессорами и многокайновыми процессорами посредством настраиваемых бэкэндов, которые позволяют использовать несколько параллельных структур программирования (таких как CUDA, TBB и OpenMP).
Cub -это более низкая, специфичная для CUDA библиотека, предназначенная для параллельных алгоритмов скорости света во всех архитектурах GPU. В дополнение к алгоритмам всего устройства, он предоставляет кооперативные алгоритмы , такие как сокращение блоков и сканирование по всему деформации, предоставляя разработчикам ядра CUDA строительные блоки для создания настраиваемых ядер скорости.
Libcudacxx является стандартной библиотекой Cuda C ++. Он предоставляет реализацию стандартной библиотеки C ++, которая работает как в коде хоста, так и в коде устройства. Кроме того, он предоставляет абстракции для аппаратных функций, специфичных для CUDA, таких как примитивы синхронизации, управление кешем, атомику и многое другое.
Основной целью CCCL является выполнение аналогичной роли, которую стандартная библиотека C ++ заполняет для стандартного C ++: предоставить общепринятые инструменты скорости света для разработчиков CUDA C ++, что позволяет им сосредоточиться на решении проблем, которые имеют значение. Объединение этих проектов - первый шаг к достижению этой цели.
Это простой пример, демонстрирующий использование функциональности CCCL из тяги, Cub и Libcudacxx.
Он показывает, как использовать THRUST/CUB/LIBCUDACXX для реализации простого ядра с параллельным сокращением. Каждый блок потока вычисляет сумму подмножества массива с использованием cub::BlockReduce . Сумма каждого блока затем уменьшается до одного значения, используя атомное добавление через cuda::atomic_ref из libcudacxx.
Затем он показывает, как может быть сделано одно и то же сокращение, используя алгоритм reduce , и сравнивает результаты.
Попробуйте жить на 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 ;
}Все в CCCL только для заголовка. Поэтому пользователям нужно только беспокоиться о том, как они получают файлы заголовков и как они включают их в свою систему сборки.
Самый простой способ начать использование CCCL - через инструментарий CUDA, который включает заголовки CCCL. Когда вы компилируете с nvcc , он автоматически добавляет заголовки CCCL к вашему пути включения, чтобы вы могли просто #include любой заголовок CCCL в вашем коде без дополнительной конфигурации.
Если компиляция с другим компилятором, вам нужно будет обновить систему вашей системы сборки, включающий путь поиска, чтобы указать на заголовки CCCL в вашей установке CTK (например, /usr/local/cuda/include ).
# include < thrust/device_vector.h >
# include < cub/cub.cuh >
# include < cuda/std/atomic > Пользователям, которые хотят оставаться на переднем крае разработки CCCL, рекомендуется использовать CCCL из GitHub. Использование более новой версии CCCL с более старой версией Cuda Toolkit поддерживается, но не наоборот. Для получения полной информации о совместимости между CCCL и The Cuda Toolkit, см. В нашу поддержку платформы.
Все в CCCL только для заголовка, поэтому клонирование и включение его в простом проекте так же просто, как и следующее:
git clone https://github.com/NVIDIA/cccl.git
nvcc -Icccl/thrust -Icccl/libcudacxx/include -Icccl/cub main.cu -o mainПримечание Использование
-I, а не-isystem, чтобы избежать столкновений с заголовками CCCL, неявно включеннымиnvccиз CUDA Toolkit. Все заголовки CCCL используют#pragma system_header, чтобы гарантировать, что предупреждения все еще будут замолчать, как будто используя-isystem, см. #527 для получения дополнительной информации.
Минимальная сборка, которая генерирует только правила установки, может быть настроена с использованием предустановки 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 Чтобы включить экспериментальные библиотеки в установку, используйте install-unstable и каталог сборки.
Чтобы установить только экспериментальные библиотеки, используйте install-unstable-only для установки и строительство.
CCCL также предоставляет пакеты Conda каждого выпуска через канал conda-forge :
conda config --add channels conda-forge
conda install cccl Это установит последнюю CCCL в среду Conda $CONDA_PREFIX/include/ и $CONDA_PREFIX/lib/cmake/ каталоги. Он обнаруживается CMAKE через find_package(CCCL) и может использоваться любыми компиляторами в среде Conda. Для получения дополнительной информации см. Это введение в Conda-Forge.
Если вы хотите использовать ту же версию CCCL, которая поставлялась с определенным инструментом CUDA, например, CUDA 12.4, вы можете установить CCCL с:
conda config --add channels conda-forge
conda install cuda-cccl cuda-version=12.4 MetaPackage cuda-cccl устанавливает версию cccl , которая поставляется с инструментарием CUDA, соответствующей cuda-version . Если вы хотите обновиться до последней cccl после установки cuda-cccl , удалите cuda-cccl перед обновлением cccl :
conda uninstall cuda-cccl
conda install -c conda-forge ccclОбратите внимание, что есть также пакеты Conda с такими именами, как
cuda-cccl_linux-64. Эти пакеты содержат версии CCCL, отправленные как часть инструментария CUDA, но предназначены для внутреннего использования инструментом CUDA. Вместо этого установитеccclилиcuda-ccclдля совместимости с компиляторами Conda. Для получения дополнительной информации см. Рецепт CCCL Conda-Forge.
CCCL использует CMAKE для всей инфраструктуры сборки и установки, включая тесты, а также цели для связи в других проектах CMAKE. Следовательно, Cmake - это рекомендуемый способ интеграции CCCL в другой проект.
Для полного примера того, как это сделать, используя диспетчер пакетов Cmake, см. Наш базовый пример проекта.
Другие системы сборки должны работать, но тестируется только Cmake. Взносы для упрощения интеграции CCCL в другие системы сборки приветствуются.
Заинтересованы в улучшении CCCL? Ознакомьтесь с нашим руководством для получения комплексного обзора всего, что вам нужно знать, чтобы настроить среду разработки, внести изменения, запустить тесты и отправить PR.
Цель: в этом разделе описывается, где пользователи могут ожидать, что CCCL будет успешно компилировать и работать.
В целом, CCCL должен работать везде, где поддерживается инструментарий CUDA, однако дьявол находится в деталях. Приведенные ниже разделы описывают детали поддержки и тестирования для различных версий инструментария CUDA, компиляторов хоста и диалектов C ++.
Краткое содержание:
Пользователям CCCL рекомендуется использовать последние улучшения и «Live at Head», всегда используя новейшую версию CCCL. Для беспрепятственного опыта вы можете обновить CCCL независимо от всего инструментария CUDA. Это возможно, потому что CCCL поддерживает обратную совместимость с последним выпуском патча каждого незначительного выпуска CTK как из текущей, так и предыдущей серии основных версий. В некоторых исключительных случаях минимально поддерживаемая второстепенная версия релиза Cuda Toolkit, возможно, должна быть новой, чем самый старый выпуск в его основной серии версий. Например, CCCL требует минимальной поддерживаемой версии 11.1 из серии 11.x из -за неизбежной проблемы компилятора, присутствующей в CTK 11.0.
Когда выпущен новый крупный CTK, мы отдаем поддержку самой старой основной версии.
| CCCL версия | Поддерживает версию инструментов CUDA |
|---|---|
| 2.x. | 11.1 - 11.8, 12.x (только последние выпуски патча) |
| 3.x (будущее) | 12.x, 13.x (только последние выпуски патча) |
Код с благополучием с использованием последней CCCL должен успешно скомпилировать и работать с любой поддерживаемой версией CTK. Исключения могут возникнуть для новых функций, которые зависят от новых функций CTK, поэтому эти функции не будут работать на более старых версиях CTK. Например, поддержка C ++ 20 не была добавлена в nvcc до CUDA 12.0, поэтому функции CCCL, которые зависят от C ++ 20, не будут работать с CTK 11.x.
Пользователи могут интегрировать более новую версию CCCL в более старый CTK, но не наоборот. Это означает, что более старая версия CCCL не совместима с более новой CTK. Другими словами, CCCL никогда не совместима с инструментом CUDA.
В таблице ниже суммируется совместимость CTK и CCCL:
| CTK версия | Включена версия CCCL | Желаемый CCCL | Поддерживается? | Примечания |
|---|---|---|---|---|
CTK XY | CCCL MAJOR.MINOR | CCCL MAJOR.MINOR+n | ✅ | Некоторые новые функции могут не сработать |
CTK XY | CCCL MAJOR.MINOR | CCCL MAJOR+1.MINOR | ✅ | Возможные перерывы; Некоторые новые функции могут быть недоступны |
CTK XY | CCCL MAJOR.MINOR | CCCL MAJOR+2.MINOR | CCCL поддерживает только две основные версии CTK | |
CTK XY | CCCL MAJOR.MINOR | CCCL MAJOR.MINOR-n | CCCL не совместим | |
CTK XY | CCCL MAJOR.MINOR | CCCL MAJOR-n.MINOR | CCCL не совместим |
Для получения дополнительной информации о версии CCCL, совместимости API/ABI и разрывах изменения см. В разделе «Версия ниже».
Если не указано иное, CCCL поддерживает все те же операционные системы, что и инструментарий CUDA, которые задокументированы здесь:
Если не указано иное, CCCL поддерживает те же компиляторы хоста, что и последний инструментарий CUDA, которые задокументированы здесь:
При использовании более старых инструментов CUDA мы также поддерживаем только компиляторы хоста последнего инструментария CUDA, но, по крайней мере, самый последний компилятор хоста любого поддерживаемого более старого инструментария CUDA.
Мы можем сохранить поддержку дополнительных компиляторов и принять соответствующие исправления от сообщества с разумными исправлениями. Но мы не будем инвестировать значительное время в тридневные или исправления проблем для старых компиляторов.
В духе «Вы поддерживаете только то, что вы тестируете», посмотрите наш обзор CI для получения дополнительной информации о том, что мы тестируем.
Если не указано иное, CCCL поддерживает все одинаковые архитектуры/вычислительные возможности GPU, что и инструментарий CUDA, которые задокументированы здесь: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#compute-capapability
Обратите внимание, что некоторые функции могут поддерживать только определенные архитектуры/вычислительные возможности.
Стратегия тестирования CCCL достигает баланса между тестированием как можно большим количеством конфигураций и поддержанием разумного времени CI.
Для версий инструментов CUDA тестирование проводится как в самых старых, так и в новейших поддерживаемых версиях. Например, если последняя версия инструментария CUDA составляет 12,3, тесты проводятся против 11.1 и 12.3. Для каждой версии CUDA сборки завершаются против всех поддерживаемых компиляторов хоста со всеми поддерживаемыми диалектами C ++.
Стратегия тестирования и матрица постоянно развиваются. Матрица, определенная в файле ci/matrix.yaml является окончательным источником истины. Для получения дополнительной информации о нашем трубопроводе CI см. Здесь.
Цель: в этом разделе описывается, как CCCL версирует, гарантии API/ABI и рекомендациям по совместимости для минимизации головных болей обновления.
Краткое содержание
cub:: или thrust:: именаcuda:: имена могут произойти в любое время, но будут отражены путем увеличения версии ABI, которая встроена в встроенное пространство имен для всех символов cuda:: . Многочисленные версии ABI могут быть поддержаны одновременно.ПРИМЕЧАНИЕ. До слияния тяги, Cub и Libcudacxx в этот репозиторий каждая библиотека была самостоятельно версирована в соответствии с семантическими версиями. Начиная с выпуска 2.1, все три библиотеки синхронизировали свои версии выпуска в их отдельных репозиториях. Двигаясь вперед, CCCL будет продолжать выпускать под одной семантической версией, а 2.2.0 - первый выпуск из репозитория NVIDIA/CCCL.
Разрывное изменение - это изменение явно поддерживаемой функциональности между выпущенными версиями, которые потребовали бы, чтобы пользователь выполнял работу, чтобы перейти на более новую версию.
В пределах, любое изменение может где -то разбить кого -то. В результате не все возможные изменения разрыва источника считаются нарушающими изменения в публичном API, которые гарантируют увеличение основной семантической версии.
Приведенные ниже разделы описывают детали нарушения изменений в API и ABI CCCL.
Общественный API CCCL является полностью в том, что функциональность преднамеренно подвергается предоставлению полезности библиотеки.
Другими словами, публичный API CCCL выходит за рамки просто функциональных подписей и включает (но не ограничивается):
Кроме того, публичный API CCCL не включает ни одно из следующих:
_ или __detail , включая detail:: пространство имен или макросdetail/ каталоге или его подчиненииВ общем, цель состоит в том, чтобы не нарушать что -либо в публичном API. Такие изменения вносятся только в том случае, если они предлагают пользователям лучшую производительность, более легкие для понимания API и/или более последовательные API.
Любое нарушение изменения публичного API потребует выталкивания основного номера версии CCCL. В соответствии с совместимостью с малой версией CUDA изменение разрыва API и CCCL -основные варки версии будут происходить только с новой крупной версией версии CUDA Toolkit.
Все, что не является частью публичного API, может измениться в любое время без предупреждения.
Общественный API всех компонентов CCCL имеет единую семантическую версию MAJOR.MINOR.PATCH .
Поддерживается только последняя выпущенная версия. Как правило, функции и исправления ошибок не возвращаются в ранее выпущенную версию или филиалы.
Предпочтительным методом запроса версии является использование CCCL_[MAJOR/MINOR/PATCH_]VERSION как описано ниже. Для обратной совместимости определения версий TRUST/CUB/LibcudACXXX доступны и всегда будут соответствовать CCCL_VERSION . Обратите внимание, что тяга/Cub использует схему MMMmmmpp , тогда как CCCL и Libcudacxx используют MMMmmmppp .
| Cccl | libcudacxx | Толкать | Нога | |
|---|---|---|---|---|
| Заголовок | <cuda/version> | <cuda/std/version> | <thrust/version.h> | <cub/version.h> |
| Основная версия | CCCL_MAJOR_VERSION | _LIBCUDACXX_CUDA_API_VERSION_MAJOR | THRUST_MAJOR_VERSION | CUB_MAJOR_VERSION |
| Незначительная версия | CCCL_MINOR_VERSION | _LIBCUDACXX_CUDA_API_VERSION_MINOR | THRUST_MINOR_VERSION | CUB_MINOR_VERSION |
| Патч/подчиненная версия | CCCL_PATCH_VERSION | _LIBCUDACXX_CUDA_API_VERSION_PATCH | THRUST_SUBMINOR_VERSION | CUB_SUBMINOR_VERSION |
| Согласованная версия | CCCL_VERSION (MMMmmmppp) | _LIBCUDACXX_CUDA_API_VERSION (MMMmmmppp) | THRUST_VERSION (MMMmmmpp) | CUB_VERSION (MMMmmmpp) |
Двоичный интерфейс приложения (ABI) - это набор правил для:
ABI библиотеки включает, но не ограничивается:
Разбиваемое изменение ABI - это любое изменение, которое приводит к изменению ABI функции или типа в публичном API. Например, добавление нового элемента данных в структуру является разбившимся изменением ABI, поскольку он меняет размер типа.
В CCCL гарантии об ABI заключаются в следующем:
thrust:: и cub:: Пространства имен могут сломаться ABI в любое время без предупреждения.thrust:: и cub:: Символы включают архитектуры CUDA, используемые для компиляции. Следовательно, thrust:: или cub:: Символ может иметь другой ABI, если:-x cu ) против C ++ Source ( -x cpp )cuda:: также могут сломать ABI в любое время. Тем не менее, cuda:: Символы внедряют номер версии ABI, который увеличивается всякий раз, когда происходит разрыв ABI. Несколько версий ABI могут быть поддержаны одновременно, и поэтому пользователи имеют возможность вернуться к предыдущей версии ABI. Для получения дополнительной информации см. Здесь.Кто должен заботиться об Аби?
В целом, пользователям CCCL должны беспокоиться только о проблемах ABI только при создании или использовании бинарного артефакта (например, общей библиотеки), API, прямо или косвенно включает типы, предоставляемые CCCL.
Например, рассмотрим, был ли libA.so построен с использованием CCCL версии X , а его публичный API включает в себя функцию, такую как:
void foo (cuda::std::optional< int >); Если другая библиотека, libB.so , составлена с использованием версии CCCL Y и использует foo из libA.so , то это может выйти из строя, если между версией X и Y произошел разрыв ABI. В отличие от изменений разрыва API, разрывы ABI обычно не требуют изменений кода и требуют только перекомпилирования всего, чтобы использовать одну и ту же версию ABI.
Чтобы узнать больше об ABI и почему это важно, посмотрите, что такое ABI, и что C ++ с этим делать?.
Как упомянуто выше, не все возможные изменения разрыва источника представляют собой нарушающее изменение, которое потребовало бы увеличения основного номера версии API CCCL.
Пользователям рекомендуется придерживаться следующих руководящих принципов, чтобы минимизировать риск сбоев случайно в зависимости от частей CCCL, которые не являются частью публичного API:
thrust:: , cub:: , nv:: или cuda:: Пространства имен, если исключение не отмечено для конкретного символа, например, cuda::std::iterator_traitsthrust:: , cub:: , cuda:: , или nv:: имена.thrust:: , cub:: , cuda:: или nv:: имена._ , __ или с detail в любом месте его имени, включая detail:: пространство имен или макрос#include файл заголовка, который объявляет этот символ. Другими словами, не полагайтесь на заголовки, неявно включенные другими заголовками.Части этого раздела были вдохновлены руководящими принципами совместимости Abseil.
Мы сделаем все возможное, чтобы уведомить пользователей до внесения каких -либо прерывищихся изменений в публичном API, ABI или изменять поддерживаемые платформы и компиляторы.
При необходимости, в форме программных предупреждений, которые могут быть отключены.
Период деформации будет зависеть от воздействия изменения, но обычно будет длиться не менее 2 незначительных версий.
Вскоре!
Подробный обзор трубопровода CI см. Ci-Overview.md.
Проекты, которые связаны с миссией CCCL, сделать CUDA более восхитительным:
Ваш проект использует CCCL? Откройте пиар, чтобы добавить свой проект в этот список!
