Library › C › CUDA

Volume C · The TensorCore.si Library

CUDA

Programming thousands of GPU cores at once.

CUDA is NVIDIA's platform and C++-based language for running code on GPUs. Every major AI framework uses it to drive Tensor Cores, and writing your own kernels is how the fastest AI code, such as FlashAttention, is made.

Prologue

Why learn CUDA?

A CPU has a few powerful cores; a GPU has thousands of simpler ones. CUDA lets you split a problem into many small pieces that all run at the same moment. Once you can think this way, you understand why AI runs on GPUs and how to make it faster.

Direct access to Tensor Cores and GPU memory, the hardware that powers modern AI.
Every framework sits on it: PyTorch, TensorFlow, JAX, vLLM and TensorRT call CUDA kernels.
Huge speed-ups: well-written kernels run tens to hundreds of times faster than CPU code on parallel problems.
Valuable skill: GPU kernel engineers are among the most sought-after people in AI.

Prologue

Set up and run your first program

You need an NVIDIA GPU, a recent NVIDIA driver and the CUDA Toolkit from developer.nvidia.com, which includes the nvcc compiler. No GPU? Free cloud notebooks such as Google Colab and Kaggle provide one.

terminal

nvidia-smi          # shows your GPU and driver
nvcc --version      # shows the CUDA compiler
nvcc hello.cu -o hello
./hello

hello.cu

#include <cstdio>

__global__ void hello() {
    printf("Hello from GPU thread %d\n", threadIdx.x);
}

int main() {
    hello<<<1, 4>>>();          // 1 block of 4 threads
    cudaDeviceSynchronize();    // wait for the GPU to finish
    return 0;
}

Tip__global__ marks a kernel: a function that runs on the GPU and is launched from the CPU with the <<<blocks, threads>>> syntax.

Chapter 1

Host, device and kernels

A CUDA program has two sides. The host is the CPU and its memory; the device is the GPU and its memory. The host prepares data, copies it to the device, launches kernels and copies results back.

Functions are marked by where they run: __global__ (runs on the GPU, called from the CPU), __device__ (runs on the GPU, called from GPU code) and __host__ (normal CPU code).

model.cu

__device__ float square(float x) { return x * x; }   // GPU helper

__global__ void square_all(float* data, int n) {      // kernel
    int i = threadIdx.x;
    if (i < n) data[i] = square(data[i]);
}

// host code launches it:  square_all<<<1, 256>>>(d_data, 256);

Chapter 2

Threads, blocks and grids

Threads are grouped into blocks (up to 1,024 threads each), and blocks form a grid. Each thread computes its own global index from blockIdx, blockDim and threadIdx, then works on its own piece of the data.

On the hardware, blocks run on streaming multiprocessors (SMs) and threads execute in groups of 32 called warps. Choose block sizes that are multiples of 32, such as 128 or 256.

threads.cu

__global__ void scale(float* x, float a, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;   // global index
    if (i < n) x[i] *= a;                            // guard the tail
}

// launch enough blocks to cover n elements:
// int threads = 256;
// int blocks = (n + threads - 1) / threads;
// scale<<<blocks, threads>>>(d_x, 2.0f, n);

Key ideaAlways check i < n. The last block usually has more threads than there is data left.

Chapter 3

GPU memory

Allocate device memory with cudaMalloc, copy with cudaMemcpy and release it with cudaFree. Unified memory (cudaMallocManaged) gives one pointer usable on both sides, which is simpler while learning.

The memory hierarchy matters for speed: registers (fastest, per thread), shared memory (per block), L2 cache, then global memory (HBM or GDDR, the largest and slowest).

memory.cu

#include <cstdio>

int main() {
    const int n = 1 << 20;                  // about 1 million floats
    size_t bytes = n * sizeof(float);
    float* h = (float*)malloc(bytes);       // host memory
    for (int i = 0; i < n; ++i) h[i] = 1.0f;

    float* d;
    cudaMalloc(&d, bytes);                               // device memory
    cudaMemcpy(d, h, bytes, cudaMemcpyHostToDevice);     // CPU -> GPU
    // ... launch kernels that use d ...
    cudaMemcpy(h, d, bytes, cudaMemcpyDeviceToHost);     // GPU -> CPU
    cudaFree(d);
    free(h);
    printf("done\n");
}

Chapter 4

Your first complete kernel: vector add

This full program adds two arrays of a million numbers on the GPU, each thread handling one element. It uses unified memory to keep the code short.

vector_add.cu

#include <cstdio>

__global__ void add(const float* a, const float* b, float* c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) c[i] = a[i] + b[i];
}

int main() {
    const int n = 1 << 20;
    float *a, *b, *c;
    cudaMallocManaged(&a, n * sizeof(float));
    cudaMallocManaged(&b, n * sizeof(float));
    cudaMallocManaged(&c, n * sizeof(float));
    for (int i = 0; i < n; ++i) { a[i] = 1.0f; b[i] = 2.0f; }

    int threads = 256, blocks = (n + threads - 1) / threads;
    add<<<blocks, threads>>>(a, b, c, n);
    cudaDeviceSynchronize();

    printf("c[0] = %.1f, c[n-1] = %.1f\n", c[0], c[n - 1]);   // 3.0, 3.0
    cudaFree(a); cudaFree(b); cudaFree(c);
}

Chapter 5

Shared memory and tiling

Global memory is slow compared with the GPU's maths units. Shared memory is a small, fast scratchpad shared by the threads in a block. Matrix multiplication loads tiles of the inputs into shared memory once and reuses them many times.

This classic tiled kernel is the stepping stone to understanding how cuBLAS and Tensor Core kernels reach high speed.

matmul_tiled.cu

#define TILE 16

__global__ void matmul(const float* A, const float* B, float* C, int N) {
    __shared__ float As[TILE][TILE];
    __shared__ float Bs[TILE][TILE];
    int row = blockIdx.y * TILE + threadIdx.y;
    int col = blockIdx.x * TILE + threadIdx.x;
    float sum = 0.0f;

    for (int t = 0; t < (N + TILE - 1) / TILE; ++t) {
        int ac = t * TILE + threadIdx.x, br = t * TILE + threadIdx.y;
        As[threadIdx.y][threadIdx.x] = (row < N && ac < N) ? A[row * N + ac] : 0.0f;
        Bs[threadIdx.y][threadIdx.x] = (br < N && col < N) ? B[br * N + col] : 0.0f;
        __syncthreads();                       // wait until the tile is loaded
        for (int k = 0; k < TILE; ++k) sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        __syncthreads();                       // wait before overwriting the tile
    }
    if (row < N && col < N) C[row * N + col] = sum;
}

// launch: dim3 threads(TILE, TILE); dim3 blocks((N + TILE - 1) / TILE, (N + TILE - 1) / TILE);
//         matmul<<<blocks, threads>>>(dA, dB, dC, N);

Chapter 6

Synchronisation and atomics

__syncthreads() makes every thread in a block wait at the same point. When many threads update the same value, use atomic operations such as atomicAdd so updates are not lost. Faster still is a reduction inside each block followed by one atomic per block.

reduce.cu

__global__ void sum(const float* x, float* total, int n) {
    __shared__ float partial[256];
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    partial[threadIdx.x] = (i < n) ? x[i] : 0.0f;
    __syncthreads();
    for (int s = blockDim.x / 2; s > 0; s >>= 1) {     // tree reduction
        if (threadIdx.x < s) partial[threadIdx.x] += partial[threadIdx.x + s];
        __syncthreads();
    }
    if (threadIdx.x == 0) atomicAdd(total, partial[0]); // one atomic per block
}
// launch with 256 threads per block

Chapter 7

Error checking

CUDA calls return an error code and kernel launches fail silently unless you check. Wrap every call in a macro and check cudaGetLastError() after each launch. This saves hours of debugging.

check.cuh

#include <cstdio>
#include <cstdlib>

#define CUDA_CHECK(call) do {                                        \
    cudaError_t err = (call);                                        \
    if (err != cudaSuccess) {                                        \
        fprintf(stderr, "CUDA error %s at %s:%d\n",                  \
                cudaGetErrorString(err), __FILE__, __LINE__);        \
        exit(1);                                                     \
    }                                                                \
} while (0)

// usage:
// CUDA_CHECK(cudaMalloc(&d, bytes));
// kernel<<<blocks, threads>>>(...);
// CUDA_CHECK(cudaGetLastError());
// CUDA_CHECK(cudaDeviceSynchronize());

Chapter 8

Streams and overlapping work

A stream is a queue of GPU work. Work in different streams can overlap, so the GPU can copy one batch while computing another. Use pinned host memory (cudaMallocHost) and cudaMemcpyAsync for real overlap.

streams.cu

cudaStream_t s1, s2;
cudaStreamCreate(&s1);
cudaStreamCreate(&s2);

// copy and compute two halves at the same time
cudaMemcpyAsync(d_a, h_a, half_bytes, cudaMemcpyHostToDevice, s1);
cudaMemcpyAsync(d_b, h_b, half_bytes, cudaMemcpyHostToDevice, s2);
process<<<blocks, threads, 0, s1>>>(d_a, n / 2);
process<<<blocks, threads, 0, s2>>>(d_b, n / 2);

cudaStreamSynchronize(s1);
cudaStreamSynchronize(s2);
cudaStreamDestroy(s1);
cudaStreamDestroy(s2);

Chapter 9

Programming Tensor Cores

Tensor Cores multiply small matrix tiles in one instruction. The WMMA API exposes them in CUDA C++: one warp of 32 threads loads 16×16 tiles of FP16 data, multiplies them and accumulates the result in FP32.

Production kernels use lower-level instructions (mma, and on Hopper and Blackwell wgmma and tcgen05) through libraries like CUTLASS. Read what is a Tensor Core for the hardware side.

wmma.cu

#include <cuda_fp16.h>
#include <mma.h>
using namespace nvcuda;

// C (16x16, float) = A (16x16, half) x B (16x16, half), computed by one warp
__global__ void wmma_16x16(const half* a, const half* b, float* c) {
    wmma::fragment<wmma::matrix_a, 16, 16, 16, half, wmma::row_major> fa;
    wmma::fragment<wmma::matrix_b, 16, 16, 16, half, wmma::col_major> fb;
    wmma::fragment<wmma::accumulator, 16, 16, 16, float> fc;

    wmma::fill_fragment(fc, 0.0f);
    wmma::load_matrix_sync(fa, a, 16);
    wmma::load_matrix_sync(fb, b, 16);
    wmma::mma_sync(fc, fa, fb, fc);           // the Tensor Core operation
    wmma::store_matrix_sync(c, fc, 16, wmma::mem_row_major);
}

// launch with exactly one warp:  wmma_16x16<<<1, 32>>>(d_a, d_b, d_c);
// compile for Volta or newer:    nvcc -arch=sm_70 wmma.cu

Chapter 10

CUDA libraries

You rarely need to write everything yourself. NVIDIA's libraries are tuned for every GPU generation: cuBLAS (matrix maths), cuDNN (neural network layers), NCCL (multi-GPU communication), Thrust (C++ algorithms on the GPU), CUTLASS (templates for custom Tensor Core kernels) and cuRAND, cuFFT and more.

thrust_demo.cu

#include <thrust/device_vector.h>
#include <thrust/reduce.h>
#include <thrust/sort.h>
#include <cstdio>

int main() {
    thrust::device_vector<float> v(1 << 20, 1.0f);   // lives on the GPU
    v[0] = 5.0f;
    thrust::sort(v.begin(), v.end());                 // GPU sort
    float total = thrust::reduce(v.begin(), v.end()); // GPU sum
    printf("sum = %.0f\n", total);
}

Chapter 11

Profiling and debugging

Speed comes from measuring, not guessing. Nsight Systems shows the timeline of kernels and copies; Nsight Compute analyses one kernel in depth, including memory bandwidth and Tensor Core use. compute-sanitizer finds out-of-bounds accesses and race conditions.

Compare your kernel with the roofline model to see whether it is limited by memory or compute.

terminal

nsys profile -o report ./my_app          # timeline of the whole program
ncu --set full -o kernel ./my_app        # detailed kernel analysis
compute-sanitizer ./my_app               # memory errors and races

Part II · AI with CUDA

How AI uses CUDA

When you call torch.matmul on a GPU tensor, PyTorch launches a cuBLAS kernel that runs on the Tensor Cores. Attention, normalisation and activation layers are CUDA kernels too, many written by hand for speed.

You can write your own kernel and call it from PyTorch in a few lines. This is how researchers add new operations and how projects like FlashAttention began.

The essential libraries

cuBLAS / cuBLASLtMatrix multiplication, the core of every model.

cuDNNConvolutions, attention and normalisation layers.

NCCLFast communication between GPUs for training.

TensorRT / TensorRT-LLMOptimised inference engines.

CUTLASS / CuTeBuilding blocks for custom Tensor Core kernels.

TritonA Python language that compiles to GPU kernels.

FlashAttentionMemory-efficient attention kernels.

CUDA Python / CuPyUse CUDA from Python, NumPy-style.

Write a CUDA kernel and call it from PyTorch

load_inline compiles CUDA code on the fly and exposes it as a Python function. Here we implement ReLU, one of the simplest neural network activations.

my_relu.py

import torch
from torch.utils.cpp_extension import load_inline

cuda_src = r"""
__global__ void relu_kernel(const float* x, float* y, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) y[i] = x[i] > 0.f ? x[i] : 0.f;
}

torch::Tensor relu(torch::Tensor x) {
    auto y = torch::empty_like(x);
    int n = x.numel();
    relu_kernel<<<(n + 255) / 256, 256>>>(x.data_ptr<float>(), y.data_ptr<float>(), n);
    return y;
}
"""

ext = load_inline(name="my_relu",
                  cpp_sources="torch::Tensor relu(torch::Tensor x);",
                  cuda_sources=cuda_src,
                  functions=["relu"])

x = torch.randn(1_000_000, device="cuda")
print(torch.allclose(ext.relu(x), torch.relu(x)))   # True

Part III · Build an AI agent

Power an AI agent with CUDA

Agents are usually written in Python, JavaScript or Go, but the language model they talk to runs on CUDA. Building an inference server with CUDA support is how you give your agents a fast, private, local brain.

llama.cpp is a good example: its CUDA back end runs quantized models on NVIDIA GPUs and exposes the same OpenAI-compatible API used by the agent examples in the other volumes.

Your goal→Model plans→Tool call→Your code runs the tool→Result back to model→Final answer

Install the CUDA Toolkit and CMake.
Build llama.cpp with CUDA enabled.
Download a GGUF model, for example from Hugging Face.
Start the server with all layers on the GPU (-ngl 99) and tool-calling templates (--jinja).
Point any agent at http://localhost:8080/v1/chat/completions.

terminal

git clone https://github.com/ggml-org/llama.cpp
cd llama.cpp
cmake -B build -DGGML_CUDA=ON
cmake --build build --config Release -j

# run a model fully on the GPU with an OpenAI-compatible API
./build/bin/llama-server -m models/Llama-3.1-8B-Instruct-Q4_K_M.gguf \
    -ngl 99 --port 8080 --jinja

# test it
curl http://localhost:8080/v1/chat/completions -H "Content-Type: application/json" \
  -d '{"messages":[{"role":"user","content":"Hello from my GPU!"}]}'

Use the Python agent with URL = "http://localhost:8080/v1/chat/completions". Watch nvidia-smi while it runs to see your GPU memory and utilisation, and use the LLM inference simulator to predict speed.

Run it locallyCheck which models fit your GPU with the VRAM calculator before downloading.

Part IV · Your path

Projects to build

GPU vs CPU benchmarkTime vector add and matrix multiply on both and chart the speed-up.

Image blurApply a blur filter to a large image with one thread per pixel.

Tiled matmulOptimise matrix multiply step by step and compare with cuBLAS.

Custom PyTorch opWrite a fused GELU or LayerNorm kernel and call it from PyTorch.

Softmax kernelImplement a numerically stable softmax with block reductions.

Local LLM serverBuild llama.cpp with CUDA and connect an agent to it.

Reference

CUDA cheat sheet

Task	Syntax
Kernel	`__global__ void k(float* x) {}`
Launch	`k<<<blocks, threads>>>(d_x);`
Global index	`blockIdx.x * blockDim.x + threadIdx.x`
Blocks needed	`(n + threads - 1) / threads`
Allocate	`cudaMalloc(&d, bytes);`
Copy to GPU	`cudaMemcpy(d, h, bytes, cudaMemcpyHostToDevice);`
Unified memory	`cudaMallocManaged(&p, bytes);`
Shared memory	`__shared__ float tile[16][16];`
Barrier	`__syncthreads();`
Atomic	`atomicAdd(&total, v);`
Wait for GPU	`cudaDeviceSynchronize();`
Compile	`nvcc -O3 -arch=sm_90 app.cu -o app`

Reference

Common mistakes

No bounds check in the kernel: threads past the end of the array write to random memory.
Not checking errors: a failed launch looks like "nothing happened". Use a check macro.
Too many small copies: host-device transfers are slow. Copy once and keep data on the GPU.
Uncoalesced memory access: neighbouring threads should read neighbouring addresses.
Forgetting __syncthreads(): threads read shared memory before it is filled. But never call it inside a branch that only some threads take.

Reference

Where to go next

CUDA C++ Programming Guide — the official reference
An Even Easier Introduction to CUDA — NVIDIA's beginner tutorial
CUDA C++ Best Practices Guide
GPU MODE lectures — free community course on GPU programming

On TensorCore.si: What is a Tensor Core?, What is CUDA?, How a GPU works, Roofline simulator.

Keyboard: [ previous language · ] next language · / search the library

CUDA

Why learn CUDA?

Set up and run your first program

Host, device and kernels

Threads, blocks and grids

GPU memory

Your first complete kernel: vector add

Shared memory and tiling

Synchronisation and atomics

Error checking

Streams and overlapping work

Programming Tensor Cores

CUDA libraries

Profiling and debugging

How AI uses CUDA

The essential libraries

Write a CUDA kernel and call it from PyTorch

Power an AI agent with CUDA

CUDA learning roadmap

Beginner

Intermediate

Advanced

Projects to build

CUDA cheat sheet

Common mistakes

Where to go next