cuda

CUDA

Purpose

Guide agents through NVIDIA CUDA C/C++ development: kernel launch configuration, the memory hierarchy from registers through global memory, asynchronous execution with streams, nvcc compilation flags, Thrust library usage, and diagnosing common performance pitfalls like warp divergence and uncoalesced memory access.

When to Use

• Writing or optimizing a CUDA kernel for matrix multiply, reduction, or stencil operations
• Choosing block/grid dimensions and estimating occupancy
• Debugging slow kernels due to memory access patterns or low occupancy
• Setting up multi-stream pipelines with async cudaMemcpyAsync
• Compiling with nvcc and selecting architecture flags (-gencode)
• Using Thrust for parallel primitives instead of hand-written kernels

Workflow

1. Minimal kernel and launch

CUDA organizes work as threads grouped into blocks, blocks grouped into a grid.

Thread hierarchy
├── grid (1D/2D/3D)
│   └── block (1D/2D/3D, max 1024 threads)
│       └── thread (threadIdx, blockIdx, blockDim, gridDim)

// vector_add.cu
#include <cuda_runtime.h>
#include <stdio.h>

__global__ void vector_add(const float *a, const float *b, float *c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n)
        c[i] = a[i] + b[i];
}

int main(void) {
    const int n = 1 << 20;
    size_t bytes = n * sizeof(float);
    float *h_a, *h_b, *h_c, *d_a, *d_b, *d_c;

    cudaMalloc(&d_a, bytes);
    cudaMalloc(&d_b, bytes);
    cudaMalloc(&d_c, bytes);
    // ... host init and cudaMemcpy H2D ...

    int threads = 256;
    int blocks = (n + threads - 1) / threads;
    vector_add<<<blocks, threads>>>(d_a, d_b, d_c, n);
    cudaDeviceSynchronize();

    cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    return 0;
}

2. Memory hierarchy

Memory	Scope	Latency	Typical use
Registers	Per-thread	~1 cycle	Local scalars, loop indices
Shared (__shared__)	Per-block	~5 cycles	Tile data, halo exchange
Global	All threads	~400+ cycles	Large arrays, coalesced access
Constant (__constant__)	Read-only, cached	Fast broadcast	Kernel parameters, lookup tables
Texture	Cached 2D access	Cached	Image sampling, irregular reads