Next steps: from writing kernels to knowing the GPU

After these five stages you already have the foundation to read most open-source CUDA kernels. The path splits from here, depending on the problem you want to solve. This lesson teaches no new syntax. It gives you a map.

Tensor Core: where modern GPU compute actually lives

A blunt fact: A100 FP32 compute is 19.5 TFLOPS; Tensor Core BF16 is 312 TFLOPS, a factor of 16. If you are doing matmul without Tensor Cores, you are using less than 7% of the card. There are three programming levels: abstraction goes down, control goes up.

1#include <mma.h>2using namespace nvcuda;3 4// One warp cooperates on a 16x16x16 MMA5__global__ void wmmaGemm(const half* A, const half* B, float* C,6                         int M, int N, int K) {7    // A fragment is a tile spread across the warp's 32 thread registers; layout is hardware-defined8    wmma::fragment<wmma::matrix_a, 16, 16, 16, half, wmma::row_major> a;9    wmma::fragment<wmma::matrix_b, 16, 16, 16, half, wmma::col_major> b;10    wmma::fragment<wmma::accumulator, 16, 16, 16, float> acc;11 12    wmma::fill_fragment(acc, 0.0f);13 14    int warpM = (blockIdx.x * blockDim.x + threadIdx.x) / warpSize;15    int warpN = blockIdx.y * blockDim.y + threadIdx.y;16 17    for (int k = 0; k < K; k += 16) {18        wmma::load_matrix_sync(a, A + warpM * 16 * K + k, K);19        wmma::load_matrix_sync(b, B + k * N + warpN * 16, N);20        wmma::mma_sync(acc, a, b, acc);      // this instruction is the Tensor Core21    }22 23    wmma::store_matrix_sync(C + warpM * 16 * N + warpN * 16, acc, N,24                            wmma::mem_row_major);25}

Triton: when to use it

1import triton2import triton.language as tl3 4@triton.jit5def add_kernel(x_ptr, y_ptr, out_ptr, n, BLOCK: tl.constexpr):6    pid = tl.program_id(0)                      # like blockIdx.x7    offs = pid * BLOCK + tl.arange(0, BLOCK)    # a whole tile, not one element8    mask = offs < n                             # bounds expressed as a mask9 10    x = tl.load(x_ptr + offs, mask=mask)11    y = tl.load(y_ptr + offs, mask=mask)12    tl.store(out_ptr + offs, x + y, mask=mask)13 14# No threadIdx, no __syncthreads, no shared-memory management.15# The Triton compiler decides those from BLOCK.

Advanced reading list

1. 1.CUDA C++ Programming Guide: official docs. Do not read cover to cover; use it as a dictionary. Especially Performance Guidelines and the Compute Capabilities appendix.
2. 2.CUDA C++ Best Practices Guide: more practical than the Programming Guide. Optimization advice is ordered by payoff.
3. 3.CUTLASS source: how production GEMM is organized. Start from the basic samples in examples/; do not start with include/cutlass/gemm/.
4. 4.FlashAttention paper + source: the textbook fusion case. Get the online-softmax derivation first, then read the kernel.
5. 5.Nsight Compute Kernel Profiling Guide: once every metric's definition is clear, profiling gets twice as fast.
6. 6.PTX ISA docs: look up the instruction you do not recognize. Do not read cover to cover.
7. 7.Volkov, Better Performance at Lower Occupancy (GTC 2010): the classic on the ILP vs occupancy trade-off. Still current.

Projects to practice on

• ▸Hand-write a GEMM and try to rank: from the naive version all the way to near-cuBLAS. This is the best overall exercise. Aim for 80% of cuBLAS.
• ▸Implement FlashAttention: fuse QK^T, softmax, and V into one kernel; use online softmax so the attention matrix is never materialized.
• ▸Write a fused PyTorch op: fused LayerNorm + residual, then measure end-to-end gain on a real model.
• ▸Enter a kernel-optimization contest: benches like KernelBench give a fixed problem and a leaderboard, so the feedback loop is short.
• ▸Read an inference engine's kernel directory: vLLM, llama.cpp's CUDA backend, TensorRT-LLM. Real production samples.