Best Practices

39.1 Design principles

Think in data parallelism: one thread per output element is a good starting point.

Keep data on the GPU; minimize transfers; overlap the rest.

Coalesce, reuse (shared memory, registers), and avoid divergence.

Use libraries first (cuBLAS, cuDNN, CUB/Thrust, cuFFT); write custom kernels for fusion and special cases.

Profile-driven optimization (Nsight Systems → Nsight Compute → roofline).

Check errors everywhere; validate against CPU references; run compute-sanitizer in CI.

Write portable code: no hard-coded SM counts or warp assumptions; query device properties; compile for multiple architectures + PTX.

Manage resources with RAII (or Thrust containers) to avoid leaks.

Use 64-bit indexing for arrays larger than 2³¹ elements.

Document launch configurations and assumptions (block size multiples, shared memory sizes, alignment).

39.2 Common pitfalls

forgetting bounds checks (if (i < n)) always guard global accesses ignoring launch errors cudaGetLastError() after launches reading results before synchronization sync stream/device or use events __syncthreads() in divergent branches restructure so all threads reach the barrier missing __syncthreads() between shared-memory write and read accidental double precision (2.0 literals, sqrt on floats in some contexts) use 2.0f, sqrtf integer overflow in blockIdx.x * blockDim.x cast to size_t for huge arrays uncoalesced (AoS, strided) accesses SoA, transposes via shared memory tiny kernels / tiny grids fuse kernels, increase work per launch, CUDA graphs cudaMalloc/cudaFree in loops allocate once, pools, cudaMallocAsync pageable memory with async copies warp-synchronous code without __syncwarp/_sync intrinsics use explicit sync primitives (Volta+) comparing GPU and CPU floats exactly relative/absolute tolerances compiling for the wrong architecture

-arch=native or proper -gencode list

39.3 Production concerns

Version compatibility: driver must support the toolkit version (forward compatibility packages, minor-version compatibility in CUDA 11+/12+).

Containerization: NVIDIA Container Toolkit (docker run --gpus all nvcr.io/nvidia/cuda:12.6.0-devel-ubuntu22.04), NGC containers.

Monitoring: nvidia-smi, DCGM (Data Center GPU Manager), power/thermal limits, ECC errors.

Testing: unit tests with GoogleTest + CUDA, deterministic seeds, CI on GPU runners.

Licensing: CUDA Toolkit EULA allows redistribution of runtime libraries per its terms.