cuda

name	cuda
description	Comprehensive reference documentation and skill for NVIDIA CUDA C++ - the parallel computing platform and programming model for GPU acceleration. Covers CUDA Programming Guide (Release 13.2) and CUDA C++ Best Practices Guide (Release 13.2). Includes programming model, memory management, asynchronous execution, CUDA graphs, cooperative groups, advanced synchronization, async data copies, TMA unit, L2 cache control, green contexts, virtual memory, IPC, multi-GPU, driver API, math functions, device-callable APIs, compute capabilities, C++ language support, deployment, performance optimization, and CUDA profiler tools (Visual Profiler, nvprof, NVTX).
version	13.2

CUDA C++ - Parallel Computing Platform & Programming Model

Overview

CUDA (Compute Unified Device Architecture) is NVIDIA's parallel computing platform and programming model. It enables dramatic increases in computing performance by harnessing the power of GPUs for general-purpose processing (GPGPU). CUDA C++ extends C++ with GPU-specific keywords, built-in variables, and runtime APIs.

Supported Hardware: NVIDIA GPUs with Compute Capability 5.0+ (full feature set requires CC 7.0+)

Supported Platforms: Linux, Windows, WSL

CUDA Toolkit Version: 13.2 (March 2026)

Key Architecture Concepts

• Host: CPU + system memory
• Device: GPU + device memory
• Kernel: Function launched for parallel GPU execution
• Thread Block: Group of threads executing on a single SM, can cooperate via shared memory
• Grid: Collection of thread blocks executing a kernel
• Warp: Group of 32 threads executing in SIMT fashion
• SM (Streaming Multiprocessor): GPU compute unit with registers, shared memory, and functional units
• Thread Block Cluster (CC 9.0+): Group of thread blocks on a single GPC with distributed shared memory

Quick Reference

Minimal Kernel Example

__global__ void vecAdd(float* A, float* B, float* C, int N) {
    int idx = threadIdx.x + blockDim.x * blockIdx.x;
    if (idx < N) C[idx] = A[idx] + B[idx];
}

// Launch
int threads = 256;
int blocks = (N + threads - 1) / threads;
vecAdd<<<blocks, threads>>>(d_A, d_B, d_C, N);
cudaDeviceSynchronize();

Memory Allocation

// Unified Memory
float* data;
cudaMallocManaged(&data, N * sizeof(float));
cudaFree(data);

// Explicit Device Memory
cudaMalloc(&d_data, N * sizeof(float));
cudaMemcpy(d_data, h_data, N * sizeof(float), cudaMemcpyHostToDevice);
cudaFree(d_data);

// Stream-Ordered (CUDA 11.0+)
cudaMallocAsync(&d_data, N * sizeof(float), stream);
cudaFreeAsync(d_data, stream);

GPU Timing

cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, 0);
kernel<<<grid, block>>>(...);
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
float ms;
cudaEventElapsedTime(&ms, start, stop);

Error Checking Macro

#define CUDA_CHECK(expr) do { \
    cudaError_t err = expr; \
    if (err != cudaSuccess) { \
        fprintf(stderr, "CUDA Error: %s:%d: %s\n", \
            __FILE__, __LINE__, cudaGetErrorString(err)); \
        exit(EXIT_FAILURE); \
    } \
} while(0)

Compilation

nvcc -arch=sm_80 -O2 mykernel.cu -o myapp
nvcc -gencode=arch=compute_70,code=sm_70 \
     -gencode=arch=compute_80,code=sm_80 \
     -gencode=arch=compute_90,code=sm_90 \
     mykernel.cu -o myapp

Memory Hierarchy at a Glance

Memory	Location	Cached	Access	Scope	Lifetime
Register	On-chip (SM)	N/A	R/W	1 thread	Kernel
Shared	On-chip (SM)	N/A	R/W	Thread block	Kernel
L1 Cache	On-chip (SM)	Yes	R/W	Thread block	Kernel
L2 Cache	On-chip (GPU)	Yes	R/W	All threads	Kernel
Local	Off-chip (DRAM)	L1+L2	R/W	1 thread	Kernel
Global	Off-chip (DRAM)	L1+L2	R/W	All threads + host	Application
Constant	Off-chip (DRAM)	Yes	R	All threads + host	Application
Texture	Off-chip (DRAM)	Yes	R	All threads + host	Application

Code Examples

Official CUDA samples from NVIDIA/cuda-samples, organized by category:

0_Introduction (48 samples)

Basic CUDA concepts and runtime API usage. Key samples:

• vectorAdd - Basic kernel launch and memory operations
• matrixMul - Matrix multiplication with shared memory tiling
• simpleCooperativeGroups - Cooperative groups fundamentals
• simpleStreams - CUDA streams for concurrent execution
• simpleMultiGPU - Multi-GPU programming basics
• simpleCudaGraphs (in 3_CUDA_Features) - CUDA Graph capture and execution

2_Concepts_and_Techniques (34 samples)

Parallel algorithms and optimization techniques. Key samples:

• reduction - Parallel reduction with shared memory
• scan - Parallel prefix sum (scan)
• shfl_scan - Warp-level shuffle operations for scan
• histogram - Histogram computation with atomics
• sortingNetworks - Sorting network implementations
• radixSortThrust - Radix sort using Thrust

3_CUDA_Features (26 samples)

Advanced CUDA features. Key samples:

• simpleCudaGraphs - CUDA Graph API usage
• jacobiCudaGraphs - Iterative solver with CUDA Graphs
• cudaTensorCoreGemm - Tensor Core GEMM (FP16)
• bf16TensorCoreGemm - BFloat16 Tensor Core GEMM
• tf32TensorCoreGemm - TF32 Tensor Core GEMM
• globalToShmemAsyncCopy - Async copy to shared memory (TMA)
• cdpAdvancedQuicksort - CUDA Dynamic Parallelism
• graphMemoryNodes - Graph memory allocation nodes

4_CUDA_Libraries (36 samples)

CUDA library integration. Key samples:

• simpleCUBLAS - cuBLAS matrix operations
• matrixMulCUBLAS - Matrix multiplication via cuBLAS
• simpleCUFFT - cuFFT Fast Fourier Transform
• conjugateGradient - Iterative solver with cuSPARSE/cuBLAS

6_Performance (7 samples)

Performance optimization techniques:

• transpose - Optimized matrix transpose (coalesced access patterns)
• alignedTypes - Memory alignment for bandwidth
• UnifiedMemoryPerf - Unified Memory performance analysis
• cudaGraphsPerfScaling - CUDA Graphs performance scaling

Full source available in the examples/ directory with 206 samples total.

High-Priority Best Practices

1. Profile first - Use Nsight Compute/Systems to identify bottlenecks
2. Coalesce global memory accesses - Consecutive threads access consecutive addresses
3. Minimize host-device transfers - Keep data on device as long as possible
4. Use shared memory - Avoid redundant global memory loads
5. Avoid warp divergence - Keep threads in a warp on the same control flow path
6. Use effective bandwidth as a performance metric
7. Overlap computation with transfers using streams and async operations

Documentation Structure

Core Programming Model

• 01-introduction-and-programming-model - CUDA overview, GPU hardware model, threads/blocks/grids/warps, SIMT, compute capability
• 02-getting-started - CUDA C++ basics, kernels, NVCC compilation, memory, error handling, device/host functions
• 03-memory-management - Unified memory, explicit memory, page-locked memory, memory spaces, managed memory paradigms

Asynchronous Execution

• 04-asynchronous-execution - CUDA streams, events, callbacks, stream ordering, default stream, priorities
• 05-cuda-graphs - Graph structure, building (API + stream capture), instantiation, execution, updating, conditional nodes, memory nodes, device graph launch

Memory & Synchronization

• 06-stream-ordered-memory-allocator - cudaMallocAsync, cudaFreeAsync, memory pools, IPC pools
• 07-cooperative-groups - CG handles, thread_block, cluster_group, grid_group, tiled_partition, collective operations
• 08-advanced-synchronization - Async barriers, pipelines, programmatic dependent launch, memory sync domains

Data Movement

• 09-async-data-copies - LDGSTS, TMA (1D + multi-dim), STAS, tensor maps, bank swizzling, work stealing
• 10-l2-cache-control - L2 persistence, access policy windows, hit ratio tuning

Execution Contexts

• 11-green-contexts - SM partitioning, resource descriptors, work queues, execution context streams
• 12-virtual-memory-and-ipc - VM management, unicast/multicast sharing, IPC, compressible memory
• 13-multi-gpu-programming - Device enumeration, P2P transfers, peer access, consistency

APIs & Language

• 14-cuda-driver-api - Context, module, kernel execution, runtime/driver interop
• 15-compute-capabilities - CC specs, features, device/SM/memory info per CC, tensor core types
• 16-cpp-language-support - C++11/14/17/20 features, restrictions, libcu++, lambda expressions
• 17-language-extensions - Execution/memory space specifiers, built-in types, atomics, warp functions, vector types, DPX
• 18-math-functions - Standard math, intrinsic functions, non-standard functions, half precision math
• 19-device-callable-apis - Memory barrier primitives, pipeline primitives, CG API reference, device runtime (CDP)

Memory Model & Execution

• 20-cuda-memory-model - Thread scopes, atomicity, data races, forward progress, execution model
• 21-environment-variables - All CUDA env vars: device enumeration, JIT, execution, module loading, error log

Interoperability

• 22-interop-and-external-resources - Vulkan, Direct3D, OpenGL, NVSCI interop, external semaphores
• 23-driver-entry-point-access - cuGetProcAddress, versioning, typedefs, per-thread default stream

Deployment

• 24-deployment-and-compatibility - Building for compatibility, runtime redistribution, nvidia-smi, NVML, lazy loading

Best Practices Guide

• 25-best-practices-overview - APOD cycle, profiling, assessing applications, verification, debugging
• 26-memory-optimizations - Host-device transfers, coalescing, shared memory optimization, L2 cache, allocation
• 27-execution-configuration - Occupancy calculation, thread/block heuristics, concurrent kernels, multiple contexts
• 28-instruction-optimization - Arithmetic throughput table, control flow, math libraries, compiler flags, nvcc switches

Extended Features

• 29-extended-gpu-memory - EGM, NVLink-C2C, NUMA considerations
• 30-lazy-loading-and-error-log - Module lazy loading, CUDA error log management

Profiling Tools

• 31-cuda-profiler - Visual Profiler (nvvp), nvprof CLI, NVTX API, focused profiling, remote profiling, MPI/MPS profiling, dependency analysis, metrics reference, warp state analysis, migration to Nsight tools