The NVIDIA CUDA Core Compute Libraries (CCCL) gives pleasant and environment friendly abstractions for CUDA builders in C++ and Python. It options:<\/p>\n
Parallel algorithms \u2013 Host-launched algorithms together with type, scan and scale back that take away the necessity to write customized kernels for widespread operations\u00a0<\/p>\n
Cooperative algorithms \u2013 System-side algorithms corresponding to block-wide or warp-wide reductions or scans that simplify customized kernel improvement\u00a0<\/p>\n
Language idiomatic CUDA abstractions \u2013 Basic abstractions for CUDA-specific operations together with reminiscence allocation, useful resource administration, and {hardware} options\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0<\/p>\n
This publish introduces a brand new group of performance in CCCL that gives modernized C++ abstractions for elementary CUDA programming mannequin ideas that make CUDA C++ improvement safer and extra handy.\u00a0\u00a0<\/p>\n
What’s CCCL runtime?<\/h2>\n
NVIDIA CCCL runtime is a brand new set of idiomatic C++ APIs that implement core CUDA performance: stream administration, reminiscence allocation, kernel launches, and extra.\u00a0<\/p>\n
The acquainted NVIDIA CUDA runtime was initially developed as a comfort layer on prime of the CUDA driver API. The brand new CCCL runtime goals to be another with the identical purpose, however with an up to date design aligned with fashionable C++. Determine 1, under, reveals the connection between the three CUDA API surfaces talked about above:<\/p>\n
![\"Both](\"https:\/\/developer-blogs.nvidia.com\/wp-content\/uploads\/2026\/06\/image1-9.webp\")
CCCL runtime is a set of headers inside CCCL, corresponding to , , and . It leverages fashionable C++ options to supply extra handy and sturdy abstractions than what was attainable inside the C supply compatibility constraints of the standard CUDA runtime API.\u00a0<\/p>\n
We additionally took the chance to include classes realized over 20 years of CUDA evolution into the API design. Even with all these modifications, CCCL runtime gives compatibility helpers that permit builders undertake it incrementally with out rewriting surrounding code that makes use of the CUDA runtime API.<\/p>\n
As CUDA applications develop extra complicated,\u00a0 with a number of libraries sharing units, streams, and reminiscence, the necessity for APIs that compose cleanly and make dependencies specific turns into extra urgent. That’s the house CCCL runtime is designed to fill.<\/p>\n
The code<\/h2>\n
Right here is the basic vectorAdd instance carried out with the brand new CCCL runtime APIs. If you happen to\u2019ve written CUDA earlier than, the general construction shall be acquainted: Deal with what\u2019s totally different. Don\u2019t attempt to perceive all the things directly, the remainder of this publish will stroll by means of this instance to elucidate the semantics and design decisions behind CCCL runtime.<\/p>\n
\n#embrace
\n#embrace
\n#embrace
\n#embrace
\n#embrace <\/p>\n
struct kernel {
\n template
\n __device__ void operator()(Config config,
\n cuda::std::span A,
\n cuda::std::span B,
\n cuda::std::span C) {
\n auto tid = cuda::gpu_thread.rank(cuda::grid, config);
\n if (tid < A.dimension())
\n C[tid] = A[tid] + B[tid];
\n }
\n}; <\/p>\n
int foremost() {
\n \/\/ 1. Gadgets and streams
\n cuda::device_ref machine = cuda::units[0];
\n cuda::stream stream{machine}; <\/p>\n
\/\/ 2. Reminiscence allocation
\n auto pool = cuda::device_default_memory_pool(machine); <\/p>\n
int num_elements = 1000;
\n auto A = cuda::make_buffer(stream, pool, num_elements, 1);
\n auto B = cuda::make_buffer(stream, pool, num_elements, 2);
\n auto C = cuda::make_buffer(stream, pool, num_elements, cuda::no_init); <\/p>\n
\/\/ 3. Kernel launch
\n constexpr int threads_per_block = 256;
\n auto config = cuda::distribute(num_elements); <\/p>\n
cuda::launch(stream, config, kernel{}, A, B, C); <\/p>\n
\/\/ Make the CPU thread look ahead to the GPU work to complete.
\n stream.sync();
\n return 0;
\n}\n<\/p><\/div>\n
The instance may be damaged down into the next three foremost sections:<\/p>\n
1.) Gadgets and streams<\/h2>\n
Contemplate the creation of a stream utilizing the CUDA Runtime API as the next code snippet reveals.<\/p>\n
\ncudaStreamCreate(&stream); \/\/ related to whichever machine occurs to be “present”\n<\/div>\n
Notice this creates a stream, however the stream is related to whichever machine is present when cudaStreamCreate known as.\u00a0 Primarily based on this name alone, you don\u2019t know which machine the stream is related to.<\/p>\n
Distinction that with utilizing CCCL runtime API as illustrated by the code snippet that follows.<\/p>\n
\ncuda::stream stream{machine};\n<\/div>\n
The above code snippet reveals the way to create a stream on a selected machine. The primary line illustrates a core design precept: CCCL runtime makes use of devoted sorts as a substitute of uncooked identifiers. A tool is a device_ref, not a plain integer; a stream is an object, not an opaque pointer. Robust typing throughout the API helps catch errors at compile time slightly than chasing them at runtime.\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0<\/p>\n
The second line illustrates one other precept: making dependencies specific. In each CCCL runtime and the CUDA runtime API, a stream is related to a tool. The distinction is how. Right here, the cuda::stream constructor takes the machine as an specific argument whereas with the CUDA runtime API the stream is related to whichever machine is lively when the stream is created.<\/p>\n
Express dependencies allow native reasoning. You’ll be able to learn a perform and perceive what it does with out monitoring the worldwide state. Additionally they enhance composability: When a number of libraries are used, none of them want to avoid wasting and restore implicit state throughout calls to keep away from interfering with one another.\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0<\/p>\n
A associated consequence is that CCCL runtime doesn\u2019t expose the default stream. Managing the which means of the default stream requires monitoring the present machine, which is precisely the type of implicit state we’re shifting away from. Whereas a default stream from the CUDA runtime API can nonetheless be wrapped into CCCL runtime sorts, its utilization is discouraged; something involving the default stream must be dealt with by means of the CUDA runtime API instantly. With no default stream within the API, the notion of a \u201cblocking stream\u201d not applies, so all CCCL runtime streams are created as non-blocking.<\/p>\n
Useful resource possession: Proudly owning sorts and refs<\/h3>\n
Following the instance of std::string and std::string_view, many CUDA objects have two sorts in CCCL runtime: an proudly owning sort and a non-owning sort with a _ref suffix; cuda::stream owns the underlying cudaStream_t deal with and destroys it in its destructor. The cuda::stream_ref holds the deal with with out managing its lifetime and is trivially copyable.\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0<\/p>\n
The _ref sorts are important for composability with current code. If a stream deal with\u2019s lifetime is managed elsewhere, cudaStream_t implicitly converts to cuda::stream_ref, and the uncooked deal with may be retrieved with .get(). To switch possession, cuda::stream::from_native_handle wraps a uncooked deal with into the proudly owning sort, and .launch() relinquishes possession again.\u00a0\u00a0<\/p>\n
\n cuda::stream_ref non_owning{deal with};
\n assert(deal with == non_owning.get());<\/p>\n
cuda::stream proudly owning = cuda::stream::from_native_handle(deal with);
\n assert(deal with == proudly owning.get());
\n assert(deal with == proudly owning.launch());
\n}\n<\/p><\/div>\n
The identical sample applies to occasions, reminiscence swimming pools, and different CUDA objects: cuda::device_ref has no proudly owning counterpart as a result of there is no such thing as a machine state to personal.<\/p>\n
2.) Reminiscence allocation <\/h2>\n\nauto pool = cuda::device_default_memory_pool(machine);<\/p>\nauto A = cuda::make_buffer(stream, pool, num_elements, 1);
\nauto B = cuda::make_buffer(stream, pool, num_elements, 2);
\nauto C = cuda::make_buffer(stream, pool, num_elements, cuda::no_init);\n<\/p><\/div>\n
The following part demonstrates asynchronously allocating and initializing machine reminiscence. Right here we see the following design precept: APIs are asynchronous by default. Slightly than distinguishing synchronous and asynchronous variants by title, CCCL runtime makes use of a easy conference: If an API takes a stream as its first argument, it operates in stream order. We don\u2019t plan to supply synchronous counterparts for APIs which have each variants within the CUDA runtime API. \u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0<\/p>\n
Reminiscence allocation is the place this issues most in apply. Stream-ordered reminiscence administration through reminiscence swimming pools has been obtainable since CUDA 11.2 (defined right here), and CUDA 13.0 expanded it to managed and host reminiscence. Reminiscence pooling and fewer frequent synchronization factors are generally important to succeed in most efficiency, and stream-ordered reminiscence administration composes naturally with the remainder of the asynchronous programming mannequin. To convey these tips, CCCL runtime makes reminiscence swimming pools and stream-ordered allocation the default. On older CUDA variations and platforms, the place newer reminiscence pool sorts usually are not but supported, we offer non-stream-ordered allocation as a fallback, however plan to take away it as soon as pool assist is common.\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0<\/p>\n
Within the snippet above, we first question the default reminiscence pool for a given machine, passing it as an specific argument slightly than counting on cudaMallocAsync\u2018s implicit machine choice. The instance makes use of the default pool which must be most well-liked the place attainable, however CCCL runtime additionally permits creating separate pool objects when totally different pool settings are wanted.<\/p>\n
The pool reference is then used to create three buffers utilizing the brand new cuda::make_buffer. It takes a stream as its first argument to sign stream-ordered operation. Every buffer submits three operations to that stream: allocation from the required pool, initialization, and ultimately deallocation when the buffer goes out of scope.<\/p>\n
Initialization is necessary until explicitly opted out with cuda::no_init, as with buffer C which shall be overwritten by the kernel. Uninitialized machine reminiscence is a typical supply of hard-to-diagnose bugs, so we selected to require an specific opt-out slightly than making it the silent default. Enter buffers A and B have all components initialized to 1 and a couple of, respectively. Buffers assist further initialization modes as effectively, for instance from one other buffer or a spread.<\/p>\n
Buffer lifetime and deallocation<\/h3>\n
The stream handed to make_buffer is saved contained in the buffer and used for deallocation when the buffer is destroyed. This implies the buffer ought to typically maintain the stream that corresponds to its utilization, in order that computation is correctly ordered with deallocation. It’s attainable to alter the stream later with .set_stream() or manually set off destruction on a selected stream with .destroy(), however the default habits is designed to do the appropriate factor within the widespread case.<\/p>\n
\n{
\n auto pool = cuda::device_default_memory_pool(machine);
\n \/\/ Equal to cudaMallocFromPoolAsync on the stream, probably together with initialization pushed into the stream as effectively. Saves the stream for future deallocation
\n auto buffer = cuda::make_buffer(allocation_stream, pool, \/*… *\/);<\/p>\n \/\/ buffer utilization…
\n}
\n\/\/ Closing bracket will name cudaFreeAsync on allocation_stream, there may be additionally buffer.destroy(which_stream) to maintain the habits specific\n<\/p><\/div>\n
3.) Kernel launch<\/h2>\n\nstruct kernel {
\n template
\n __device__ void operator()(Config config,
\n cuda::std::span A,
\n cuda::std::span B,
\n cuda::std::span C) {
\n auto tid = cuda::gpu_thread.rank(cuda::grid, config);
\n if (tid < A.dimension())
\n C[tid] = A[tid] + B[tid];
\n }
\n};<\/p>\n\/\/ …<\/p>\n
constexpr int threads_per_block = 256;
\nauto config = cuda::distribute(num_elements);<\/p>\n
cuda::launch(stream, config, kernel{}, A, B, C);\n<\/p><\/div>\n
The ultimate part demonstrates configuring and launching the kernel on the GPU with cuda::launch.\u00a0\u00a0\u00a0\u00a0\u00a0<\/p>\n
cuda::launch takes three teams of arguments:<\/p>\n
The stream to run on<\/p>\n
A configuration object that encodes the thread hierarchy (block and grid sizes) together with different launch choices. Right here, cuda::distribute creates a configuration that launches at the least num_elements threads grouped into blocks of threads_per_block. This replaces the widespread sample many CUDA builders are accustomed to of (N + block_size – 1) \/ block_size<\/p>\n
The\u00a0 kernel and its arguments<\/p>\n
Compile-time configuration stream<\/h2>\n
Probably the most novel facet of cuda::launch is the way it strikes compile-time data from the host launch web site into machine code by means of the kind system. For instance, discover how the\u00a0block dimension is offered as a template argument to cuda::distribute, which suggests it’s encoded within the configuration object\u2019s sort. <\/p>\n
When the kernel accepts that configuration as its first argument, cuda::launch passes it by means of mechanically. Contained in the kernel, this static data is accessible once we compute the rank of the calling thread contained in the grid:<\/p>\n
\nauto tid = cuda::gpu_thread.rank(cuda::grid, config);\n<\/div>\nAs a result of the block dimension is thought at compile time, the rank calculation can use solely the x dimension and skip the runtime block-size question solely. It is a easy instance, however the mechanism generalizes.\u00a0 The CCCL documentation reveals additional circumstances the place configuration-embedded data is used to specialize machine code.<\/p>\n
Typically kernel implementation makes assumptions concerning the precise form of the grid and\/or block. Compile time data within the configuration object permits kernel authors to implement checks to make sure alignment of the kernel and the decision web site in these circumstances.<\/p>\n
\ntemplate
\n__global__ void kernel(Config conf) {
\n \/\/ Be sure that the block is one dimensional with 256 threads
\n static_assert(cuda::gpu_thread.static_dims(cuda::block, conf).x == 256);
\n static_assert(cuda::gpu_thread.static_dims(cuda::block, conf).y == 1);
\n static_assert(cuda::gpu_thread.static_dims(cuda::block, conf).z == 1);
\n}\n<\/div>\nKernel functors<\/h2>\n
You might have observed the kernel is a struct with a __device__ operator() slightly than a __global__ perform. Whereas cuda::launch helps current __global__ capabilities, we additionally launched kernel functors: sorts with a __device__-annotated name operator. The sensible benefit is that template arguments are deduced mechanically, whereas __global__ capabilities used with cuda::launch require specific instantiation.<\/p>\n
auto A = cuda::make_buffer(stream, pool, num_elements, 1);
\nauto B = cuda::make_buffer(stream, pool, num_elements, 2);
\nauto C = cuda::make_buffer(stream, pool, num_elements, cuda::no_init);\n<\/p><\/div>\n
The following part demonstrates asynchronously allocating and initializing machine reminiscence. Right here we see the following design precept: APIs are asynchronous by default. Slightly than distinguishing synchronous and asynchronous variants by title, CCCL runtime makes use of a easy conference: If an API takes a stream as its first argument, it operates in stream order. We don\u2019t plan to supply synchronous counterparts for APIs which have each variants within the CUDA runtime API. \u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0<\/p>\n
Reminiscence allocation is the place this issues most in apply. Stream-ordered reminiscence administration through reminiscence swimming pools has been obtainable since CUDA 11.2 (defined right here), and CUDA 13.0 expanded it to managed and host reminiscence. Reminiscence pooling and fewer frequent synchronization factors are generally important to succeed in most efficiency, and stream-ordered reminiscence administration composes naturally with the remainder of the asynchronous programming mannequin. To convey these tips, CCCL runtime makes reminiscence swimming pools and stream-ordered allocation the default. On older CUDA variations and platforms, the place newer reminiscence pool sorts usually are not but supported, we offer non-stream-ordered allocation as a fallback, however plan to take away it as soon as pool assist is common.\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0<\/p>\n
Within the snippet above, we first question the default reminiscence pool for a given machine, passing it as an specific argument slightly than counting on cudaMallocAsync\u2018s implicit machine choice. The instance makes use of the default pool which must be most well-liked the place attainable, however CCCL runtime additionally permits creating separate pool objects when totally different pool settings are wanted.<\/p>\n
The pool reference is then used to create three buffers utilizing the brand new cuda::make_buffer. It takes a stream as its first argument to sign stream-ordered operation. Every buffer submits three operations to that stream: allocation from the required pool, initialization, and ultimately deallocation when the buffer goes out of scope.<\/p>\n
Initialization is necessary until explicitly opted out with cuda::no_init, as with buffer C which shall be overwritten by the kernel. Uninitialized machine reminiscence is a typical supply of hard-to-diagnose bugs, so we selected to require an specific opt-out slightly than making it the silent default. Enter buffers A and B have all components initialized to 1 and a couple of, respectively. Buffers assist further initialization modes as effectively, for instance from one other buffer or a spread.<\/p>\n
Buffer lifetime and deallocation<\/h3>\n
The stream handed to make_buffer is saved contained in the buffer and used for deallocation when the buffer is destroyed. This implies the buffer ought to typically maintain the stream that corresponds to its utilization, in order that computation is correctly ordered with deallocation. It’s attainable to alter the stream later with .set_stream() or manually set off destruction on a selected stream with .destroy(), however the default habits is designed to do the appropriate factor within the widespread case.<\/p>\n
\n auto pool = cuda::device_default_memory_pool(machine);
\n \/\/ Equal to cudaMallocFromPoolAsync on the stream, probably together with initialization pushed into the stream as effectively. Saves the stream for future deallocation
\n auto buffer = cuda::make_buffer(allocation_stream, pool, \/*… *\/);<\/p>\n
\/\/ buffer utilization…
\n}
\n\/\/ Closing bracket will name cudaFreeAsync on allocation_stream, there may be additionally buffer.destroy(which_stream) to maintain the habits specific\n<\/p><\/div>\n
3.) Kernel launch<\/h2>\n\nstruct kernel {
\n template
\n __device__ void operator()(Config config,
\n cuda::std::span A,
\n cuda::std::span B,
\n cuda::std::span C) {
\n auto tid = cuda::gpu_thread.rank(cuda::grid, config);
\n if (tid < A.dimension())
\n C[tid] = A[tid] + B[tid];
\n }
\n};<\/p>\n\/\/ …<\/p>\n
constexpr int threads_per_block = 256;
\nauto config = cuda::distribute(num_elements);<\/p>\n
cuda::launch(stream, config, kernel{}, A, B, C);\n<\/p><\/div>\n
The ultimate part demonstrates configuring and launching the kernel on the GPU with cuda::launch.\u00a0\u00a0\u00a0\u00a0\u00a0<\/p>\n
cuda::launch takes three teams of arguments:<\/p>\n
The stream to run on<\/p>\n
A configuration object that encodes the thread hierarchy (block and grid sizes) together with different launch choices. Right here, cuda::distribute creates a configuration that launches at the least num_elements threads grouped into blocks of threads_per_block. This replaces the widespread sample many CUDA builders are accustomed to of (N + block_size – 1) \/ block_size<\/p>\n
The\u00a0 kernel and its arguments<\/p>\n
Compile-time configuration stream<\/h2>\n
Probably the most novel facet of cuda::launch is the way it strikes compile-time data from the host launch web site into machine code by means of the kind system. For instance, discover how the\u00a0block dimension is offered as a template argument to cuda::distribute, which suggests it’s encoded within the configuration object\u2019s sort. <\/p>\n
When the kernel accepts that configuration as its first argument, cuda::launch passes it by means of mechanically. Contained in the kernel, this static data is accessible once we compute the rank of the calling thread contained in the grid:<\/p>\n
\nauto tid = cuda::gpu_thread.rank(cuda::grid, config);\n<\/div>\nAs a result of the block dimension is thought at compile time, the rank calculation can use solely the x dimension and skip the runtime block-size question solely. It is a easy instance, however the mechanism generalizes.\u00a0 The CCCL documentation reveals additional circumstances the place configuration-embedded data is used to specialize machine code.<\/p>\n
Typically kernel implementation makes assumptions concerning the precise form of the grid and\/or block. Compile time data within the configuration object permits kernel authors to implement checks to make sure alignment of the kernel and the decision web site in these circumstances.<\/p>\n
\ntemplate
\n__global__ void kernel(Config conf) {
\n \/\/ Be sure that the block is one dimensional with 256 threads
\n static_assert(cuda::gpu_thread.static_dims(cuda::block, conf).x == 256);
\n static_assert(cuda::gpu_thread.static_dims(cuda::block, conf).y == 1);
\n static_assert(cuda::gpu_thread.static_dims(cuda::block, conf).z == 1);
\n}\n<\/div>\nKernel functors<\/h2>\n
You might have observed the kernel is a struct with a __device__ operator() slightly than a __global__ perform. Whereas cuda::launch helps current __global__ capabilities, we additionally launched kernel functors: sorts with a __device__-annotated name operator. The sensible benefit is that template arguments are deduced mechanically, whereas __global__ capabilities used with cuda::launch require specific instantiation.<\/p>\n
\n template
\n __device__ void operator()(Config config,
\n cuda::std::span A,
\n cuda::std::span B,
\n cuda::std::span C) {
\n auto tid = cuda::gpu_thread.rank(cuda::grid, config);
\n if (tid < A.dimension())
\n C[tid] = A[tid] + B[tid];
\n }
\n};<\/p>\n
\/\/ …<\/p>\n
constexpr int threads_per_block = 256;
\nauto config = cuda::distribute(num_elements);<\/p>\n
cuda::launch(stream, config, kernel{}, A, B, C);\n<\/p><\/div>\n
The ultimate part demonstrates configuring and launching the kernel on the GPU with cuda::launch.\u00a0\u00a0\u00a0\u00a0\u00a0<\/p>\n
cuda::launch takes three teams of arguments:<\/p>\n
The stream to run on<\/p>\n
A configuration object that encodes the thread hierarchy (block and grid sizes) together with different launch choices. Right here, cuda::distribute creates a configuration that launches at the least num_elements threads grouped into blocks of threads_per_block. This replaces the widespread sample many CUDA builders are accustomed to of (N + block_size – 1) \/ block_size<\/p>\n
The\u00a0 kernel and its arguments<\/p>\n
Compile-time configuration stream<\/h2>\n
Probably the most novel facet of cuda::launch is the way it strikes compile-time data from the host launch web site into machine code by means of the kind system. For instance, discover how the\u00a0block dimension is offered as a template argument to cuda::distribute, which suggests it’s encoded within the configuration object\u2019s sort. <\/p>\n
When the kernel accepts that configuration as its first argument, cuda::launch passes it by means of mechanically. Contained in the kernel, this static data is accessible once we compute the rank of the calling thread contained in the grid:<\/p>\n
As a result of the block dimension is thought at compile time, the rank calculation can use solely the x dimension and skip the runtime block-size question solely. It is a easy instance, however the mechanism generalizes.\u00a0 The CCCL documentation reveals additional circumstances the place configuration-embedded data is used to specialize machine code.<\/p>\n
Typically kernel implementation makes assumptions concerning the precise form of the grid and\/or block. Compile time data within the configuration object permits kernel authors to implement checks to make sure alignment of the kernel and the decision web site in these circumstances.<\/p>\n
\n__global__ void kernel(Config conf) {
\n \/\/ Be sure that the block is one dimensional with 256 threads
\n static_assert(cuda::gpu_thread.static_dims(cuda::block, conf).x == 256);
\n static_assert(cuda::gpu_thread.static_dims(cuda::block, conf).y == 1);
\n static_assert(cuda::gpu_thread.static_dims(cuda::block, conf).z == 1);
\n}\n<\/div>\n
Kernel functors<\/h2>\n
You might have observed the kernel is a struct with a __device__ operator() slightly than a __global__ perform. Whereas cuda::launch helps current __global__ capabilities, we additionally launched kernel functors: sorts with a __device__-annotated name operator. The sensible benefit is that template arguments are deduced mechanically, whereas __global__ capabilities used with cuda::launch require specific instantiation.<\/p>\n