### Stream Compaction Example Setup

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

This C++ code sets up a device function for stream compaction using cooperative groups. It calculates thread-specific work distribution and initializes a loop to process input data.

```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) {

```

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

### Example: Odd/Even Partitioning with binary_partition

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

This example demonstrates how to use `binary_partition` to divide a 32-thread tile into two groups: one for threads processing odd numbers and another for threads processing even numbers from an input array.

```cuda
_global__  void oddEven(int *inputArr) {
 auto block = cg::this_thread_block();
 auto tile32 = cg::tiled_partition<32>(block);

 // inputArr contains random integers
 int elem = inputArr[block.thread_rank()];
 // after this, tile32 is split into 2 groups,
 // a subtile where elem&1 is true and one where its false
 auto subtile = cg::binary_partition(tile32, (elem & 1));
}
```

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

### Example of cudaMemPrefetchAsync and Kernel Launch

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

Demonstrates prefetching data to a specific device, launching a kernel, and then prefetching data back to the host.

```c
  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
  location = {.type = cudaMemLocationTypeHost};
  cudaMemPrefetchAsync(data, dataSizeBytes, location, flags, s);

  cudaStreamSynchronize(s);

  // use data on CPU
  use_data(data, dataSizeBytes);
  cudaFree(data);
}
// test-prefetch-managed-end
```

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

### CUDA Source File Example

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

A simple CUDA C++ program demonstrating a kernel launch and synchronization.

```cuda-c++
#include <stdio.h>
 __global__  void kernel() {
 printf("Hello from kernel\n");
}

void kernel_launcher() {
 kernel<<<1, 1>>>();
 cudaDeviceSynchronize();
}

int main() {
 kernel_launcher();
 return 0;
}
```

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

### Host-Side Kernel Launch Configuration

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

Example of launching a kernel from the host, specifying shared memory size.

```cuda-c++
void host_launch(int *data) {
 permute<<< 1, 256, 256*sizeof(int) >>>(256, data);
}
```

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

### Family-Specific Compatibility Table Examples

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/compute-capabilities.html

Examples of family-specific compilation targets and their compatible compute capabilities, illustrating the 'f' suffix usage for feature sets common to GPU families.

```text
compute_100f
```

```text
compute_103f
```

```text
compute_110f
```

```text
compute_120f
```

```text
compute_121f
```

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

### Example of cudaMemAdvise with System Allocator

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

Shows how to use cudaMemAdvise with memory allocated via malloc, setting read-mostly hints and prefetching for read duplication.

```c
void test_advise_sam(cudaStream_t stream) {
  char *dataPtr;
  size_t dataSize = 64 * threadsPerBlock;  // 16 KiB

  // Allocate memory using malloc or cudaMallocManaged
  dataPtr = (char*)malloc(dataSize);

  // Set the advice on the memory region
  cudaMemLocation loc = {.type = cudaMemLocationTypeDevice, .id = myGpuId};
  cudaMemAdvise(dataPtr, dataSize, cudaMemAdviseSetReadMostly, loc);

  int outerLoopIter = 0;
  while (outerLoopIter < maxOuterLoopIter) {
    // The data is written by the CPU each outer loop iteration
    init_data(dataPtr, dataSize);

    // The data is made available to all GPUs by prefetching.
    // Prefetching here causes read duplication of data instead
    // of data migration
    cudaMemLocation location;
    location.type = cudaMemLocationTypeDevice;
    for (int device = 0; device < maxDevices; device++) {
      location.id = device;
      const unsigned int flags = 0;
      cudaMemPrefetchAsync(dataPtr, dataSize, location, flags, stream);
    }

    // The kernel only reads this data in the inner loop
    int innerLoopIter = 0;
    while (innerLoopIter < maxInnerLoopIter) {
      mykernel<<<32, threadsPerBlock, 0, stream>>>((const char *)dataPtr, dataSize);
      innerLoopIter++;
    }
    outerLoopIter++;
  }

  free(dataPtr);
}
```

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

### Per-Thread Allocation Example Kernel

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

Demonstrates per-thread memory allocation and deallocation within a CUDA kernel using malloc and free. This example also includes setting the heap size before kernel launch.

```cpp
#include <stdlib.h>
#include <stdio.h>

__global__ void single_thread_allocation_kernel() {
    size_t size = 123;
    char*  ptr  = (char*) malloc(size);
    memset(ptr, 0, size);
    printf("Thread %d got pointer: %p\n", threadIdx.x, ptr);
    free(ptr);
}

int main() {
    // Set a heap size of 128 megabytes.
    // Note that this must be done before any kernel is launched.
    cudaDeviceSetLimit(cudaLimitMallocHeapSize, 128 * 1024 * 1024);

```

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

### Prefetching Data for Managed Memory

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

Example demonstrating `cudaMemPrefetchAsync` with managed memory, encouraging data migration to the GPU before kernel execution and back to the CPU afterwards.

```c
void test_prefetch_managed(const cudaStream_t& s) {
  // initialize data on CPU
  char *data;
  cudaMallocManaged(&data, 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 uinsigned 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);
  cudaFree(data);
}
```

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

### Kernel Invocation Example

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

This snippet demonstrates how to launch a CUDA kernel using the driver API, including setting up thread and block dimensions and passing arguments.

```APIDOC
## Kernel Invocation

### Description
Invokes a previously loaded CUDA kernel with specified grid and block dimensions, and arguments.

### Method
`cuLaunchKernel`

### Parameters
- `kernel` (CUfunction) - Handle to the kernel function.
- `gridDimX`, `gridDimY`, `gridDimZ` (unsigned int) - Dimensions of the grid.
- `blockDimX`, `blockDimY`, `blockDimZ` (unsigned int) - Dimensions of the block.
- `sharedMemBytes` (unsigned int) - Shared memory size in bytes.
- `hStream` (CUstream) - Handle to the CUDA stream.
- `args` (void**) - Array of pointers to kernel arguments.
- `extra` (void*) - Optional extra arguments.

### Request Example
```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);
```
```

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

### Example: Min of two unsigned 32-bit integers with index

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

Shows how to use `__vibmin_u32` to find the minimum of two unsigned 32-bit integers and determine which of the original parameters was smaller.

```cpp
unsigned a = 9;
unsigned b = 6;
bool     smaller_value;
unsigned min_value = __vibmin_u32(a, b, &smaller_value); // min_value is 6, smaller_value is true
```

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

### CUDA Graph Tail Self-Launch Example

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

Demonstrates a graph relaunching itself using `cudaGetCurrentGraphExec` and `cudaStreamGraphTailLaunch`. Ensure `relaunchCount` is managed appropriately to avoid infinite loops.

```c++
__device__ int relaunchCount = 0;

__global__ void relaunchSelf() {
    int relaunchMax = 100;

    if (threadIdx.x == 0) {
        if (relaunchCount < relaunchMax) {
            cudaGraphLaunch(cudaGetCurrentGraphExec(), cudaStreamGraphTailLaunch);
        }

        relaunchCount++;
    }
}
```

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

### Access New CUDA Driver API with Older Toolkit

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/driver-entry-point-access.html

Demonstrates how to retrieve and use a new CUDA driver API (e.g., `cuFoo` from CUDA 12.0) when using an older CUDA toolkit (e.g., 11.3) but having a compatible newer driver installed. It involves manually defining the function prototype and using `cuGetProcAddress` to get the function pointer.

```c
int main()
{
 // Assuming we have CUDA 12.0 driver installed.

 // Manually define the prototype as cudaTypedefs.h in CUDA 11.3 does not have the cuFoo typedef
 typedef CUresult (CUDAAPI *PFN_cuFoo)(...);
 PFN_cuFoo pfn_cuFoo = NULL;
 CUdriverProcAddressQueryResult driverStatus;

 // Get the address for cuFoo API using cuGetProcAddress. Specify CUDA version as
 // 12000 since cuFoo was introduced then or get the driver version dynamically
 // using cuDriverGetVersion
 int driverVersion;
 cuDriverGetVersion(&driverVersion);
 CUresult status = cuGetProcAddress("cuFoo", &pfn_cuFoo, driverVersion, CU_GET_PROC_ADDRESS_DEFAULT, &driverStatus);

 if (status == CUDA_SUCCESS && pfn_cuFoo) {
 pfn_cuFoo(...);
 }
 else {
 printf("Cannot retrieve the address to cuFoo - driverStatus = %d. Check if the latest driver for CUDA 12.0 is installed.\n", driverStatus);
 assert(0);
 }

 // rest of code here

}
```

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

### Example CUDA Error Log Output

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

This is an example of an error message generated by the Error Log Management system when attempting to dump logs to an unallocated buffer.

```text
[22:21:32.099][25642][CUDA][E][cuLogsDumpToMemory] buffer cannot be NULL
```

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

### Compiling Dynamic Parallelism Example

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

Command to build the 'Hello World' dynamic parallelism CUDA program. Ensure to use appropriate architecture flags and link the CUDA device runtime library.

```bash
$ nvcc -arch=sm_75 -rdc=true hello_world.cu -o hello -lcudadevrt
```

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

### Programmatic Event Trigger at Block Start (True)

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

Configure a launch attribute for a programmatic event, specifying that the trigger should occur at block start. This is used for CUDA graph dependencies.

```c++
cudaLaunchAttribute attribute;
attribute.id = cudaLaunchAttributeProgrammaticEvent;
attribute.val.programmaticEvent.triggerAtBlockStart = 1;

```

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

### Prefetching Data for System Allocator

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

Example demonstrating how to use `cudaMemPrefetchAsync` to encourage data migration to the GPU before kernel execution and back to the CPU afterwards when using system-allocated memory.

```c
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);
}
```

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

### Programmatic Event Trigger at Block Start (False)

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

Configure a launch attribute for a programmatic event, specifying that the trigger should not occur at block start. This is relevant for CUDA graph dependencies.

```c++
cudaLaunchAttribute attribute;
attribute.id = cudaLaunchAttributeProgrammaticEvent;
attribute.val.programmaticEvent.triggerAtBlockStart = 0;

```

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

### CUDA Graphs: Capturing, Instantiating, and Launching

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

Demonstrates the process of capturing a sequence of kernel launches into a CUDA graph, instantiating it, and then repeatedly launching the graph for efficient execution. This approach minimizes CPU overhead for repetitive tasks.

```c++
#define N 500000 // tuned such that kernel takes a few microseconds

// A very lightweight kernel
__global__ void shortKernel(float * out_d, float * in_d){
    int idx=blockIdx.x*blockDim.x+threadIdx.x;
    if(idx<N) out_d[idx]=1.23*in_d[idx];
}

bool graphCreated=false;
cudaGraph_t graph;
cudaGraphExec_t instance;

// The graph will be executed NSTEP times
for(int istep=0; istep<NSTEP; istep++){
    if(!graphCreated){
        // Capture the graph
        cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);

        // Launch NKERNEL kernels
        for(int ikrnl=0; ikrnl<NKERNEL; ikrnl++){
            shortKernel<<<blocks, threads, 0, stream>>>(out_d, in_d);
        }

        // End the capture
        cudaStreamEndCapture(stream, &graph);

        // Instantiate the graph
        cudaGraphInstantiate(&instance, graph, NULL, NULL, 0);
        graphCreated=true;
    }

    // Launch the graph
    cudaGraphLaunch(instance, stream);

    // Synchronize the stream
    cudaStreamSynchronize(stream);
}

```

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

### CUDA Graph Tail Launch Setup

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

Sets up a graph for tail launch, where one graph launches another upon completion. Requires stream and graph creation, instantiation, and upload.

```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);
}
```

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

### Warp Shuffle Function Example

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

Demonstrates the use of warp shuffle functions for collective operations within a warp. This example shows a warp-reduce operation where values are aggregated across threads in a warp.

```C++
int main() {
    warp_reduce_kernel<<<1, 32>>>();
    cudaDeviceSynchronize();
    return 0;
}
```
```

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

### CUDA Kernel for Unified Memory Examples

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

A simple CUDA kernel that prints the first 8 characters of an input character array. This kernel is used across various unified memory examples.

```cuda
__global__ void kernel(const char* type, const char* data) {
  static const int n_char = 8;
  printf("%s - first %d characters: '", type, n_char);
  for (int i = 0; i < n_char; ++i) printf("%c", data[i]);
  printf("'\n");
}
```

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

### Hello World Dynamic Parallelism Example

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

A basic CUDA program demonstrating dynamic parallelism. It includes a parent kernel that launches a child kernel and a tail kernel into the cudaStreamTailLaunch stream, followed by host synchronization.

```cuda
#include <stdio.h>

__global__ void childKernel()
{
    printf("Hello ");
}

__global__ void tailKernel()
{
    printf("World!\n");
}

__global__ void parentKernel()
{
    // launch child
    childKernel<<<1,1>>>();
    if (cudaSuccess != cudaGetLastError()) {
        return;
    }

    // launch tail into cudaStreamTailLaunch stream
    // implicitly synchronizes: waits for child to complete
    tailKernel<<<1,1,0,cudaStreamTailLaunch>>>();

}

int main(int argc, char *argv[])
{
    // launch parent
    parentKernel<<<1,1>>>();
    if (cudaSuccess != cudaGetLastError()) {
        return 1;
    }

    // wait for parent to complete
    if (cudaSuccess != cudaDeviceSynchronize()) {
        return 2;
    }

    return 0;
}
```

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

### CUDA Driver API Initialization and Setup

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

This snippet demonstrates the initialization of the CUDA driver API, device enumeration, context creation, module loading, memory allocation on the device, and obtaining a function handle for kernel execution.

```APIDOC
## CUDA Driver API Initialization and Setup

### Description
This example shows the basic steps to initialize the CUDA driver API, select a device, create a context, load a CUDA module (PTX or binary), allocate memory on the device, and get a handle to a kernel function.

### Initialization
```c
// Initialize the CUDA driver API
cuInit(0);
```

### Device Handling
```c
// 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);
```

### Context Management
```c
// Create context
CUcontext cuContext;
cuCtxCreate(&cuContext, 0, cuDevice);
```

### Module and Function Loading
```c
// Create module from binary file
CUmodule cuModule;
cuModuleLoad(&cuModule, "VecAdd.ptx");

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

### Device Memory Management
```c
// 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);
```
```

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

### Invalid Warp Shuffle Usage Example

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

Illustrates an example of incorrect usage for warp shuffle functions. This snippet is intended to highlight common pitfalls or configurations that lead to undefined behavior, such as incorrect lane IDs or masks.

```cpp
int laneId = threadIdx.x % warpSize;
int value  = ...
```

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

### Create Streams and Launch Kernels on Specific Devices

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/03-advanced/multi-gpu-systems.html

Illustrates creating streams on different devices and launching kernels using these streams. A kernel launch to a stream not associated with the current device will fail.

```c
cudaSetDevice(0); // Set device 0 as current
cudaStream_t s0;
cudaStreamCreate(&s0); // Create stream s0 on device 0
MyKernel<<<100, 64, 0, s0>>>(); // Launch kernel on device 0 in s0

cudaSetDevice(1); // Set device 1 as current
cudaStream_t s1;
cudaStreamCreate(&s1); // Create stream s1 on device 1
MyKernel<<<100, 64, 0, s1>>>(); // Launch kernel on device 1 in s1

// This kernel launch will fail, since stream s0 is not associated to device 1:
MyKernel<<<100, 64, 0, s0>>>(); // Launch kernel on device 1 in s0
```

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

### Create, Instantiate, and Launch CUDA Graph

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

A complete example demonstrating stream capture to create a graph, followed by instantiation and launching. Ensure the stream is properly initialized before capture.

```c
cudaGraph_t graph;

cudaStreamBeginCapture(stream);

kernel_A<<< ..., stream >>>(...);
kernel_B<<< ..., stream >>>(...);
libraryCall(stream);
kernel_C<<< ..., stream >>>(...);

cudaStreamEndCapture(stream, &graph);

cudaGraphExec_t graphExec;
cudaGraphInstantiate(&graphExec, graph, NULL, NULL, 0);
cudaGraphLaunch(graphExec, stream);
```

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

### CUDA Warp Active Mask Example - Predicate Check

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

This example demonstrates an incorrect usage of `__activemask()` to check if at least one thread's predicate evaluates to true. The `at_least_one` variable's value is non-deterministic.

```cpp
// Check whether at least one thread's predicate evaluates to true
if (pred) {
    // Invalid: the value of 'at_least_one' is non-deterministic
    // and could vary between executions.
    at_least_one = __activemask() > 0;
}
```

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

### Launch Kernel with Thread Block Clusters on Green Context Stream

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

Configure and launch a kernel with thread block clusters on a specified green context stream. This example demonstrates setting launch attributes, including cluster dimensions, and using occupancy APIs to determine potential cluster sizes and active cluster counts.

```c++
// Assume cudaStream_t gc_stream has already been created and a __global__ void cluster_kernel exists.

// Uncomment to support non portable cluster size, if possible
// CUDA_CHECK(cudaFuncSetAttribute(cluster_kernel, cudaFuncAttributeNonPortableClusterSizeAllowed, 1))

cudaLaunchConfig_t config = {0};
config.gridDim = grid_dim; // has to be a multiple of cluster dim.
config.blockDim = block_dim;
config.dynamicSmemBytes = expected_dynamic_shared_mem;

cudaLaunchAttribute attribute[1];
attribute[0].id = cudaLaunchAttributeClusterDimension;
attribute[0].val.clusterDim.x = 1;
attribute[0].val.clusterDim.y = 1;
attribute[0].val.clusterDim.z = 1;
config.attrs = attribute;
config.numAttrs = 1;

config.stream=gc_stream; // Need to pass the CUDA stream that will be used for that kernel

int max_potential_cluster_size = 0;
// the next call will ignore cluster dims in launch config
CUDA_CHECK(cudaOccupancyMaxPotentialClusterSize(&max_potential_cluster_size, cluster_kernel, &config));
std::cout << "max potential cluster size is " << max_potential_cluster_size << " for CUDA stream gc_stream" << std::endl;

// Could choose to update launch config's clusterDim with max_potential_cluster_size.
// Doing so would result in a successful cudaLaunchKernelEx call for the same kernel and launch config.

int num_clusters= 0;
CUDA_CHECK(cudaOccupancyMaxActiveClusters(&num_clusters, cluster_kernel, &config));
std::cout << "Potential max. active clusters count is " << num_clusters << std::endl;
```

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

### Example of `__syncwarp()` for Shared Memory Access

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

This example demonstrates synchronizing threads within a warp using `__syncwarp()` to safely access a shared memory array. It ensures that data written by one thread is visible to others before proceeding.

```cpp
__global__ void example_syncwarp(int* input_data, int* output_data) {
    if (threadIdx.x < warpSize) {
        __shared__ int shared_data[warpSize];
        shared_data[threadIdx.x] = input_data[threadIdx.x];

        __syncwarp(); // equivalent to __syncwarp(0xFFFFFFFF)
        if (threadIdx.x == 0)
            output_data[0] = shared_data[1];
    }
}
```

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

### CUDA Kernels for Dynamic Parallelism Example

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

Demonstrates parent, child, and tail launch kernels. The child kernel is launched from the parent, and a tail kernel is launched into `cudaStreamTailLaunch` to ensure modifications are visible after the child grid exits.

```cuda
__global__ void tail_launch(int *data) {
   data[threadIdx.x] = data[threadIdx.x]+1;
}
```

```cuda
__global__ void child_launch(int *data) {
   data[threadIdx.x] = data[threadIdx.x]+1;
}
```

```cuda
__global__ void parent_launch(int *data) {
   data[threadIdx.x] = threadIdx.x;

   __syncthreads();

   if (threadIdx.x == 0) {
       child_launch<<< 1, 256 >>>(data);
       tail_launch<<< 1, 256, 0, cudaStreamTailLaunch >>>(data);
   }
}
```

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

### Invalid Warp Intrinsics Usage Examples

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

Illustrates common mistakes in using warp synchronization intrinsics, such as incorrect mask assignments or divergent mask usage across threads. These examples highlight conditions that lead to kernel hangs or undefined behavior.

```cuda
if (threadIdx.x < 4) {           // threads 0, 1, 2, 3 are active
    __all_sync(0b0000011, pred); // WRONG, threads 2, 3 are active but not set in mask
    __all_sync(0b1111111, pred); // WRONG, threads 4, 5, 6 are not active but set in mask
}

// WRONG, participating threads have a different and overlapping mask
__all_sync(threadIdx.x == 0 ? 1 : 0xFFFFFFFF, pred);
```

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

### Device Memory Heap Management API

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

APIs for getting and setting the size of the device memory heap.

```APIDOC
## Device Memory Heap Management API

### Description
Manages the size of the device memory heap, which must be specified before allocating or freeing memory in device code.

### Functions
- `cudaDeviceGetLimit(size_t* size, cudaLimitMallocHeapSize)`: Retrieves the current heap size limit.
- `cudaDeviceSetLimit(cudaLimitMallocHeapSize, size_t size)`: Sets the device memory heap size.

### Behavior
- The heap size must be set before any program that allocates or frees memory in device code.
- A default heap of eight megabytes is allocated if not explicitly specified.
- The heap size granted will be at least `size` bytes.
- The actual memory allocation occurs when a module is loaded.
- If memory allocation fails, a `CUDA_ERROR_SHARED_OBJECT_INIT_FAILED` error is generated.
- The heap size cannot be changed after a module has been loaded and does not dynamically resize.
- Memory reserved for the device heap is in addition to memory allocated via host-side CUDA API calls.
```

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

### Unsupported Local Memory Atomic Operation

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

Example showing an atomic operation applied to local memory, which is disallowed.

```cpp
 // Not permitted to be applied to local memory
__device__ void foo() {
   unsigned a = 1, b;
   __nv_atomic_load(&a, &b, __NV_ATOMIC_RELAXED, __NV_THREAD_SCOPE_SYSTEM);
}
```

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

### Unsupported Host Function Atomic Operation

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

Example demonstrating an atomic operation called from a host function, which is not permitted.

```cpp
 // Not permitted in a host function
 __host__ void bar() {
     unsigned u1 = 1, u2 = 2;
     __nv_atomic_load(&u1, &u2, __NV_ATOMIC_RELAXED, __NV_THREAD_SCOPE_SYSTEM);
 }
```

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

### Unsupported Function Address Taking

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

Illustrates that the address of atomic built-in functions cannot be taken, for example, as a template default argument.

```cpp
 // Not permitted as a template default argument.
 // The function address cannot be taken.
 template<void *F = __nv_atomic_load_n>
 class X {
     void *f = F; // The function address cannot be taken.
 };
```

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

### Dynamic Runtime Version with cuGetProcAddress

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/driver-entry-point-access.html

This example demonstrates using cuGetProcAddress with a dynamically obtained runtime CUDA version. Be cautious, as mismatches between the typedef and the resolved function pointer can lead to undefined behavior.

```c
#include <cudaTypedefs.h>

CUuuid uuid;
CUdevice dev;
CUresult status;
int cudaVersion;
CUdriverProcAddressQueryResult driverStatus;

status = cuDeviceGet(&dev, 0); // Get device 0
// handle status

status = cuDriverGetVersion(&cudaVersion);
// handle status

PFN_cuDeviceGetUuid pfn_cuDeviceGetUuid;
status = cuGetProcAddress("cuDeviceGetUuid", &pfn_cuDeviceGetUuid, cudaVersion, CU_GET_PROC_ADDRESS_DEFAULT, &driverStatus);
if(CUDA_SUCCESS == status && pfn_cuDeviceGetUuid) {
 // pfn_cuDeviceGetUuid points to ???
}
```

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

### Kernel Launch with Max Threads per Block

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

Demonstrates a common but incorrect way to launch a kernel using a macro that relies on `__CUDA_ARCH__` in host code. This can lead to unexpected thread counts.

```cpp
// Host code
MyKernel<<<blocksPerGrid, MY_KERNEL_MAX_THREADS>>>(...);

```

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

### External Variable Setup and Teardown

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

Code for setting up and tearing down memory for an external global variable, often managed by a third-party library.

```c
/** This may be a non-CUDA file */
char* ext_data;
static const char global_string[] = "Hello World";

void __attribute__ ((constructor)) setup(void) {
  ext_data = (char*)malloc(sizeof(global_string));
  strncpy(ext_data, global_string, sizeof(global_string));
}

void __attribute__ ((destructor)) tear_down(void) {
  free(ext_data);
}
```

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

### Unified Memory Example

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

Demonstrates the use of unified memory for allocating buffers and launching a kernel. Unified memory simplifies memory management by allowing the CPU and GPU to access the same memory space. cudaDeviceSynchronize() is used to ensure kernel completion before proceeding.

```cuda-cpp
//unified-memory-begin
void unifiedMemExample(int vectorLength) {
    // Pointers to memory vectors
    float* A = nullptr;
    float* B = nullptr;
    float* C = nullptr;
    float* comparisonResult = (float*)malloc(vectorLength*sizeof(float));

    // Use unified memory to allocate buffers
    cudaMallocManaged(&A, vectorLength*sizeof(float));
    cudaMallocManaged(&B, vectorLength*sizeof(float));
    cudaMallocManaged(&C, vectorLength*sizeof(float));

    // Initialize vectors on the host
    initArray(A, vectorLength);
    initArray(B, vectorLength);

    // Launch the kernel. Unified memory will make sure A, B, and C are
    // accessible to the GPU
    int threads = 256;
    int blocks = cuda::ceil_div(vectorLength, threads);
    vecAdd<<<blocks, threads>>>(A, B, C, vectorLength);

    // Wait for the kernel to complete execution
    cudaDeviceSynchronize();

    // Perform computation serially on CPU for compari
```

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

### Unsupported Constructor Initialization List Atomic Operation

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

Example showing an atomic operation called within a constructor initialization list, which is not permitted.

```cpp
 // Not permitted to be called in a constructor initialization list.
 class Y {
     int a;
 public:
     __device__ Y(int *b): a(__nv_atomic_load_n(b, __NV_ATOMIC_RELAXED)) {}
 };
```

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

### Launch Sibling Graph with CUDA

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

Demonstrates how to launch a graph as a sibling to the launching graph's parent environment using `cudaGraphLaunch` with `cudaStreamGraphFireAndForgetAsSibling`. This setup is useful for creating independent execution paths.

```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);
}

```

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

### Calculating Number of Blocks for Kernel Launch (Manual)

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

Demonstrates a common method for calculating the number of blocks required for a kernel launch when the vector length might not be a multiple of the threads per block. This manual calculation uses integer division with an offset to achieve a ceiling effect.

```cuda-c++
// vectorLength is an integer storing number of elements in the vector
int threads = 256;
int blocks = (vectorLength + threads-1)/threads;
vecAdd<<<blocks, threads>>>(devA, devB, devC, vectorLength);
```

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

### Dumping Error Logs to a File

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

Writes the error log buffer to a specified file. If an iterator is provided, logs are dumped starting from that entry.

```APIDOC
## Dumping Error Logs to a File

### Description
Dumps the error log buffer to a specified file. If `iterator` is NULL, the entire buffer is dumped (up to 100 entries). If `iterator` is not NULL, logs are dumped starting from that entry, and the `iterator` is updated to the current end of the logs.

### Method
`CUresult cuLogsDumpToFile(CUlogIterator *iterator, const char *pathToFile, unsigned int flags)`

### Parameters
#### Path Parameters
- `iterator` (CUlogIterator *) - Optional - If NULL, dumps the entire buffer. If not NULL, dumps from this entry and updates the iterator.
- `pathToFile` (const char *) - Required - The path to the file where logs will be written.
- `flags` (unsigned int) - Required - Currently must be 0. Additional options are reserved for future releases.
```

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

### Initialize CUDA Driver API and Create Context

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

Initializes the CUDA driver API, retrieves the device count, selects a device, and creates a CUDA context for the calling host thread. Ensure `cuInit()` is called before any other driver API function.

```c
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");

```

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

### Sequential Grid Launches in Main Function with Tail Launch

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

Demonstrates how a subsequent grid launch (P2) waits for the completion of a grid launched with the tail launch stream (C within P1).

```cuda
__global__  void P1( ... ) {
 C<<< ... , cudaStreamTailLaunch >>>( ... );
}

 __global__  void P2( ... ) {
}
```

```c++
int main ( ... ) {
 ...
 P1<<< ... >>>( ... );
 P2<<< ... >>>( ... );
 ...
}
```

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

### Calling Kernel with Global-Scope Extern Variable

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

Example of calling the CUDA kernel with an extern global variable, demonstrating interaction with memory managed by other libraries.

```c
// declared in separate file, see below
extern char* ext_data;

void test_extern() {
  kernel<<<1, 1>>>("extern", ext_data);
  ASSERT(cudaDeviceSynchronize() == cudaSuccess,
    "CUDA failed with '%s'", cudaGetErrorString(cudaGetLastError()));
}
```

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

### Configuring and Launching Kernels with Programmatic Stream Serialization

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

Sets up the launch configuration for the secondary kernel with the programmatic stream serialization attribute and launches both kernels. The secondary kernel must be launched using the extensible launch API.

```cuda-c
cudaLaunchAttribute attribute[1];
attribute[0].id = cudaLaunchAttributeProgrammaticStreamSerialization;
attribute[0].val.programmaticStreamSerializationAllowed = 1;
configSecondary.attrs = attribute;
configSecondary.numAttrs = 1;

primary_kernel<<<grid_dim, block_dim, 0, stream>>>();
cudaLaunchKernelEx(&configSecondary, secondary_kernel);
```

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

### Baseline Compiler Target Example

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/compute-capabilities.html

The baseline compilation target does not allow the use of architecture-specific features and is compatible with all devices of the specified compute capability and later.

```text
compute_100
```

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

### Compiling and Linking for Tensor Core Fragment Compatibility

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

Demonstrates the compilation and linking process for ensuring tensor core fragment compatibility across different architectures. Incorrect linking can lead to runtime issues.

```bash
// sm_70 fragment layout
$> nvcc -dc -arch=compute_70 -code=sm_70 fragA.cu -o fragA.o
// sm_75 fragment layout
$> nvcc -dc -arch=compute_75 -code=sm_75 fragB.cu -o fragB.o
// Linking the two together
$> nvcc -dlink -arch=sm_75 fragA.o fragB.o -o frag.o
```