02 — C++ for CUDA: The Pragmatic Subset

“You do not need to be a C++ expert. You need to be a C++ reader who can write kernels without leaking, crashing, or being scared of template<typename T>.”

The Trap to Avoid

C++ is the largest programming language in production. If you try to learn all of it before writing CUDA, you will spend six months on move semantics and never write a kernel. The trap is symmetrical to the CUTLASS trap in Phase 2: premature depth in the tools.

What you actually need, in priority order:

  1. Pointer/reference/value distinctions. Because CUDA kernels take pointers, and misunderstanding what T& vs T* vs const T& means at the ABI level is how you get memory corruption.

  2. RAII. Because CUDA resources (device pointers, streams, events, cuBLAS handles) leak silently. RAII wrappers save you.

  3. const correctness. Because it’s how you communicate intent to the compiler and to CUTLASS’s template machinery.

  4. Basic templates. Because CUDA kernels, cuBLAS, and CUTLASS are all templated.

  5. CMake at survival level. Because that’s how you build CUDA projects.

  6. Move semantics at working level. Because you’ll see std::move in modern libraries and need to not misuse it.

  7. Iterators and the STL basics. Because C++ code is unreadable otherwise.

That’s it. You do not need to learn: variadic templates from scratch (you’ll pick them up from context), template metaprogramming (CUTLASS forces this later, and it’s better learned in context), custom allocators in depth, exception hierarchies, or the entire Boost universe.

Canonical Resource

Bjarne Stroustrup, A Tour of C++ (3rd edition, 2022).

  • 240 pages. Written by the language’s designer as a fast tour for people who already program.

  • Chapters to read: 1–8 (basics through classes), 13 (utilities: unique_ptr, shared_ptr, optional, span), 6 (essential operations: RAII, move), 7 (templates, at reading level).

  • Skip for now: chapters on ranges, coroutines, modules, concepts (you’ll pick these up from context later).

Buy the book. It is not free but it is short. The next-best free alternatives are Herb Sutter’s talks on YouTube (“Modern C++: What You Need to Know” — CppCon) and https://learncpp.com (thorough, longer than you need, but the reference is excellent when you have a specific question).

Modern C++ Idioms You Must Recognize

RAII (Resource Acquisition Is Initialization)

Every resource is owned by an object whose destructor releases it. This is how modern C++ avoids the “forgot to free” bug. For CUDA:

class DeviceBuffer {
    void* ptr_ = nullptr;
    size_t bytes_ = 0;
public:
    explicit DeviceBuffer(size_t bytes) : bytes_(bytes) {
        cudaError_t err = cudaMalloc(&ptr_, bytes);
        if (err != cudaSuccess) throw std::runtime_error(cudaGetErrorString(err));
    }
    ~DeviceBuffer() { if (ptr_) cudaFree(ptr_); }  // No-throw destructor.
    // Rule of five: delete copy, define move.
    DeviceBuffer(const DeviceBuffer&) = delete;
    DeviceBuffer& operator=(const DeviceBuffer&) = delete;
    DeviceBuffer(DeviceBuffer&& o) noexcept : ptr_(o.ptr_), bytes_(o.bytes_) { o.ptr_ = nullptr; }
    DeviceBuffer& operator=(DeviceBuffer&& o) noexcept {
        if (this != &o) { cudaFree(ptr_); ptr_ = o.ptr_; bytes_ = o.bytes_; o.ptr_ = nullptr; }
        return *this;
    }
    void* data() const { return ptr_; }
    size_t bytes() const { return bytes_; }
};

Internalize the Rule of Five: if you write a destructor, you probably need copy-ctor, copy-assign, move-ctor, move-assign (or = delete the ones you don’t want). This is not optional.

unique_ptr / shared_ptr

auto buf = std::make_unique<float[]>(1024);  // Heap-allocated, RAII, no leak.

Use unique_ptr by default. shared_ptr only when you genuinely need shared ownership — it has a reference-count cost and hides ownership.

const correctness

const T& means “I promise not to modify this.” Use it aggressively on function parameters. It documents intent, enables optimizations, and lets you pass rvalues (temporaries) without a copy.

constexpr

Compile-time evaluation. Matters more than you’d think in CUDA — tile sizes, block dimensions, loop unroll factors are constexpr values that the compiler inlines and optimizes around.

Templates at reading level

You should be able to read this without flinching:

template <typename T, int BLOCK_M, int BLOCK_N, int BLOCK_K>
__global__ void sgemm_tiled(const T* A, const T* B, T* C, int M, int N, int K) {
    __shared__ T As[BLOCK_M][BLOCK_K];
    __shared__ T Bs[BLOCK_K][BLOCK_N];
    // ... you get the idea
}

You don’t need to write template metaprograms. You need to read them.

Move semantics (working level)

std::move(x) casts x to an rvalue reference so it can be moved from. It doesn’t move anything by itself. Rule of thumb: use std::move when passing a value you own into a function that takes by value or by &&, and you no longer need it afterwards.

CMake Survival Guide

You will not master CMake. Nobody has. You will learn the incantations that work for CUDA projects.

Minimum viable CMakeLists.txt for a CUDA project:

cmake_minimum_required(VERSION 3.24)
project(my_kernels LANGUAGES CXX CUDA)

set(CMAKE_CXX_STANDARD 17)
set(CMAKE_CUDA_STANDARD 17)
set(CMAKE_CUDA_ARCHITECTURES 80 89 90)  # A100=80, RTX 4090=89, H100=90; adjust to your GPU.

add_executable(sgemm sgemm.cu main.cpp)
target_compile_options(sgemm PRIVATE
    $<$<COMPILE_LANGUAGE:CUDA>:-lineinfo -O3 --use_fast_math>
)
target_link_libraries(sgemm PRIVATE CUDA::cudart CUDA::cublas)

Commands you’ll run 100× a week:

cmake -B build -S . -DCMAKE_BUILD_TYPE=Release
cmake --build build -j
./build/sgemm

Use -G Ninja for faster incremental builds. Learn add_library, target_include_directories, target_link_libraries and stop there — anything more, ask the docs.

Modern CMake reading: https://cliutils.gitlab.io/modern-cmake/ (free HTML book, well-maintained).

The Reading Drill: One llama.cpp File

The fastest way to become fluent in the C++ dialect of ML systems is to read one real file until nothing surprises you.

Recommended file: ggml-quants.c (or one of its split-out headers in the current tree) in llama.cpp — the CPU quantization kernels. https://github.com/ggerganov/llama.cpp

(Note: the repo has been renamed to https://github.com/ggml-org/llama.cpp; both URLs redirect. It has ~90k+ stars and daily commits — anything but suspicious.)

Alternative: common/common.cpp for lighter warmup, or src/llama.cpp for the deep end.

Protocol:

  1. Open the file.

  2. Read top to bottom.

  3. Every line you don’t understand → note it in a questions.md.

  4. Answer each question via docs, cppreference (cppreference.com is the reference — bookmark it), or code search.

  5. Do not move on until questions.md is empty.

Budget: 6–10 hours for one meaty file. This exercise trades “I’ve read a C++ book” for “I can navigate real C++ code”, which is the actually valuable skill.

Common CUDA-Adjacent C++ Gotchas

Both for future reference and to inoculate you against them:

  1. Kernel launch is asynchronous. kernel<<<>>>() returns immediately. Time it wrong and you’ll “measure” 0.001 ms for a kernel that takes 50 ms. Always cudaDeviceSynchronize() or use cudaEvent_t timers.

  2. Host and device pointers are not interchangeable. Dereferencing a cudaMalloc’d pointer from CPU code is undefined behavior (segfault if you’re lucky).

  3. Struct-of-arrays vs array-of-structs. SoA is almost always faster on GPUs for coalescing reasons. This is why frameworks split KV cache into separate K and V tensors instead of one packed struct.

  4. __restrict__ on kernel pointer parameters tells the compiler pointers don’t alias, enabling optimizations. Use it habitually.

  5. Compile with -lineinfo (not -G, which is full debug and 10× slower). -lineinfo gives Nsight Compute source-line-to-SASS mapping without disabling optimizations.

Exit Deliverable

  1. A small CMake+CUDA project skeleton committed to your portfolio: builds a hello-world kernel and a DeviceBuffer RAII class, with a smoke test.

  2. llama_cpp_annotated.md: one llama.cpp file (your choice) with your line-by-line notes and answered questions.

  3. Your own raii_wrappers.h header with RAII wrappers for cudaMalloc, cudaStream_t, cudaEvent_t, and cublasHandle_t. You will reuse this in every project in Phases 2–4.