| 寄稿者ガイド | 開発容器 | 不和 | ゴッドボルト | Githubプロジェクト | ドキュメント |
|---|
Cuda Core Compute Libraries(CCCL)へようこそ。私たちの使命は、CUDAをより楽しいものにすることです。
このリポジトリは、3つの重要なCUDA C ++ライブラリを単一の便利なリポジトリに統合します。
CCCLの目標は、CUDA C ++開発者に、安全で効率的なコードの作成を容易にするビルディングブロックを提供することです。これらのライブラリをまとめると、開発プロセスが合理化され、CUDA C ++の力を活用する能力が広がります。これらのプロジェクトを統一する決定の詳細については、こちらの発表を参照してください。
CUDAコアコンピューティングライブラリ(CCCL)のコンセプトは、CUDA開発者に高品質で高性能で使いやすいC ++抽象化を提供するために、長年にわたって独立して開発された推力、カブ、およびLibCudacxxプロジェクトから有機的に成長しました。当然のことながら、3つのプロジェクトの中で多くの重複があり、コミュニティがそれらを単一のリポジトリに統合することでより良いサービスを提供することが明らかになりました。
スラストは、C ++並列アルゴリズムライブラリで、C ++標準ライブラリへの並列アルゴリズムの導入に影響を与えました。スラストのハイレベルインターフェイスは、プログラマーの生産性を大幅に向上させ、複数の並列プログラミングフレームワーク(CUDA、TBB、OpenMPなど)を使用できる構成可能なバックエンドを介して、GPUとマルチコアCPU間のパフォーマンスポータビリティを可能にします。
Cubは、すべてのGPUアーキテクチャにわたって、速度の平行アルゴリズム用に設計された下位レベルのCUDA固有のライブラリです。デバイス全体のアルゴリズムに加えて、ブロック全体の削減やワープ全体のスキャンなどの共同アルゴリズムを提供し、CUDAカーネル開発者にビルディングブロックを提供し、ライトの速度、カスタムカーネルを作成します。
libcudacxxは、CUDA C ++標準ライブラリです。ホストコードとデバイスコードの両方で機能するC ++標準ライブラリの実装を提供します。さらに、同期プリミティブ、キャッシュコントロール、アトミックなどなど、CUDA固有のハードウェア機能の抽象化を提供します。
CCCLの主な目標は、標準のC ++ライブラリが標準のC ++に記入する同様の役割を埋めることです。CUDAC++開発者に汎用の速度の速度ツールを提供し、重要な問題の解決に集中できるようにします。これらのプロジェクトを統合することは、その目標を実現するための最初のステップです。
これは、スラスト、カブ、およびlibcudacxxからのCCCL機能の使用を示す簡単な例です。
Thrust/Cub/libcudacxxを使用して、単純な並列削減カーネルを実装する方法を示しています。各スレッドブロックは、 cub::BlockReduceを使用して配列のサブセットの合計を計算します。各ブロックの合計は、libcudacxxのcuda::atomic_refを介したアトミックアドを使用して単一の値に縮小されます。
次に、スラストのreduceアルゴリズムを使用して同じ削減を行う方法を示し、結果を比較します。
ゴッドボルトでライブで試してみてください!
# 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の使用を開始する最も簡単な方法は、CCCLヘッダーを含むCUDAツールキットを介してです。 nvccをコンパイルすると、CCCLヘッダーがインクルードパスに自動的に追加されるため、追加の構成が必要ない場合にCCCLヘッダーをコード内に#includeだけです。
別のコンパイラとコンパイルする場合、CTKインストールのCCCLヘッダーを指すための検索パスを含むビルドシステムを更新する必要があります(例えば、 /usr/local/cuda/include )。
# include < thrust/device_vector.h >
# include < cub/cub.cuh >
# include < cuda/std/atomic > CCCL開発の最先端にとどまりたいユーザーは、GitHubのCCCLを使用することをお勧めします。 CCCLの新しいバージョンのCCCLを使用して、古いバージョンのCUDAツールキットを使用することはサポートされていますが、その逆ではありません。 CCCLとCUDA Toolkitの互換性に関する完全な情報については、プラットフォームサポートを参照してください。
CCCLのすべてはヘッダーのみであるため、単純なプロジェクトにクローン化して含めることは、次のように簡単です。
git clone https://github.com/NVIDIA/cccl.git
nvcc -Icccl/thrust -Icccl/libcudacxx/include -Icccl/cub main.cu -o mainCCUDAツールキットの
nvccが暗黙的に含めるCCCLヘッダーとの衝突を回避するために、-I-isystemなく使用します。すべての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-forgeチャンネルを介して各リリースのコンドラパッケージも提供します。
conda config --add channels conda-forge
conda install ccclこれにより、最新のCCCLがConda環境の$CONDA_PREFIX/include/ and $CONDA_PREFIX/lib/cmake/ directoriesにインストールされます。 Cmakeはfind_package(CCCL)を介して発見可能であり、Conda環境の任意のコンパイラが使用できます。詳細については、この紹介の紹介を参照してください。
特定のCUDAツールキットを搭載した同じCCCLバージョンを使用する場合は、例えばCUDA 12.4を使用します。CCCLをインストールできます。
conda config --add channels conda-forge
conda install cuda-cccl cuda-version=12.4 cuda-ccclメタパッケージは、 cuda-versionに対応するCUDAツールキットに出荷されたccclバージョンをインストールします。 cccl更新する前に、 cuda-ccclをインストールした後に最新のccclを更新する場合は、 cuda-ccclをアンインストールします。
conda uninstall cuda-cccl
conda install -c conda-forge cccl注
cuda-cccl_linux-64などの名前のコンドラパッケージもあります。これらのパッケージには、CUDAツールキットの一部として出荷されたCCCLバージョンが含まれていますが、CUDAツールキットが内部使用するために設計されています。代わりに、ccclまたはcuda-ccclをインストールして、CONDAコンパイラと互換性があります。詳細については、CCCLコンドラフォージレシピを参照してください。
CCCLは、他のCMAKEプロジェクトでリンクするテストやターゲットなど、すべてのビルドおよびインストールインフラストラクチャにCMAKEを使用します。したがって、CMAKEは、CCCLを別のプロジェクトに統合するための推奨される方法です。
Cmakeパッケージマネージャーを使用してこれを行う方法の完全な例については、基本的な例プロジェクトを参照してください。
他のビルドシステムは機能するはずですが、Cmakeのみがテストされます。 CCCLの統合を他のビルドシステムに簡素化するための貢献は大歓迎です。
CCCLの改善に貢献することに興味がありますか?開発環境をセットアップし、変更を加え、テストを実行し、PRを送信するために知っておく必要があるすべての概念の概要については、寄稿ガイドをご覧ください。
目的:このセクションでは、ユーザーがCCCLがコンパイルして正常に実行することを期待できる場所について説明します。
一般に、CCCLはCUDAツールキットがサポートされている至る所であらゆる場所で動作するはずですが、悪魔は詳細にあります。以下のセクションでは、CUDAツールキット、ホストコンパイラ、C ++方言のさまざまなバージョンのサポートとテストの詳細について説明します。
まとめ:
CCCLユーザーは、常に最新バージョンのCCCLを使用することにより、最新の拡張機能を活用し、「Live at Head」を利用することをお勧めします。シームレスなエクスペリエンスをするには、CCCLをCCCL全体とは独立してアップグレードできます。これは、CCCLが現在および以前のメジャーバージョンシリーズの両方からのすべてのマイナーCTKリリースの最新のパッチリリースとの後方互換性を維持するためです。いくつかの例外的なケースでは、CUDAツールキットリリースの最小限のマイナーバージョンは、メジャーバージョンシリーズ内で最も古いリリースよりも新しいものである必要がある場合があります。たとえば、CCCLは、CTK 11.0に存在する避けられないコンパイラの問題により、11.xシリーズの11.1の最小サポートバージョンを必要とします。
新しいメジャーCTKがリリースされると、最も古いサポートされたメジャーバージョンのサポートをドロップします。
| CCCLバージョン | CUDAツールキットバージョンをサポートしています |
|---|---|
| 2.x | 11.1-11.8、12.x(最新のパッチのみリリース) |
| 3.x(将来) | 12.x、13.x(最新のパッチのみリリース) |
最新のCCCLを使用した行儀コードは、サポートされているCTKバージョンでコンパイルして正常に実行する必要があります。新しいCTK機能に依存する新機能の例外が発生する可能性があるため、これらの機能はCTKの古いバージョンでは機能しません。たとえば、C ++ 20のサポートはCUDA 12.0までnvccに追加されなかったため、C ++ 20に依存するCCCL機能は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は2つの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はCUDAツールキットと同じGPUアーキテクチャ/計算機能をすべてサポートしています。これはここに記録されています:https://docs.nvidia.com/cuda/cuda/cuda-c-programming-guide/index.html#compute-capability
一部の機能には、特定のアーキテクチャ/計算機能のみをサポートする場合があることに注意してください。
CCCLのテスト戦略は、できるだけ多くの構成をテストすることと、合理的なCI時間を維持することとのバランスをとっています。
CUDAツールキットバージョンの場合、テストは、最古のバージョンと最新バージョンの両方に対して行われます。たとえば、CUDAツールキットの最新バージョンが12.3の場合、テストは11.1と12.3に対して実施されます。 CUDAバージョンごとに、サポートされているすべてのC ++方言を持つすべてのサポートされているホストコンパイラーに対してビルドが完了します。
テスト戦略とマトリックスは常に進化しています。 ci/matrix.yamlファイルで定義されているマトリックスは、真実の決定的な源です。 CIパイプラインの詳細については、こちらをご覧ください。
目的:このセクションでは、CCCLのバージョン化、API/ABI安定性の保証、およびアップグレードヘッジを最小限に抑える互換性ガイドラインについて説明します。
まとめ
cub::またはthrust::名前空間のエンティティのABI安定性に依存しないでくださいcuda:: NamespaceのシンボルのABIブレイキング変更はいつでも発生する可能性がありますが、すべてのcuda::のインラインネームスペースに埋め込まれたABIバージョンを増加させることで反映されます。複数のABIバージョンが同時にサポートされる場合があります。注:このリポジトリにスラスト、カブ、およびlibcudacxxをマージする前に、各ライブラリはセマンティックバージョンに従って独立してバージョンされていました。 2.1リリースから始めて、3つのライブラリはすべて、リリースバージョンを別々のリポジトリで同期しました。前進すると、CCCLは1つのセマンティックバージョンで引き続きリリースされ、2.2.0はNVIDIA/CCCLリポジトリからの最初のリリースです。
壊れた変更とは、新しいバージョンにアップグレードするためにユーザーが作業を行う必要があるリリースされたバージョン間で明示的にサポートされている機能が明示的にサポートされる機能です。
限界では、変更はどこかに誰かを壊す可能性があります。その結果、可能な限りすべてのソースブレイク変更の変更が、主要なセマンティックバージョンに衝突する必要があるパブリックAPIの変更を破壊すると見なされるわけではありません。
以下のセクションでは、CCCLのAPIとABIの壊れた変更の詳細について説明します。
CCCLのパブリックAPIは、ライブラリのユーティリティを提供するために意図的に露出した機能全体です。
言い換えれば、CCCLのパブリックAPIは機能の署名だけを超えており、以下を含みます(ただしません)。
さらに、CCCLのパブリックAPIには以下のいずれも含まれていません。
_または__が付いたプレフィックスの記号detailを含むdetail::名前空間またはマクロを含むdetail/ディレクトリまたはサブディレクトリに含まれるヘッダーファイル一般に、目標は、公共のAPIで何かを壊すことを避けることです。このような変更は、ユーザーがより良いパフォーマンス、理解しやすいAPI、および/またはより一貫したAPIを提供する場合にのみ行われます。
パブリックAPIへの壊れた変更には、CCCLのメジャーバージョン番号がぶつかる必要があります。 CUDAマイナーバージョンの互換性に沿って、APIの壊れた変更とCCCLメジャーバージョンのバンプは、Cuda Toolkitの新しいメジャーバージョンリリースとのみ発生します。
パブリックAPIの一部ではないものは、警告なしにいつでも変更される場合があります。
すべてのCCCLのコンポーネントのパブリックAPIは、 MAJOR.MINOR.PATCHの統一されたセマンティックバージョンを共有しています。
最近リリースされたバージョンのみがサポートされています。原則として、機能とバグの修正は、以前にリリースされたバージョンまたはブランチにバックポートされていません。
バージョンをクエリするための優先方法は、以下に説明するようにCCCL_[MAJOR/MINOR/PATCH_]VERSIONを使用することです。後方互換性のために、スラスト/カブ/libcudacxxxバージョンの定義が利用可能であり、常にCCCL_VERSIONと一致します。 Thrust/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の壊れた変更は、パブリックAPIの関数またはタイプのABIの変更をもたらす変更です。たとえば、新しいデータメンバーを構造体に追加することは、タイプのサイズを変更するため、ABIの壊れた変化です。
CCCLでは、ABIに関する保証は次のとおりです。
thrust::およびcub:: namespacesは、警告なしにいつでもABIを壊す可能性があります。thrust:: cub::シンボルのABIには、コンパイルに使用されるCUDAアーキテクチャが含まれています。したがって、 thrust:: cub::-x cu )vs C ++ソース( -x cpp )としてコンパイルcuda:: namespaceのシンボルは、いつでもABIを壊す場合があります。ただし、 cuda:: Symbolsは、ABIブレイクが発生するたびに増加するABIバージョン番号を埋め込みました。複数のABIバージョンが同時にサポートされる場合があるため、ユーザーは以前のABIバージョンに戻すオプションがあります。詳細については、こちらをご覧ください。誰がABIを気にするべきですか?
一般に、CCCLユーザーは、APIがCCCLが提供するタイプを直接または間接的に含むバイナリアーティファクト(共有ライブラリなど)を構築または使用する際にABIの問題を心配する必要があります。
たとえば、 libA.soがCCCLバージョンXを使用して構築されたかどうかを検討し、そのパブリックAPIには次のような関数が含まれています。
void foo (cuda::std::optional< int >);別のライブラリlibB.soがCCCLバージョンYを使用してコンパイルされ、 libA.soのfooを使用する場合、バージョンXとYの間にABIブレークがあった場合、これは失敗する可能性があります。 APIの侵害の変更とは異なり、ABIブレイクは通常、コードの変更を必要とせず、同じABIバージョンを使用するためにすべてを再コンパイルする必要があります。
ABIの詳細とそれが重要な理由については、ABIとは何か、C ++はそれについて何をすべきかを参照してください。
前述のように、すべての可能なソース壊しの変更が、CCCLのAPIメジャーバージョン番号の増加を必要とする壊れた変化を構成するわけではありません。
ユーザーは、パブリックAPIの一部ではないCCCLの一部に応じて、誤って混乱のリスクを最小限に抑えるために、次のガイドラインを遵守することをお勧めします。
thrust:: 、 cub:: 、 nv:: 、またはcuda:: namepacesに宣言を追加したり、テンプレートを専門化したりしないでくださいcuda::std::iterator_traitsthrust:: 、 cub:: 、 cuda:: 、またはnv:: namespaces。thrust:: 、 cub:: 、 cuda:: 、またはnv:: namespaces。_ 、 __でプレフィックスしたシンボルを直接参照しないでください。または、その名前のdetail:: namespaceまたはmacroを含むdetailどこにでも参照しないでください#include 。言い換えれば、他のヘッダーに暗黙的に含まれるヘッダーに依存しないでください。このセクションの一部は、Abseilの互換性ガイドラインに触発されました。
パブリックAPI、ABI、またはサポートされているプラットフォームとコンパイラを変更する前に、ユーザーに通知するために最善を尽くします。
必要に応じて、非難は、無効になる可能性のあるプログラム的な警告の形でもたらされます。
非推奨期間は、変化の影響に依存しますが、通常、少なくとも2つのマイナーバージョンリリースは持続します。
近日公開!
CIパイプラインの詳細な概要については、ci-overview.mdを参照してください。
CCCLの任務に関連するプロジェクトは、CUDAをより楽しくすることです。
あなたのプロジェクトはCCCLを使用していますか? PRを開いて、このリストにプロジェクトを追加してください!