---
description: CUDA kernel and host code conventions
globs: ["**/*.cu", "**/*.cuh", "**/CMakeLists.txt"]
alwaysApply: false
---

# CUDA C++

- Build with CMake's CUDA language support (`project(x LANGUAGES CXX CUDA)`), set `CMAKE_CUDA_ARCHITECTURES` explicitly from the target GPUs, and match the toolkit version in the README/CI.
- Wrap every runtime API call and check kernel launches: `CUDA_CHECK(cudaMalloc(...))`, and after launches `CUDA_CHECK(cudaGetLastError())` (plus a sync in debug builds). No unchecked calls.
- Prefer RAII wrappers or `thrust`/CUB containers over raw `cudaMalloc`/`cudaFree` pairs; never leak device memory on error paths.
- Before writing a custom kernel, check whether cuBLAS, cuDNN, CUB, Thrust or CUTLASS already does it.
- Kernels: compute a global index and guard bounds (`if (i >= n) return;`), or use grid-stride loops. Block size is a multiple of 32 (start at 128 or 256) and is a named constant.
- Memory access: make adjacent threads read adjacent addresses (coalescing); use structure-of-arrays for hot data; use `__restrict__` and `const` on pointers.
- Shared memory tiles need `__syncthreads()` on both sides of the reuse; avoid bank conflicts with padding. Never call `__syncthreads()` inside divergent branches.
- Use streams and `cudaMemcpyAsync` with pinned host memory for overlap; don't use the legacy default stream implicitly in multi-stream code.
- Avoid atomics in hot loops when a warp/block reduction (`__shfl_down_sync`, CUB `BlockReduce`) works.
- Floating point: state tolerance in tests; don't compare GPU and CPU results with `==`.

## Process
Correctness first: a CPU reference and a test for sizes that aren't multiples of the block size. Then profile with Nsight Compute/Systems before optimizing.

## Done when
It builds without warnings, tests pass, and `compute-sanitizer` (memcheck and racecheck) is clean on the test binary.