Home › CUDA / GPU › CUDA C++ kernels Cursor rules
CUDA C++ kernels Cursor rules
Cursor rules for CUDA C++: error checking every API call, coalesced memory access, launch configuration, streams, and correctness tests before tuning.
CursorCUDA / GPU
How to use
Save this as .cursor/rules/cuda.mdc in your repository root, commit it, and your agent will pick it up automatically. Edit the commands to match your project.
Note: Review and test before use. Provided as is, without warranty. See Terms.
.cursor/rules/cuda.mdcDownload ---
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.