CUDA Programming Guide

https://docs.nvidia.com/cuda/cuda-programming-guide
Admin
CUDA is a parallel computing platform and programming model that enables developers to harness GPU...
Tokens:592,602
Snippets:2,948
Trust Score:10
Update:2 days ago
Show doc for...
Tokens:
Raw
### Example using cudaMemAdvise

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/unified-memory.html

Shows a code fragment that uses cudaMemAdvise to hint data usage. The example includes initialization, prefetching to the GPU, kernel launch, and prefetching back to the CPU.

```cpp
  init_data(data, dataSizeBytes);
  cudaMemLocation location = {.type = cudaMemLocationTypeDevice, .id = myGpuId};

  // encourage data to move to GPU before use
  const unsigned int flags = 0;
  cudaMemPrefetchAsync(data, dataSizeBytes, location, flags, s);

  // use data on GPU
  const uinsigned num_blocks = (dataSizeBytes + threadsPerBlock - 1) / threadsPerBlock;
  mykernel<<<num_blocks, threadsPerBlock, 0, s>>>(data, dataSizeBytes);

  // encourage data to move back to CPU
```

--------------------------------

### CUDA_DEVICE_ORDER examples

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/environment-variables.html

Examples of setting CUDA_DEVICE_ORDER to control device enumeration order.

```bash
CUDA_DEVICE_ORDER=FASTEST_FIRST
CUDA_DEVICE_ORDER=PCI_BUS_ID
nvidia-smi --query-gpu=name,pci.bus_id # Get list of PCI bus IDs
```

--------------------------------

### CUDA_VISIBLE_DEVICES examples

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/environment-variables.html

Examples of setting CUDA_VISIBLE_DEVICES to control GPU visibility and enumeration order.

```bash
nvidia-smi -L # Get list of GPU UUIDs
CUDA_VISIBLE_DEVICES=0,1
CUDA_VISIBLE_DEVICES=GPU-8932f937-d72c-4106-c12f-20bd9faed9f6
CUDA_VISIBLE_DEVICES=MIG-GPU-8932f937-d72c-4106-c12f-20bd9faed9f6/1/2
```

--------------------------------

### Reduce Operators Usage Example

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/device-callable-apis.html

Example showing how to use cg::reduce with various function objects, including hardware intrinsic dispatch on CC 8.0+ and fallback behavior for unsupported types.

```cpp
{
 // cg::plus<int> is specialized within cg::reduce and calls __reduce_add_sync(...) on CC 8.0+
 cg::reduce(tile, (int)val, cg::plus<int>());

 // cg::plus<float> fails to match with an accelerator and instead performs a standard shuffle based reduction
 cg::reduce(tile, (float)val, cg::plus<float>());

 // While individual components of a vector are supported, reduce will not use hardware intrinsics for the following
 // It will also be necessary to define a corresponding operator for vector and any custom types that may be used
 int4 vec = {...};
 cg::reduce(tile, vec, cg::plus<int4>())

 // Finally lambdas and other function objects cannot be inspected for dispatch
 // and will instead perform shuffle based reductions using the provided function object.
 cg::reduce(tile, (int)val, [](int l, int r) -> int {return l + r;});
}
```

--------------------------------

### Driver API host code for vector addition

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/03-advanced/driver-api.html

This example shows the host code for a vector addition kernel using the CUDA driver API. It initializes the API, gets a device, creates a context, loads a module from a PTX file, allocates device memory, copies data, and gets a function handle.

```cpp
int main()
{
 int N = ...;
 size_t size = N * sizeof(float);

 // Allocate input vectors h_A and h_B in host memory
 float* h_A = (float*)malloc(size);
 float* h_B = (float*)malloc(size);

 // Initialize input vectors
 ...

 // Initialize
 cuInit(0);

 // Get number of devices supporting CUDA
 int deviceCount = 0;
 cuDeviceGetCount(&deviceCount);
 if (deviceCount == 0) {
 printf("There is no device supporting CUDA.\n");
 exit (0);
 }

 // Get handle for device 0
 CUdevice cuDevice;
 cuDeviceGet(&cuDevice, 0);

 // Create context
 CUcontext cuContext;
 cuCtxCreate(&cuContext, 0, cuDevice);

 // Create module from binary file
 CUmodule cuModule;
 cuModuleLoad(&cuModule, "VecAdd.ptx");

 // Allocate vectors in device memory
 CUdeviceptr d_A;
 cuMemAlloc(&d_A, size);
 CUdeviceptr d_B;
 cuMemAlloc(&d_B, size);
 CUdeviceptr d_C;
 cuMemAlloc(&d_C, size);

 // Copy vectors from host memory to device memory
 cuMemcpyHtoD(d_A, h_A, size);
 cuMemcpyHtoD(d_B, h_B, size);

 // Get function handle from module
 CUfunction vecAdd;
 cuModuleGetFunction(&vecAdd, cuModule, "VecAdd");
```

--------------------------------

### nvcc with include, library, and macro flags

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/nvcc.html

Shows how to specify include directories, library paths, link against libraries, and define macros using nvcc.

```bash
nvcc example.cu -I path_to_include/ -L path_to_library/ -lcublas -o <output_file>
```

--------------------------------

### Broadcasting example in C++ (ct::iota, ct::tile)

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/writing-tile-kernels.html

Shows rank promotion and singleton stretching: an 8x2 tile is promoted to 1x8x2 and broadcast with a 4x1x2 tile to produce a 4x8x2 tile.

```C++
auto x = ct::iota<ct::tile<int, ct::shape<8, 2>>>();      // 8x2   (rank 2)
auto y = ct::iota<ct::tile<int, ct::shape<4, 1, 2>>>();   // 4x1x2 (rank 3)
auto z = x + y;                                           // x promoted to 1x8x2, then broadcasts to 4x8x2

```

--------------------------------

### Load Module and Get Kernel Handle

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/03-advanced/driver-api.html

Loads a CUDA module from a PTX file and retrieves a handle to a kernel function. This is a typical setup for using the driver API to execute device code.

```c
CUmodule cuModule;
cuModuleLoad(&cuModule, "myModule.ptx");
CUfunction myKernel;
cuModuleGetFunction(&myKernel, cuModule, "MyKernel");
```

--------------------------------

### Launching kernel with parameter buffer

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/03-advanced/driver-api.html

Shows how to set up the extra array with CU_LAUNCH_PARAM_BUFFER_POINTER and CU_LAUNCH_PARAM_BUFFER_SIZE, then call cuLaunchKernel with the parameter buffer.

```c
void* extra[] = {
 CU_LAUNCH_PARAM_BUFFER_POINTER, paramBuffer,
 CU_LAUNCH_PARAM_BUFFER_SIZE, &paramBufferSize,
 CU_LAUNCH_PARAM_END
};
cuLaunchKernel(cuFunction,
 blockWidth, blockHeight, blockDepth,
 gridWidth, gridHeight, gridDepth,
 0, 0, 0, extra);
```

--------------------------------

### Incorrect dynamic shared memory partitioning due to misalignment

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/writing-cuda-kernels.html

Shows an incorrect partitioning example where array1 is not aligned to 4 bytes because it starts at index 127 of the short array, causing misalignment for the float pointer.

```CUDA C
extern  __shared__  float array[];
short* array0 = (short*)array;
float* array1 = (float*)&array0[127];
```

--------------------------------

### Get available workqueue config resources using cudaDeviceGetDevResource

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/green-contexts.html

Shows how to populate a cudaDevResource struct with workqueue config resources for a GPU device. After a successful cudaDeviceGetDevResource call, the wqConcurrencyLimit and sharingScope can be reviewed. When starting from a device, wqConcurrencyLimit matches CUDA_DEVICE_MAX_CONNECTIONS or its default.

```cpp
int current_device = 0; // assume device ordinal of 0
CUDA_CHECK(cudaSetDevice(current_device));

cudaDevResource initial_WQ_config_resources = {};
CUDA_CHECK(cudaDeviceGetDevResource(current_device /* GPU device */,
 &initial_WQ_config_resources /* device resource to populate */,
 cudaDevResourceTypeWorkqueueConfig /* resource type*/));

std::cout << "Initial WQ config. resources: " << std::endl;
std::cout << " - WQ concurrency limit: " << initial_WQ_config_resources.wqConfig.wqConcurrencyLimit << std::endl;
std::cout << " - WQ sharing scope: " << initial_WQ_config_resources.wqConfig.sharingScope << std::endl;
```

--------------------------------

### Execution.Model.API.3: Single cudaStreamQuery call

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cuda-cpp-execution-model.html

This example is similar to API.2 but adds a single cudaStreamQuery call before the spin-loop. The allowed outcome is still that eventually no thread makes progress because a single query API call does not guarantee the device thread starts; only repeated calls do.

```cpp
// Example: Execution.Model.API.3
// Allowed outcome: eventually, no thread makes progress.
// Rationale: same as Example.Model.API.2, with the addition that a single CUDA query API call does not guarantee
// the device thread eventually starts, only repeated CUDA query API calls do (see Execution.Model.API.4).
cuda::atomic<int, cuda::thread_scope_system> flag = 0;
__global__ void producer() { flag.store(1); }
int main() {
    cudaHostRegister(&flag, sizeof(flag));
    producer<<<1,1>>>();
    (void)cudaStreamQuery(0);
    while (flag.load() == 0);
    return cudaDeviceSynchronize();
}
```

--------------------------------

### Broadcasting example in Python (ct.full)

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/writing-tile-kernels.html

Shows rank promotion and singleton stretching: an 8x2 tile is promoted to 1x8x2 and broadcast with a 4x1x2 tile to produce a 4x8x2 tile.

```Python
x = ct.full((8, 2),    3, dtype=ct.int32)   # 8x2   (rank 2)
y = ct.full((4, 1, 2), 5, dtype=ct.int32)   # 4x1x2 (rank 3)
z = x + y                                    # x promoted to 1x8x2, then broadcasts to 4x8x2

```

--------------------------------

### Querying block position in Python with ct.bid() and ct.num_blocks()

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/writing-tile-kernels.html

Python example showing how to get the block index and total number of blocks along an axis using ct.bid(axis) and ct.num_blocks(axis). The axis argument is 0, 1, or 2, and the returned values are int32 scalars.

```Python
@ct.kernel
def my_kernel(a, b, c):
    bid_x = ct.bid(0)          # block index along axis 0
    bid_y = ct.bid(1)          # block index along axis 1
    num_x = ct.num_blocks(0)   # total blocks along axis 0

```

--------------------------------

### Tail Launch Example: launchTailGraph and graphSetup

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/cuda-graphs.html

Shows how to launch a device graph from within a kernel using cudaStreamGraphTailLaunch. The kernel launchTailGraph launches the given graph, and graphSetup creates, instantiates, uploads, and launches the graphs. Tail launches execute one at a time in enqueue order.

```CUDA C++
__global__ void launchTailGraph(cudaGraphExec_t graph) {
    cudaGraphLaunch(graph, cudaStreamGraphTailLaunch);
}

void graphSetup() {
    cudaGraphExec_t gExec1, gExec2;
    cudaGraph_t g1, g2;

    // Create, instantiate, and upload the device graph.
    create_graph(&g2);
    cudaGraphInstantiate(&gExec2, g2, cudaGraphInstantiateFlagDeviceLaunch);
    cudaGraphUpload(gExec2, stream);

    // Create and instantiate the launching graph.
    cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);
    launchTailGraph<<<1, 1, 0, stream>>>(gExec2);
    cudaStreamEndCapture(stream, &g1);
    cudaGraphInstantiate(&gExec1, g1);

    // Launch the host graph, which will in turn launch the device graph.
    cudaGraphLaunch(gExec1, stream);
}
```

--------------------------------

### Querying block position in C++ with ct::bid() and ct::num_blocks()

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/writing-tile-kernels.html

C++ example showing how to get the block index and total number of blocks in each dimension using ct::bid() and ct::num_blocks(). The block index is used to determine which portion of data the block processes.

```C++
#include "cuda_tile.h"

__tile_global__ void my_kernel(float* a, float* b, float* c) {
    namespace ct = cuda::tiles;
    int bid_x = ct::bid().x;          // block index along .x
    int bid_y = ct::bid().y;          // block index along .y
    int num_x = ct::num_blocks().x;   // total blocks along .x
}

```

--------------------------------

### Build and Run Commands for Unified and Explicit Memory Examples

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/intro-to-cuda-cpp.html

Commands to compile and run the unified memory and explicit memory vector addition examples with nvcc. The output shows the expected success messages for default and custom vector lengths.

```bash
$ nvcc vecAdd_unifiedMemory.cu -o vecAdd_unifiedMemory
$./vecAdd_unifiedMemory
Unified Memory: CPU and GPU answers match
$./vecAdd_unifiedMemory 4096
Unified Memory: CPU and GPU answers match
$ nvcc vecAdd_explicitMemory.cu -o vecAdd_explicitMemory
$./vecAdd_explicitMemory
Explicit Memory: CPU and GPU answers match
$./vecAdd_explicitMemory 4096
Explicit Memory: CPU and GPU answers match
```

--------------------------------

### Launch Kernel with cuLaunchKernel

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/03-advanced/driver-api.html

Shows how to launch a kernel using the driver API. It sets up grid and block dimensions, prepares arguments, and calls cuLaunchKernel. The full code is available in the vectorAddDrv CUDA sample.

```c
 // Invoke kernel
 int threadsPerBlock = 256;
 int blocksPerGrid =
 (N + threadsPerBlock - 1) / threadsPerBlock;
 void* args[] = { &d_A, &d_B, &d_C, &N };
 cuLaunchKernel(vecAdd,
 blocksPerGrid, 1, 1, threadsPerBlock, 1, 1,
 0, 0, args, 0);
```

--------------------------------

### Execution.Model.API.1: cudaDeviceSynchronize with __syncthreads

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cuda-cpp-execution-model.html

This example shows a simple kernel with a __syncthreads barrier and a cudaDeviceSynchronize call. The program terminates and returns cudaSuccess because CUDA guarantees that if the device is empty, cudaDeviceSynchronize eventually ensures a device thread makes progress, leading to all threads starting and reaching the barrier.

```cpp
// Example: Execution.Model.API.1
// Outcome: if no other device threads (e.g., from other processes) are making progress,
// this program terminates and returns cudaSuccess.
// Rationale: CUDA guarantees that if the device is empty:
// - `cudaDeviceSynchronize` eventually ensures that at least one device-thread makes progress, which implies that eventually `hello_world` grid and one of its device-threads start.
// - All thread-block threads eventually start (due to "if a device thread makes progress, all other threads in its thread-block cluster eventually make progress").
// - Once all threads in thread-block arrive at `__syncthreads` barrier, all waiting threads are unblocked.
// - Therefore all device threads eventually exit the `hello_world`` grid.
// - And `cudaDeviceSynchronize`` eventually unblocks.
__global__ void hello_world() { __syncthreads(); }
int main() {
    hello_world<<<1,2>>>();
    return (int)cudaDeviceSynchronize();
}
```

--------------------------------

### Non-blocking event check with cudaEventQuery

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/asynchronous-execution.html

Use cudaEventQuery to non-blockingly check if an event has completed, allowing CPU work to overlap with GPU execution. This example launches two kernels, records an event, and polls the event while doing CPU work, then starts a device-to-host copy when the event completes.

```cpp
cudaEvent_t event;
cudaStream_t stream1;
cudaStream_t stream2;

size_t size = LARGE_NUMBER;
float *d_data;

// Create some data
cudaMalloc(&d_data, size);
float *h_data = (float *)malloc(size);

// create the streams
cudaStreamCreate(&stream1);   // Processing stream
cudaStreamCreate(&stream2);   // Copying stream
bool copyStarted = false;

//  create the event
cudaEventCreate(&event);

// launch kernel1 into the stream
kernel1<<<grid, block, 0, stream1>>>(d_data, size);
// enqueue an event following kernel1
cudaEventRecord(event, stream1);

// launch kernel2 into the stream
kernel2<<<grid, block, 0, stream1>>>();

// while the kernels are running do some work on the CPU
// but check if kernel1 has completed because then we will start
// a device to host copy in stream2
while ( not allCPUWorkDone() || not copyStarted ) {
    doNextChunkOfCPUWork();

    // peek to see if kernel 1 has completed
    // if so enqueue a non-blocking copy into stream2
    if ( not copyStarted ) {
        if( cudaEventQuery(event) == cudaSuccess ) {
            cudaMemcpyAsync(h_data, d_data, size, cudaMemcpyDeviceToHost, stream2);
            copyStarted = true;
        }
    }
}

// wait for both streams to be done
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);

// destroy the event
cudaEventDestroy(event);

// destroy the streams and free the data
cudaStreamDestroy(stream1);
cudaStreamDestroy(stream2);
cudaFree(d_data);
free(h_data);
```

--------------------------------

### Sibling Launch Example: cudaGraphLaunch with cudaStreamGraphFireAndForgetAsSibling

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/cuda-graphs.html

This example demonstrates a sibling launch, where a device graph is launched as a child of the launching graph's parent environment. The kernel launchSiblingGraph uses cudaGraphLaunch with cudaStreamGraphFireAndForgetAsSibling. The graphSetup function creates, instantiates, and uploads the device graph, then captures and instantiates the launching graph, and finally launches the host graph. Sibling launches do not gate tail launches enqueued by the launching graph.

```CUDA C++
__global__ void launchSiblingGraph(cudaGraphExec_t graph) {
    cudaGraphLaunch(graph, cudaStreamGraphFireAndForgetAsSibling);
}

void graphSetup() {
    cudaGraphExec_t gExec1, gExec2;
    cudaGraph_t g1, g2;

    // Create, instantiate, and upload the device graph.
    create_graph(&g2);
    cudaGraphInstantiate(&gExec2, g2, cudaGraphInstantiateFlagDeviceLaunch);
    cudaGraphUpload(gExec2, stream);

    // Create and instantiate the launching graph.
    cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);
    launchSiblingGraph<<<1, 1, 0, stream>>>(gExec2);
    cudaStreamEndCapture(stream, &g1);
    cudaGraphInstantiate(&gExec1, g1);

    // Launch the host graph, which will in turn launch the device graph.
    cudaGraphLaunch(gExec1, stream);
}
```

--------------------------------

### Kernel launch with dim3 grid and block

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/intro-to-cuda-cpp.html

Shows how to define a 2D grid and block using dim3 and launch a kernel with the <<<...>>> syntax. The grid and block dimensions are 16x16 and 8x8 respectively.

```cpp
dim3 grid(16,16); dim3 block(8,8); MatAdd<<<grid, block>>>(A, B, C);
```

--------------------------------

### Execution.Model.API.2: Spin-loop without CUDA query API

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cuda-cpp-execution-model.html

This example uses a system-scope atomic flag and a spin-loop waiting for the producer kernel to set it. The allowed outcome is that eventually no thread makes progress because the cudaDeviceSynchronize is only called after the flag is set, but CUDA only guarantees the producer starts if the synchronization API is called, so the host may spin forever.

```cpp
// Example: Execution.Model.API.2
// Allowed outcome: eventually, no thread makes progress.
// Rationale: the `cudaDeviceSynchronize` API below is only called if a device thread eventually makes progress and sets the flag.
// However, CUDA only guarantees that `producer` device thread eventually starts if the synchronization API is called.
// Therefore, the host thread may never be unblocked from the flag spin-loop.
cuda::atomic<int, cuda::thread_scope_system> flag = 0;
__global__ void producer() { flag.store(1); }
int main() {
    cudaHostRegister(&flag, sizeof(flag));
    producer<<<1,1>>>();
    while (flag.load() == 0);
    return cudaDeviceSynchronize();
}
```

--------------------------------

### Initialize CUDA by matching device UUIDs (simpleVulkan example)

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/graphics-interop.html

Use this function to select the CUDA device that corresponds to the Vulkan physical device that created the objects to be imported. It iterates over CUDA devices, compares each device's UUID with the provided Vulkan device UUID, and sets the matching device. The function returns the device index or -1 if no match is found, and exits if no CUDA devices are available or all are prohibited.

```cpp
// from the CUDA example `simpleVulkan`
int SineWaveSimulation::initCuda(uint8_t *vkDeviceUUID, size_t UUID_SIZE) {
  int current_device = 0;
  int device_count = 0;
  int devices_prohibited = 0;

  cudaDeviceProp deviceProp;
  checkCudaErrors(cudaGetDeviceCount(&device_count));

  if (device_count == 0) {
    fprintf(stderr, "CUDA error: no devices supporting CUDA.\n");
    exit(EXIT_FAILURE);
  }

  // Find the GPU which is selected by Vulkan
  while (current_device < device_count) {
    cudaGetDeviceProperties(&deviceProp, current_device);

    if ((deviceProp.computeMode != cudaComputeModeProhibited)) {
      // Compare the cuda device UUID with vulkan UUID
      int ret = memcmp((void *)&deviceProp.uuid, vkDeviceUUID, UUID_SIZE);
      if (ret == 0) {
        checkCudaErrors(cudaSetDevice(current_device));
        checkCudaErrors(cudaGetDeviceProperties(&deviceProp, current_device));
        printf("GPU Device %d: \"%s\" with compute capability %d.%d\n\n",
               current_device, deviceProp.name, deviceProp.major,
               deviceProp.minor);

        return current_device;
      }

    } else {
      devices_prohibited++;
    }

    current_device++;
  }

  if (devices_prohibited == device_count) {
    fprintf(stderr,
            "CUDA error:"
            " No Vulkan-CUDA Interop capable GPU found.\n");
    exit(EXIT_FAILURE);
  }

  return -1;
}
```

--------------------------------

### Create a CUDA Graph with a Conditional SWITCH Node

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/cuda-graphs.html

This example shows how to create a CUDA graph containing a conditional SWITCH node. The condition value is set by an upstream kernel using cudaGraphSetConditional. The body graphs are populated using the graph API, and the graph is instantiated and launched. Note that the condition value is undefined at the start of each graph execution if no default value is provided.

```cpp
__global__ void setHandle(cudaGraphConditionalHandle handle, int value)
{
    ...
    // Set the condition value to the value passed to the kernel
    cudaGraphSetConditional(handle, value);
    ...
}

void graphSetup() {
    cudaGraph_t graph;
    cudaGraphExec_t graphExec;
    cudaGraphNode_t node;
    void *kernelArgs[2];
    int value = 1;

    // Create the graph
    cudaGraphCreate(&graph, 0);

    // Create the conditional handle; because no default value is provided, the condition value is undefined at the start of each graph execution
    cudaGraphConditionalHandle handle;
    cudaGraphConditionalHandleCreate(&handle, graph);

    // Use a kernel upstream of the conditional to set the handle value
    cudaGraphNodeParams params = { cudaGraphNodeTypeKernel };
    params.kernel.func = (void *)setHandle;
    params.kernel.gridDim.x = params.kernel.gridDim.y = params.kernel.gridDim.z = 1;
    params.kernel.blockDim.x = params.kernel.blockDim.y = params.kernel.blockDim.z = 1;
    params.kernel.kernelParams = kernelArgs;
    kernelArgs[0] = &handle;
    kernelArgs[1] = &value;
    cudaGraphAddNode(&node, graph, NULL, 0, &params);

    // Create and add the conditional SWITCH node
    cudaGraphNodeParams cParams = { cudaGraphNodeTypeConditional };
    cParams.conditional.handle = handle;
    cParams.conditional.type   = cudaGraphCondTypeSwitch;
    cParams.conditional.size   = 5;
    cudaGraphAddNode(&node, graph, &node, 1, &cParams);

    // Get the body graphs of the conditional node
    cudaGraph_t *bodyGraphs = cParams.conditional.phGraph_out;

    // Populate the body graphs of the SWITCH conditional node
    ...
    cudaGraphAddNode(&node, bodyGraphs[0], NULL, 0, &params);
    ...
    cudaGraphAddNode(&node, bodyGraphs[4], NULL, 0, &params);

    // Instantiate and launch the graph
    cudaGraphInstantiate(&graphExec, graph, NULL, NULL, 0);
    cudaGraphLaunch(graphExec, 0);
    cudaDeviceSynchronize();

    // Clean up
    cudaGraphExecDestroy(graphExec);
    cudaGraphDestroy(graph);
}
```

--------------------------------

### barrier_arrive and barrier_wait for cluster shared memory initialization

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/device-callable-apis.html

Shows how to use cluster.barrier_arrive() and cluster.barrier_wait() to synchronize shared memory initialization across a cluster. The token from barrier_arrive must be passed to barrier_wait and cannot be reused.

```cpp
#include <cooperative_groups.h>

using namespace cooperative_groups;

void  __device__  init_shared_data(const thread_block& block, int *data);
void  __device__  local_processing(const thread_block& block);
void  __device__  process_shared_data(const thread_block& block, int *data);

 __global__  void cluster_kernel() {
 extern  __shared__  int array[];
 auto cluster = this_cluster();
 auto block = this_thread_block();

 // Use this thread block to initialize some shared state
 init_shared_data(block, &array[0]);

 auto token = cluster.barrier_arrive(); // Let other blocks know this block is running and data was initialized

 // Do some local processing to hide the synchronization latency
 local_processing(block);

 // Map data in shared memory from the next block in the cluster
 int *dsmem = cluster.map_shared_rank(&array[0], (cluster.block_rank() + 1) % cluster.num_blocks());

 // Make sure all other blocks in the cluster are running and initialized shared data before accessing dsmem
 cluster.barrier_wait(std::move(token));

 // Consume data in distributed shared memory
 process_shared_data(block, dsmem);
 cluster.sync();
}
```

--------------------------------

### C++ Partition View Load and Store

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/writing-tile-kernels.html

Demonstrates constructing a partition view in C++ and using its load and store methods to move tiles between global memory and tile space.

```APIDOC
## C++ Partition View Load and Store

### Description
Constructs a partition view over a tensor span and uses its load and store methods to access tiles by tile-space index.

### Methods
- `ct::tensor_span{ptr, extents}` - Constructs a tensor span from a raw pointer and extents.
- `ct::partition_view{span, shape}` - Constructs a partition view from a tensor span and tile shape.
- `view.load(idx...)` - Loads the tile at the given tile-space index.
- `view.store(tile, idx...)` - Stores the tile at the given tile-space index.

### Parameters
#### ct::tensor_span
- `ptr` (pointer) - Required - Raw pointer to the data.
- `extents` (ct::extents) - Required - Multidimensional shape of the data.

#### ct::partition_view
- `span` (ct::tensor_span) - Required - The tensor span to partition.
- `shape` (ct::shape) - Required - Fixed-size tile shape.

#### view.load
- `idx...` (integer) - Required - Tile-space index (e.g., block index).

#### view.store
- `tile` (tile) - Required - The tile to store.
- `idx...` (integer) - Required - Tile-space index.

### Example
```cpp
__tile_global__ void vec_add(float* __restrict__ a, float* __restrict__ b, float* __restrict__ out) {
    namespace ct = cuda::tiles;
    using namespace ct::literals;

    a   = ct::assume_aligned(a,   16_ic);
    b   = ct::assume_aligned(b,   16_ic);
    out = ct::assume_aligned(out, 16_ic);

    // Step 1: attach a shape to each raw pointer. 128_ic marks 128 as a compile-time constant.
    auto aSpan = ct::tensor_span{a,   ct::extents{128_ic}};
    auto bSpan = ct::tensor_span{b,   ct::extents{128_ic}};
    auto oSpan = ct::tensor_span{out, ct::extents{128_ic}};

    // Step 2: partition each span into a tile space of fixed 8-element tiles.
    auto aView = ct::partition_view{aSpan, ct::shape{8_ic}};
    auto bView = ct::partition_view{bSpan, ct::shape{8_ic}};
    auto oView = ct::partition_view{oSpan, ct::shape{8_ic}};

    int  bx    = ct::bid().x;             // this block's tile-space index along .x
    auto aTile = aView.load(bx);          // pick the bx-th tile of a
    auto bTile = bView.load(bx);
    oView.store(aTile + bTile, bx);       // write the tile back at the bx-th position of out
}
```
```

--------------------------------

### CUDA_MANAGED_FORCE_DEVICE_ALLOC examples

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/environment-variables.html

Examples of setting CUDA_MANAGED_FORCE_DEVICE_ALLOC to control physical memory allocation for managed memory.

```bash
CUDA_MANAGED_FORCE_DEVICE_ALLOC=0
CUDA_MANAGED_FORCE_DEVICE_ALLOC=1 # force device memory
```

--------------------------------

### Prefetch example with system allocator (malloc)

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/unified-memory.html

Shows prefetching a system-allocated buffer to the GPU and back to the CPU using cudaMemPrefetchAsync. The data is initialized on CPU, prefetched to the GPU, used in a kernel, then prefetched back to the host.

```cpp
void test_prefetch_sam(const cudaStream_t& s) {
  // initialize data on CPU
  char *data = (char*)malloc(dataSizeBytes);
  init_data(data, dataSizeBytes);
  cudaMemLocation location = {.type = cudaMemLocationTypeDevice, .id = myGpuId};

  // encourage data to move to GPU before use
  const unsigned int flags = 0;
  cudaMemPrefetchAsync(data, dataSizeBytes, location, flags, s);

  // use data on GPU
  const unsigned num_blocks = (dataSizeBytes + threadsPerBlock - 1) / threadsPerBlock;
  mykernel<<<num_blocks, threadsPerBlock, 0, s>>>(data, dataSizeBytes);

  // encourage data to move back to CPU
  location = {.type = cudaMemLocationTypeHost};
  cudaMemPrefetchAsync(data, dataSizeBytes, location, flags, s);

  cudaStreamSynchronize(s);

  // use data on CPU
  use_data(data, dataSizeBytes);
  free(data);
}
```

--------------------------------

### Example: Stream Compaction using exclusive_scan

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/device-callable-apis.html

Example demonstrating how to use exclusive_scan to compact data based on a predicate.

```APIDOC
## Example: Stream Compaction using exclusive_scan

### Description
This example shows how to use `cg::exclusive_scan` to compact data from an input array into an output array based on a predicate function.

### Code
```cpp
#include <cooperative_groups.h>
#include <cooperative_groups/scan.h>
namespace cg = cooperative_groups;

// put data from input into output only if it passes test_fn predicate
template<typename Group, typename Data, typename TyFn>
 __device__  int stream_compaction(Group &g, Data *input, int count, TyFn&& test_fn, Data *output) {
 int per_thread = count / g.num_threads();
 int thread_start = min(g.thread_rank() * per_thread, count);
 int my_count = min(per_thread, count - thread_start);

 // get all passing items from my part of the input
 // into a contagious part of the array and count them.
 int i = thread_start;
 while (i < my_count + thread_start) {
 if (test_fn(input[i])) {
 i++;
 }
 else {
 my_count--;
 input[i] = input[my_count + thread_start];
 }
 }

 // scan over counts from each thread to calculate my starting
 // index in the output
 int my_idx = cg::exclusive_scan(g, my_count);

 for (i = 0; i < my_count; ++i) {
 output[my_idx + i] = input[thread_start + i];
 }
 // return the total number of items in the output
 return g.shfl(my_idx + my_count, g.num_threads() - 1);
}
```
```

--------------------------------

### cuLaunchKernel - Launch Kernel with Execution Configuration

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/03-advanced/driver-api.html

Launches a kernel with a given execution configuration. Parameters can be passed as an array of pointers or via a parameter buffer using the CU_LAUNCH_PARAM_BUFFER_POINTER option.

```APIDOC
## cuLaunchKernel

### Description
Launches a kernel with a given execution configuration. Parameters are passed either as an array of pointers (next to last parameter) or as one of the extra options (last parameter) using CU_LAUNCH_PARAM_BUFFER_POINTER.

### Method
CUDA Driver API function

### Endpoint
cuLaunchKernel

### Parameters
- **f** (CUfunction) - Input - Kernel function handle.
- **gridDimX** (unsigned int) - Input - Grid dimension X.
- **gridDimY** (unsigned int) - Input - Grid dimension Y.
- **gridDimZ** (unsigned int) - Input - Grid dimension Z.
- **blockDimX** (unsigned int) - Input - Block dimension X.
- **blockDimY** (unsigned int) - Input - Block dimension Y.
- **blockDimZ** (unsigned int) - Input - Block dimension Z.
- **sharedMemBytes** (unsigned int) - Input - Dynamic shared memory size in bytes.
- **hStream** (CUstream) - Input - Stream handle.
- **kernelParams** (void**) - Input - Array of pointers to kernel parameters.
- **extra** (void**) - Input - Extra options, e.g., CU_LAUNCH_PARAM_BUFFER_POINTER.

### Request Example
```c
// Example using parameter buffer
#define ALIGN_UP(offset, alignment) \
 (offset) = ((offset) + (alignment) - 1) & ~((alignment) - 1)

char paramBuffer[1024];
size_t paramBufferSize = 0;
// ... (fill paramBuffer with parameters)
void* extra[] = {
    CU_LAUNCH_PARAM_BUFFER_POINTER, paramBuffer,
    CU_LAUNCH_PARAM_BUFFER_SIZE, &paramBufferSize,
    CU_LAUNCH_PARAM_END
};
cuLaunchKernel(f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, sharedMemBytes, hStream, NULL, extra);
```

### Response
- **err** (CUresult) - Returns CUDA_SUCCESS on success, otherwise an error code.

### Response Example
```c
// On success, err == CUDA_SUCCESS.
```
```

--------------------------------

### Vulkan extension setup for external memory and semaphores

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/graphics-interop.html

Returns the required Vulkan instance and device extensions for external memory and semaphore interoperability, including platform-specific extensions for Windows and Unix.

```cpp
  std::vector<const char *> getRequiredExtensions() const {
    std::vector<const char *> extensions;
    extensions.push_back(VK_KHR_EXTERNAL_MEMORY_CAPABILITIES_EXTENSION_NAME);
    extensions.push_back(VK_KHR_EXTERNAL_SEMAPHORE_CAPABILITIES_EXTENSION_NAME);
    return extensions;
  }

  std::vector<const char *> getRequiredDeviceExtensions() const {
    std::vector<const char *> extensions;
    extensions.push_back(VK_KHR_EXTERNAL_MEMORY_EXTENSION_NAME);
    extensions.push_back(VK_KHR_EXTERNAL_SEMAPHORE_EXTENSION_NAME);
    extensions.push_back(VK_KHR_TIMELINE_SEMAPHORE_EXTENSION_NAME);
#ifdef _WIN64
    extensions.push_back(VK_KHR_EXTERNAL_MEMORY_WIN32_EXTENSION_NAME);
    extensions.push_back(VK_KHR_EXTERNAL_SEMAPHORE_WIN32_EXTENSION_NAME);
#else
    extensions.push_back(VK_KHR_EXTERNAL_MEMORY_FD_EXTENSION_NAME);
    extensions.push_back(VK_KHR_EXTERNAL_SEMAPHORE_FD_EXTENSION_NAME);
#endif /* _WIN64 */
    return extensions;
  }
```

--------------------------------

### cudaMemGetMemPool

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/locality-domains.html

Gets the current memory pool for a given device and location. Can be used to get the current pool for a locality domain.

```APIDOC
## cudaMemGetMemPool

### Description
Gets the current memory pool for a given device and location. The location can be a device locality domain.

### Method
Runtime API function

### Endpoint
`cudaMemGetMemPool(cudaMemPool_t* memPool, int device, cudaMemLocation* location)`

### Parameters
#### Input
- **memPool** (cudaMemPool_t*) - Output - Pointer to the current memory pool.
- **device** (int) - Input - The device for which to get the current pool.
- **location** (cudaMemLocation*) - Input - The location for which to get the current pool.

### Response
#### Success Response
- **cudaError_t** - Returns `cudaSuccess` on success, or an error code on failure.

#### Response Example
```cpp
cudaError_t status = cudaMemGetMemPool(&pool, device, &location);
```
```
Tokens:
Raw