CUDA Samples Skill
Curated index of ~50 official NVIDIA CUDA Samples organized by common CUDA programming tasks. Each entry links to the canonical source, describes the demonstrated pattern, lists architecture requirements, and provides a key code snippet.
Usage: Search this file by pattern name (e.g., "reduction", "GEMM", "graph capture") or API name (e.g., "cublasSgemm", "cudaMallocAsync"). When a match is found, follow the GitHub permalink for the full implementation.
Source: NVIDIA/cuda-samples (CUDA Toolkit 13.2), 175 .cu files, ~57K lines of CUDA C++.
Quick Reference Table
| Task |
Category |
C++ Sample |
Python Sample |
| First CUDA kernel |
Getting Started |
vectorAdd |
vectorAdd |
| Device properties query |
Utilities |
deviceQuery |
deviceQuery |
| GPU timing with events |
Getting Started |
asyncAPI |
— |
| Multi-stream overlap |
Streams & Async |
simpleStreams |
streamingCopyComputeOverlap |
| Stream-ordered allocation |
Memory |
streamOrderedAllocation |
— |
| Unified Memory perf |
Memory |
UnifiedMemoryPerf |
blurImageUnifiedMemory |
| cuda.core memory model |
Memory |
— |
memoryResources |
| Thread/block indexing patterns |
Getting Started |
— |
blockwiseSum |
| IPC between processes |
Multi-GPU |
simpleIPC |
ipcMemoryPool |
| P2P between GPUs |
Multi-GPU |
simpleP2P |
simpleP2P |
| Parallel reduction |
Kernel Patterns |
reduction |
reduction, parallelReduction |
| Multi-block CG reduction |
Kernel Patterns |
reductionMultiBlockCG |
reductionMultiBlockCG |
| Warp shuffle scan |
Kernel Patterns |
shfl_scan |
— |
| Prefix sum (library) |
Kernel Patterns |
— |
prefixSum |
| Sorting networks |
Kernel Patterns |
sortingNetworks |
— |
| Radix sort (Thrust) |
Kernel Patterns |
radixSortThrust |
— |
| Histogram with atomics |
Kernel Patterns |
histogram |
parallelHistogram |
| FP16 Tensor Core GEMM |
Tensor Core |
cudaTensorCoreGemm |
— |
| BF16 Tensor Core GEMM |
Tensor Core |
bf16TensorCoreGemm |
— |
| TF32 Tensor Core GEMM |
Tensor Core |
tf32TensorCoreGemm |
— |
| Double-precision TC GEMM |
Tensor Core |
dmmaTensorCoreGemm |
— |
| Async copy global→shared (pipeline) |
Tensor Core |
globalToShmemAsyncCopy |
— |
| TMA (Hopper+) |
Tensor Core |
— |
tmaTensorMap |
| CUDA Graphs basics |
Graphs |
simpleCudaGraphs |
cudaGraphs |
| Graph with updates |
Graphs |
jacobiCudaGraphs |
— |
| Graph memory nodes |
Graphs |
graphMemoryNodes |
— |
| Conditional graph nodes |
Graphs |
graphConditionalNodes |
— |
| cuBLAS GEMM |
Libraries |
simpleCUBLAS |
— |
| Batched cuBLAS GEMM |
Libraries |
batchCUBLAS |
— |
| cuBLAS with CUDA Graphs |
Libraries |
conjugateGradientCudaGraphs |
— |
| Conjugate Gradient (cuBLAS+cuSPARSE) |
Libraries |
conjugateGradient |
— |
| cuFFT convolution |
Libraries |
simpleCUFFT |
fftSignalAnalysis |
| Matrix transpose |
Performance |
transpose |
— |
| Block size tuning |
Performance |
— |
launchConfigTuning |
| Warp-aggregated atomics |
Performance |
warpAggregatedAtomicsCG |
— |
| PyTorch custom kernel |
Framework |
— |
customPyTorchKernel |
| TensorFlow custom kernel |
Framework |
— |
customTensorFlowKernel |
| Multi-GPU gradient average |
Distributed |
— |
multiGPUGradientAverage |
| cuSOLVER linear solver |
Libraries |
cuSolverDn_LinearSolver |
— |
| JIT LTO linking |
Advanced |
— |
jitLtoLinking |
| NVTX profiling markers |
Profiling |
— |
kernelNsysProfile |
| Green context (SM partitioning) |
Advanced |
— |
greenContext |
| Image processing (NPP) |
Libraries |
histEqualizationNPP |
— |
The table above is a fast lookup index. Detailed code snippets and patterns for the most important samples are in the references/ directory, organized by topic (some entries are index-only without detailed snippets):
| Category |
File |
Contents |
| Getting Started & Memory |
samples-getting-started.md |
vectorAdd (C++/Python), asyncAPI, streamOrderedAllocation, UnifiedMemoryPerf, memoryResources |
| Streams & Async |
samples-streams-async.md |
simpleStreams, streamingCopyComputeOverlap, simpleIPC, simpleP2P |
| Reduction, Scan & Sort |
samples-reduction-scan-sort.md |
reduction (C++/Python), reductionMultiBlockCG, shfl_scan, radixSortThrust, parallelReduction |
| Tensor Core GEMM |
samples-tensor-core-gemm.md |
cudaTensorCoreGemm, bf16TensorCoreGemm, tf32TensorCoreGemm, dmmaTensorCoreGemm, globalToShmemAsyncCopy, tmaTensorMap |
| CUDA Graphs |
samples-cuda-graphs.md |
simpleCudaGraphs, cudaGraphs (Python), jacobiCudaGraphs, graphMemoryNodes |
| CUDA Libraries |
samples-cuda-libraries.md |
simpleCUBLAS, batchCUBLAS, conjugateGradient, simpleCUFFT, histEqualizationNPP |
| Performance Optimization |
samples-performance.md |
transpose, launchConfigTuning, warpAggregatedAtomicsCG, kernelNsysProfile |
| Framework Interop |
samples-framework-interop.md |
customPyTorchKernel, customTensorFlowKernel |
| Multi-GPU & Distributed |
samples-multi-gpu.md |
simpleP2P (Python), ipcMemoryPool, multiGPUGradientAverage |
| Advanced Topics |
samples-advanced-topics.md |
blockwiseSum, prefixSum, jitLtoLinking, greenContext |
Optimization Strategy → Sample Mapping
Map optimization strategies to relevant samples:
| Bottleneck Type |
Optimization Strategy |
Reference Sample(s) |
| DRAM_MEMORY_BOUND |
Block tiling, vectorized load, prefetching |
transpose (merge→coalesced), cudaTensorCoreGemm (int4 loads) |
| L1_PRESSURE_BOUND |
Shared memory padding, data transpose, fragment caching |
transpose (no bank conflict variant), reduction (shared memory tree) |
| LATENCY_BOUND |
Double buffering, ILP, loop unrolling |
bf16TensorCoreGemm (pipeline), reduction (multi-element per thread) |
| COMPUTE_BOUND |
FMA, FP16/TF32, Tensor Cores |
cudaTensorCoreGemm, tf32TensorCoreGemm, bf16TensorCoreGemm |
| OCCUPANCY_BOUND |
Adjust block size, __launch_bounds__, reduce smem |
launchConfigTuning, SHARED_MEMORY_LIMIT_64K in tensor core samples |
| MIXED_BOUND |
Narrower NCU section sets; target top metric first |
transpose (coarse-grain) |
Architecture Compatibility
| Feature |
Min SM |
Samples |
| Basic CUDA |
All |
vectorAdd, asyncAPI, simpleStreams, simpleCUBLAS, simpleCUFFT |
Warp shuffle (__shfl_*) |
3.0 |
shfl_scan, reduction (variants 5-6) |
| Cooperative groups (block-level) |
5.2 |
reduction (CG variants), warpAggregatedAtomicsCG |
| Cooperative launch (grid sync) |
6.0 |
reductionMultiBlockCG |
| WMMA (Tensor Cores) |
7.0 |
cudaTensorCoreGemm, dmmaTensorCoreGemm |
Async copy (memcpy_async) |
8.0 |
bf16TensorCoreGemm, globalToShmemAsyncCopy |
| TF32 Tensor Cores |
8.0 |
tf32TensorCoreGemm |
__reduce_add_sync |
8.0 |
reduction (variant 7) |
| TMA (Tensor Memory Accelerator) |
9.0 |
tmaTensorMap (Python) |
| Green context (SM partitioning) |
9.0 |
greenContext (Python) |
Search Tips
# Find samples by API usage (search references/ directory)
grep -r "cublasSgemm" skills/cuda-samples/references/
grep -r "cudaStreamBeginCapture\|cudaGraph" skills/cuda-samples/references/
grep -r "wmma::mma_sync\|__pipeline_memcpy_async" skills/cuda-samples/references/
# Find samples by pattern
grep -r "reduction\|scan\|histogram" skills/cuda-samples/references/
grep -r "PyTorch\|TensorFlow\|framework" skills/cuda-samples/references/
# Find samples by architecture requirement
grep -r "SM 8.0\|SM 9.0\|Hopper\|Ampere" skills/cuda-samples/references/
# Find samples by name in the quick reference table
grep -i "vectorAdd\|simpleCUBLAS\|transpose" skills/cuda-samples/SKILL.md
# Clone full repo for compilation and detailed study
git clone --depth 1 https://github.com/NVIDIA/cuda-samples.git
1---2name: cuda-samples3description: Curated index of NVIDIA official CUDA Samples with code patterns, GitHub permalinks, and key snippets. Use when looking for working CUDA code examples, GPU kernel patterns (reduction, scan, GEMM, transpose, histogram), Tensor Core WMMA usage, CUDA Graphs API, cuBLAS/cuFFT/cuSPARSE library integration, stream/event async patterns, multi-GPU P2P/IPC, or framework interop (PyTorch/TensorFlow custom kernels). Covers 50+ curated samples organized by topic. Triggers on CUDA sample, code example, kernel pattern, GEMM example, reduction example, CUDA graph example, cuBLAS example, Tensor Core WMMA, how to write CUDA kernel, CUDA code reference.4---56# CUDA Samples Skill78Curated index of ~50 official [NVIDIA CUDA Samples](https://github.com/NVIDIA/cuda-samples) organized by common CUDA programming tasks. Each entry links to the canonical source, describes the demonstrated pattern, lists architecture requirements, and provides a key code snippet.910**Usage**: Search this file by pattern name (e.g., "reduction", "GEMM", "graph capture") or API name (e.g., "cublasSgemm", "cudaMallocAsync"). When a match is found, follow the GitHub permalink for the full implementation.1112**Source**: NVIDIA/cuda-samples (CUDA Toolkit 13.2), 175 `.cu` files, ~57K lines of CUDA C++.1314---1516## Quick Reference Table1718| Task | Category | C++ Sample | Python Sample |19| ------------------------------------ | --------------- | ----------------------------- | -------------------------------- |20| First CUDA kernel | Getting Started | `vectorAdd` | `vectorAdd` |21| Device properties query | Utilities | `deviceQuery` | `deviceQuery` |22| GPU timing with events | Getting Started | `asyncAPI` | — |23| Multi-stream overlap | Streams & Async | `simpleStreams` | `streamingCopyComputeOverlap` |24| Stream-ordered allocation | Memory | `streamOrderedAllocation` | — |25| Unified Memory perf | Memory | `UnifiedMemoryPerf` | `blurImageUnifiedMemory` |26| cuda.core memory model | Memory | — | `memoryResources` |27| Thread/block indexing patterns | Getting Started | — | `blockwiseSum` |28| IPC between processes | Multi-GPU | `simpleIPC` | `ipcMemoryPool` |29| P2P between GPUs | Multi-GPU | `simpleP2P` | `simpleP2P` |30| Parallel reduction | Kernel Patterns | `reduction` | `reduction`, `parallelReduction` |31| Multi-block CG reduction | Kernel Patterns | `reductionMultiBlockCG` | `reductionMultiBlockCG` |32| Warp shuffle scan | Kernel Patterns | `shfl_scan` | — |33| Prefix sum (library) | Kernel Patterns | — | `prefixSum` |34| Sorting networks | Kernel Patterns | `sortingNetworks` | — |35| Radix sort (Thrust) | Kernel Patterns | `radixSortThrust` | — |36| Histogram with atomics | Kernel Patterns | `histogram` | `parallelHistogram` |37| FP16 Tensor Core GEMM | Tensor Core | `cudaTensorCoreGemm` | — |38| BF16 Tensor Core GEMM | Tensor Core | `bf16TensorCoreGemm` | — |39| TF32 Tensor Core GEMM | Tensor Core | `tf32TensorCoreGemm` | — |40| Double-precision TC GEMM | Tensor Core | `dmmaTensorCoreGemm` | — |41| Async copy global→shared (pipeline) | Tensor Core | `globalToShmemAsyncCopy` | — |42| TMA (Hopper+) | Tensor Core | — | `tmaTensorMap` |43| CUDA Graphs basics | Graphs | `simpleCudaGraphs` | `cudaGraphs` |44| Graph with updates | Graphs | `jacobiCudaGraphs` | — |45| Graph memory nodes | Graphs | `graphMemoryNodes` | — |46| Conditional graph nodes | Graphs | `graphConditionalNodes` | — |47| cuBLAS GEMM | Libraries | `simpleCUBLAS` | — |48| Batched cuBLAS GEMM | Libraries | `batchCUBLAS` | — |49| cuBLAS with CUDA Graphs | Libraries | `conjugateGradientCudaGraphs` | — |50| Conjugate Gradient (cuBLAS+cuSPARSE) | Libraries | `conjugateGradient` | — |51| cuFFT convolution | Libraries | `simpleCUFFT` | `fftSignalAnalysis` |52| Matrix transpose | Performance | `transpose` | — |53| Block size tuning | Performance | — | `launchConfigTuning` |54| Warp-aggregated atomics | Performance | `warpAggregatedAtomicsCG` | — |55| PyTorch custom kernel | Framework | — | `customPyTorchKernel` |56| TensorFlow custom kernel | Framework | — | `customTensorFlowKernel` |57| Multi-GPU gradient average | Distributed | — | `multiGPUGradientAverage` |58| cuSOLVER linear solver | Libraries | `cuSolverDn_LinearSolver` | — |59| JIT LTO linking | Advanced | — | `jitLtoLinking` |60| NVTX profiling markers | Profiling | — | `kernelNsysProfile` |61| Green context (SM partitioning) | Advanced | — | `greenContext` |62| Image processing (NPP) | Libraries | `histEqualizationNPP` | — |6364The table above is a fast lookup index. Detailed code snippets and patterns for the most important samples are in the [references/](references/) directory, organized by topic (some entries are index-only without detailed snippets):6566| Category | File | Contents |67| ------------------------ | --------------------------------------------------------------------------- | -------------------------------------------------------------------------------------------------------------------- |68| Getting Started & Memory | [samples-getting-started.md](references/samples-getting-started.md) | vectorAdd (C++/Python), asyncAPI, streamOrderedAllocation, UnifiedMemoryPerf, memoryResources |69| Streams & Async | [samples-streams-async.md](references/samples-streams-async.md) | simpleStreams, streamingCopyComputeOverlap, simpleIPC, simpleP2P |70| Reduction, Scan & Sort | [samples-reduction-scan-sort.md](references/samples-reduction-scan-sort.md) | reduction (C++/Python), reductionMultiBlockCG, shfl_scan, radixSortThrust, parallelReduction |71| Tensor Core GEMM | [samples-tensor-core-gemm.md](references/samples-tensor-core-gemm.md) | cudaTensorCoreGemm, bf16TensorCoreGemm, tf32TensorCoreGemm, dmmaTensorCoreGemm, globalToShmemAsyncCopy, tmaTensorMap |72| CUDA Graphs | [samples-cuda-graphs.md](references/samples-cuda-graphs.md) | simpleCudaGraphs, cudaGraphs (Python), jacobiCudaGraphs, graphMemoryNodes |73| CUDA Libraries | [samples-cuda-libraries.md](references/samples-cuda-libraries.md) | simpleCUBLAS, batchCUBLAS, conjugateGradient, simpleCUFFT, histEqualizationNPP |74| Performance Optimization | [samples-performance.md](references/samples-performance.md) | transpose, launchConfigTuning, warpAggregatedAtomicsCG, kernelNsysProfile |75| Framework Interop | [samples-framework-interop.md](references/samples-framework-interop.md) | customPyTorchKernel, customTensorFlowKernel |76| Multi-GPU & Distributed | [samples-multi-gpu.md](references/samples-multi-gpu.md) | simpleP2P (Python), ipcMemoryPool, multiGPUGradientAverage |77| Advanced Topics | [samples-advanced-topics.md](references/samples-advanced-topics.md) | blockwiseSum, prefixSum, jitLtoLinking, greenContext |7879---8081## Optimization Strategy → Sample Mapping8283Map optimization strategies to relevant samples:8485| Bottleneck Type | Optimization Strategy | Reference Sample(s) |86| ----------------- | ------------------------------------------------------- | ------------------------------------------------------------------------ |87| DRAM_MEMORY_BOUND | Block tiling, vectorized load, prefetching | `transpose` (merge→coalesced), `cudaTensorCoreGemm` (int4 loads) |88| L1_PRESSURE_BOUND | Shared memory padding, data transpose, fragment caching | `transpose` (no bank conflict variant), `reduction` (shared memory tree) |89| LATENCY_BOUND | Double buffering, ILP, loop unrolling | `bf16TensorCoreGemm` (pipeline), `reduction` (multi-element per thread) |90| COMPUTE_BOUND | FMA, FP16/TF32, Tensor Cores | `cudaTensorCoreGemm`, `tf32TensorCoreGemm`, `bf16TensorCoreGemm` |91| OCCUPANCY_BOUND | Adjust block size, `__launch_bounds__`, reduce smem | `launchConfigTuning`, `SHARED_MEMORY_LIMIT_64K` in tensor core samples |92| MIXED_BOUND | Narrower NCU section sets; target top metric first | `transpose` (coarse-grain) |9394---9596## Architecture Compatibility9798| Feature | Min SM | Samples |99| -------------------------------- | ------ | ----------------------------------------------------------------------- |100| Basic CUDA | All | `vectorAdd`, `asyncAPI`, `simpleStreams`, `simpleCUBLAS`, `simpleCUFFT` |101| Warp shuffle (`__shfl_*`) | 3.0 | `shfl_scan`, `reduction` (variants 5-6) |102| Cooperative groups (block-level) | 5.2 | `reduction` (CG variants), `warpAggregatedAtomicsCG` |103| Cooperative launch (grid sync) | 6.0 | `reductionMultiBlockCG` |104| WMMA (Tensor Cores) | 7.0 | `cudaTensorCoreGemm`, `dmmaTensorCoreGemm` |105| Async copy (`memcpy_async`) | 8.0 | `bf16TensorCoreGemm`, `globalToShmemAsyncCopy` |106| TF32 Tensor Cores | 8.0 | `tf32TensorCoreGemm` |107| `__reduce_add_sync` | 8.0 | `reduction` (variant 7) |108| TMA (Tensor Memory Accelerator) | 9.0 | `tmaTensorMap` (Python) |109| Green context (SM partitioning) | 9.0 | `greenContext` (Python) |110111---112113## Search Tips114115```bash116# Find samples by API usage (search references/ directory)117grep -r "cublasSgemm" skills/cuda-samples/references/118grep -r "cudaStreamBeginCapture\|cudaGraph" skills/cuda-samples/references/119grep -r "wmma::mma_sync\|__pipeline_memcpy_async" skills/cuda-samples/references/120121# Find samples by pattern122grep -r "reduction\|scan\|histogram" skills/cuda-samples/references/123grep -r "PyTorch\|TensorFlow\|framework" skills/cuda-samples/references/124125# Find samples by architecture requirement126grep -r "SM 8.0\|SM 9.0\|Hopper\|Ampere" skills/cuda-samples/references/127128# Find samples by name in the quick reference table129grep -i "vectorAdd\|simpleCUBLAS\|transpose" skills/cuda-samples/SKILL.md130131# Clone full repo for compilation and detailed study132git clone --depth 1 https://github.com/NVIDIA/cuda-samples.git133```
wenyi-li