GPU architecture for inference engineers

Here is the launch configuration for vLLM's fused RMSNorm, read straight out of the source:

  // For large num_tokens, use smaller blocks to increase SM concurrency.
  const bool batch_invariant_launch = vllm::vllm_is_batch_invariant();
  const int max_block_size =
      batch_invariant_launch ? 1024 : ((num_tokens < 256) ? 1024 : 256);
  dim3 grid(num_tokens);

The grid — the total parallel work handed to the GPU — is one block per token. A prefill of 2048 tokens gets 2048 blocks, spread comfortably over 132 SMs. A decode step at batch 1 gets one. The comment above it is the engineer's fingerprint: block size also has to shrink once you have enough blocks, or the SMs do not fill evenly.

That is one kernel out of roughly four hundred in a single decode step, and each one costs the CPU time to launch whether the GPU work is 0.2 microseconds or 200. Serving engineering lives in the gap between "the GPU is fast" and "the GPU is fast at the shape of work you actually gave it". This chapter builds only the hardware vocabulary you need to reason about that gap — memory hierarchy, execution model, tensor-core shapes, launch cost — and names the later chapter that spends each fact.

Think of the GPU as a very wide, very shallow machine bolted to a firehose: the compute is enormous and cheap, and moving operands to it is the entire cost. A number can live in one of four places, each roughly an order of magnitude larger and a factor of four slower than the last. Kernel engineering is the art of keeping a working set in the fastest tier that can hold it.

Two facts from that table decide the shape of every attention kernel in this book.

First: shared memory in aggregate is not small. 132 SMs × 228 KB is 30 MB of SRAM — the same order as the 50 MB L2, and available at 127.9 bytes per clock per SM. Multiply out (arithmetic on the paper's per-SM figure and the H800's 1755 MHz clock) and the machine has roughly 30 TB/s of shared-memory bandwidth against 3.35 TB/s of HBM. The paper states the L2-versus-global ratio directly for Hopper: 4.23×.

Second: the thing you want to keep there does not fit. Llama-3-8B (L=32, d=4096, h=32, h_kv=8, d_h=128) stores $2 \cdot h_{kv} \cdot d_h \cdot 2$ bytes = 4 KiB of KV per token per layer in bf16 (the derivation is §2.1's). One SM's 228 KB therefore holds 57 tokens of a single layer's KV; an 8k-token sequence needs 32 MiB for one layer, 144× too big, and would evict most of L2. Worse, a naive prefill materialises the score matrix: at S=8192, one head's $QK^\top$ in fp32 is $8192^2 \cdot 4$ bytes = 268 MB — five entire L2 caches, per head, for 32 heads. That number is why FlashAttention exists. The mechanism belongs to §3.2; the capacity gap in Figure 1 is why it pays.

A CUDA kernel is a function marked __global__ that runs once per thread. You launch it with a grid of thread blocks, and the hardware assigns blocks to SMs. Four names carry the whole model, and you can read them off a real kernel.

template <typename scalar_t, int width, bool HasWeight>
__global__ std::enable_if_t<(width == 0) || !_typeConvert<scalar_t>::exists>
fused_add_rms_norm_kernel(
    scalar_t* __restrict__ input,  // [..., hidden_size]
    const int64_t input_stride,
    scalar_t* __restrict__ residual,      // [..., hidden_size]
    const scalar_t* __restrict__ weight,  // [hidden_size], null if !HasWeight
    const float epsilon, const int num_tokens, const int hidden_size,
    const int64_t residual_stride) {
  __shared__ float s_variance;
  float variance = 0.0f;

  for (int idx = threadIdx.x; idx < hidden_size; idx += blockDim.x) {
    scalar_t z = input[blockIdx.x * input_stride + idx];
    z += residual[blockIdx.x * residual_stride + idx];
    float x = (float)z;
    variance += x * x;
    residual[blockIdx.x * residual_stride + idx] = z;
  }

  using BlockReduce = cub::BlockReduce<float, 1024>;
  __shared__ typename BlockReduce::TempStorage reduceStore;
  variance = BlockReduce(reduceStore).Reduce(variance, CubAddOp{}, blockDim.x);

  if (threadIdx.x == 0) {
    s_variance = rsqrtf(variance / hidden_size + epsilon);
  }
  __syncthreads();
  // ...

Everything you need is in those 28 lines:

• blockIdx.x — which block am I? Here, the token index, because the launcher set dim3 grid(num_tokens): one block owns one row of the activation tensor. threadIdx.x is which thread I am within that block, from 0 to blockDim.x - 1.
• for (idx = threadIdx.x; idx < hidden_size; idx += blockDim.x) is the block-stride loop: 512 threads walk a 4096-element row in 8 strides, with consecutive threads on consecutive addresses so the memory system coalesces them into wide transactions.
• __shared__ puts a variable in the SM's 228 KB scratchpad — one copy per block, visible to that block and nobody else. Here, one float holding the row's inverse RMS.
• __syncthreads() is a barrier across the block: thread 0 wrote s_variance, everyone else must wait before reading it. Omit it and you get a race that produces subtly wrong logits and no error message.

The grid-stride variant — a fixed-size grid walking an arbitrarily long array — appears in vLLM's custom all-reduce (§5.4):

template <typename T, int ngpus>
__global__ void __launch_bounds__(512, 1)
    cross_device_reduce_1stage(RankData* _dp, RankSignals sg, Signal* self_sg,
                               T* __restrict__ result, int rank, int size) {
  using P = typename packed_t<T>::P;
  using A = typename packed_t<T>::A;
  // note: we don't reorder the address so the accumulation order is the same
  // for all ranks, ensuring bitwise identical results
  auto dp = *_dp;
  barrier_at_start<ngpus>(sg, self_sg, rank);
  // do the actual reduction
  for (int idx = blockIdx.x * blockDim.x + threadIdx.x; idx < size;
       idx += gridDim.x * blockDim.x) {
    ((P*)result)[idx] = packed_reduce<P, ngpus, A>((const P**)&dp.ptrs[0], idx);
  }
  barrier_at_end<ngpus, true>(sg, self_sg, rank);
}

blockIdx.x * blockDim.x + threadIdx.x is the thread's global index and gridDim.x * blockDim.x the total thread count, so each thread hops that far and comes back. __launch_bounds__(512, 1) is a contract with the compiler — at most 512 threads per block, and keep registers low enough that 1 block stays resident per SM. Your first occupancy knob.

The warp is the real unit

Threads do not execute individually. The hardware groups them into warps of 32 that issue one instruction together, and vLLM hardcodes the constant:

#else
  #define WARP_SIZE 32
#endif

The H100 whitepaper's compute-capability table lists Threads/Warp = 32 for Volta, Ampere and Hopper alike. The #else matters because the branch above it returns 64 on AMD gfx9.

Warp shuffle instructions let participating lanes exchange registers without shared memory. Volta and later support independent thread scheduling: lanes can diverge and carry per-thread execution state, so use the correct active-lane mask and synchronization contract. That is a warp shuffle; vLLM wraps it as VLLM_SHFL_XOR_SYNC at csrc/cuda_compat.h:L46-L55 (expanding to __shfl_xor_sync on CUDA, __shfl_xor on ROCm) and uses it to finish a $q \cdot k$ dot product in $\log_2$ steps with no memory traffic:

  // Finalize the reduction across lanes.
  float qk = sum(qk_vec);
#pragma unroll
  for (int mask = THREAD_GROUP_SIZE / 2; mask >= 1; mask /= 2) {
    qk += VLLM_SHFL_XOR_SYNC(qk, mask);
  }
  return qk;
}

That butterfly — mask = 8, 4, 2, 1 for a 16-lane thread group — is the standard in-warp reduction. Any halving loop over __shfl_xor_sync is doing this.

Latency hiding, and what occupancy actually buys

An HBM load costs about 479 clocks, and nothing hides that except other work. Each SM holds up to 64 resident warps (2048 threads); its schedulers pick, every cycle, any warp whose next instruction has its operands ready. More resident warps can provide independent ready work, but even 64 warps can all wait on dependencies. Four warps may suffice for a well-pipelined workload; neither number guarantees utilization. Occupancy is that ratio — resident warps over 64 on Hopper — capped by whichever resource runs out first: 65,536 registers per SM, 228 KB of shared memory per SM, 32 blocks per SM, or too few blocks in the grid.

Decode kernels have a less flattering reason for low occupancy: there is not enough work. Take the RMSNorm launcher again, for Llama-3-8B in bf16 (hidden_size 4096):

          const int calculated_vec_size =
              std::gcd(16 / sizeof(scalar_t), hidden_size);
          const int block_size =
              std::min(hidden_size / calculated_vec_size, max_block_size);
          dim3 block(block_size);

Arithmetic on those two lines: 16 / sizeof(bf16) is 8 and gcd(8, 4096) is 8, so block_size = min(4096 / 8, 1024) = 512 threads — 16 warps. At decode batch 1 the grid is 1 block: one SM of 132 runs 16 of its 64 warp slots. The machine is at 0.76% of its SMs occupied, 25% occupancy on the one that is. Not a bug — it is what "normalise one vector of 4096 numbers" looks like on a 132-SM GPU, and it is why §1.3 works so hard to make the batch bigger.

A concrete coalescing and occupancy calculation

For 32 lanes reading adjacent float32 elements at aligned addresses base + 4*lane, the warp requests 128 contiguous bytes, covering four 32-byte sectors. A stride of 32 elements requests 32 separated sectors, or 1024 transferred sector bytes for the same 128 useful bytes, before cache reuse. Alignment and architecture-specific caching still matter.

Suppose a block has 256 threads, 64 registers per thread, and 64 KiB shared memory. Ignoring allocation granularity, registers allow floor(65536/(256*64))=4 blocks; 228 KiB shared memory allows floor(228/64)=3; the 2048-thread limit allows 8. Shared memory therefore caps residency at 3 blocks, 24 warps, or 37.5% theoretical occupancy. A grid with one block per SM reduces the realized limit further. Verify actual allocation with compiler output and Nsight Compute rather than treating this simplified arithmetic as a measurement.

Tensor cores want a rectangle, decode gives them a line

An H100 SXM5 has 132 SMs, 128 FP32 CUDA cores per SM (16,896 total), and four fourth-generation tensor cores per SM (528 total). The whitepaper's preliminary table puts peak dense BF16 tensor throughput at 1000 TFLOPS against 120 TFLOPS for non-tensor BF16 and 60 TFLOPS for non-tensor FP32; the shipping datasheet lists 1,979 TFLOPS BF16 (a with-sparsity figure, so 989.5 dense) and 67 TFLOPS FP32. Either way: a kernel that cannot use the tensor cores gives up roughly 15× of the machine.

Tensor cores are not general multipliers. They execute a fixed matrix-multiply-accumulate shape, issued per warp — here, in the inline PTX of vLLM's Marlin quantized-GEMM kernel:

// m16n8k16 tensor core mma instruction with fp16 inputs and fp32
// output/accumulation.
template <vllm::ScalarTypeId type_id, bool use_fp16_accum, int k_size = 16>
__device__ inline void mma(
// ...
#else
      float* c = reinterpret_cast<float*>(&frag_c);
      asm volatile(
          "mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 "
          "{%0,%1,%2,%3}, {%4,%5,%6,%7}, {%8,%9}, {%10,%11,%12,%13};\n"
          : "=f"(c[0]), "=f"(c[1]), "=f"(c[2]), "=f"(c[3])
          : "r"(a[0]), "r"(a[1]), "r"(a[2]), "r"(a[3]), "r"(b[0]), "r"(b[1]),
            "f"(c[0]), "f"(c[1]), "f"(c[2]), "f"(c[3]));

Read the mnemonic: m16n8k16 is M=16, N=8, K=16 — one warp, one instruction, $16 \times 8$ outputs accumulated over a K of 16. In a linear layer, M is tokens. So:

For the illustrated mapping with M=1 and one token per row, fifteen of sixteen A-fragment rows are padding: only 6.25% of the row slots carry useful data. This is a utilization calculation for that mapping, not a universal limit on every decode kernel. Hopper's illustrated warp-group shape has M=64 — the Hopper microbenchmark study shows wgmma compiling to SASS shapes like HGMMA.64x256x16, an M of 64, wasting 63 rows of 64. And 6.25% of 989 dense BF16 TFLOPS is about 62 TFLOPS, roughly the non-tensor FP32 rate: this padded mapping may offer little arithmetic advantage at batch 1. Many decode shapes are bandwidth-limited, but libraries may still select tensor-core, quantized GEMM, or CUDA-core GEMV paths. The following AMD-specific heuristic illustrates a bandwidth-oriented tiling choice; it does not establish NVIDIA kernel selection:

    if M <= 64:
        # decode (M in {1,32,64}): tiny-M GEMV is weight-BW + GPU-OCCUPANCY bound. The
        # lever is NARROW BLOCK_N=16 (maximize N-tiles so more CUs stream the weight in
        # parallel) + LARGE BLOCK_K (fewer K-iters, bigger coalesced weight loads).
        # Tuned by CUDA-graph replay latency. Optimal at both TP=4 and TP=8.

(That is an AMD gfx950 path, but the reasoning is architecture-independent: at tiny M you stop optimising for arithmetic and start optimising for how many compute units can stream the weight matrix in parallel.) Prefill is the opposite regime — a 2048-token chunk gives M=2048, the tensor cores run at their designed shape, and the layer becomes compute-bound. §1.1 is built on this split; §0.4 makes it quantitative.

Every kernel costs microseconds before it computes anything

Launching a kernel is not free. NVIDIA's own CUDA Graphs introduction measures it: on a Tesla V100 with CUDA 10.1, a kernel over 500,000 elements with 512 threads per block took 2.9 µs to execute but 9.6 µs per kernel including overheads when each launch was synchronised. Overlapping launches with execution brought the effective time to 3.8 µs; replaying the same sequence as a captured CUDA graph, 3.4 µs — cutting per-kernel overhead from 6.7 µs to 0.5 µs.

Now count the launches in a decode step. Llama-3-8B has 32 layers, and an unfused eager forward issues roughly a dozen kernels per layer: input RMSNorm, QKV projection, RoPE, the KV-cache write, attention, output projection, residual add, post-attention norm, gate and up projections, SiLU-mul, down projection, residual add. That is on the order of 400 launches per token, plus the final norm, LM head and sampling chain — arithmetic, not a measurement, and sensitive to how much your build fuses. At the blog's pipelined 0.9 µs of overhead per launch that is 360 µs per decode step before any arithmetic happens; at the synchronised 6.7 µs, 2.7 ms. Both are the same order as an entire decode step. That is "launch-bound", and it is why both engines capture CUDA graphs.

The cost is real enough that vLLM guards individual sampler launches at runtime:

        temp_np = self.temperature.np[idx_mapping_np]
        if np.all((temp_np == 0.0) | (temp_np == 1.0)):
            # No request requires temperature. Skip the kernel launch.
            return

        apply_temperature(logits, expanded_idx_mapping, self.temperature.gpu)

    def apply_min_p(
        self,
        logits: torch.Tensor,
        expanded_idx_mapping: torch.Tensor,
        idx_mapping_np: np.ndarray,
    ) -> None:
        if np.all(self.min_p.np[idx_mapping_np] == 0.0):
            # No request uses min_p. Skip the kernel launch.
            return
        apply_min_p(logits, expanded_idx_mapping, self.min_p.gpu)

A CPU-side np.all over the batch is cheaper than a no-op kernel launch — the whole economics of decode in one guard clause. The structural fix, recording the launches once and replaying them as a graph, is §8.1; both engines have a module for it: vllm/v1/cudagraph_dispatcher.py:L15-L32 and python/sglang/srt/model_executor/cuda_graph_config.py:L88-L107.

Both engines need to merge two partial attention results — the split-KV fixup that §3.3 owns. vLLM writes it in CUDA; SGLang keeps a Triton version as the portable fallback. Side by side, they show what a launch configuration is.

  // Process one pack elements per thread. for float, the
  // pack_size is 4 for half/bf16, the pack_size is 8.
  const uint threads_per_head = head_size / pack_size;
  const uint total_threads = num_tokens * num_heads * threads_per_head;

  dim3 block(NUM_THREADS);
  dim3 grid((total_threads + NUM_THREADS - 1) / NUM_THREADS);

and the kernel body that decodes its own position out of that flat grid:

  const uint pack_size = 16 / sizeof(scalar_t);
  const uint threads_per_head = head_size / pack_size;

  const uint global_idx = blockIdx.x * NUM_THREADS + threadIdx.x;
  const uint token_head_threads = num_tokens * num_heads * threads_per_head;

  if (global_idx >= token_head_threads) return;

  // global_idx -> token_idx + head_idx + pack_idx
  const uint token_head_idx = global_idx / threads_per_head;
  const uint pack_idx = global_idx % threads_per_head;

  const uint token_idx = token_head_idx / num_heads;
  const uint head_idx = token_head_idx % num_heads;

Three things generalise. One: NUM_THREADS is constexpr uint NUM_THREADS = 128 at csrc/libtorch_stable/attention/merge_attn_states.cu:L268, a compile-time constant, so the index arithmetic folds away. Two: the ceiling division leaves the last block partially idle, which is why line 42 has a bounds guard — every flat-index kernel needs one. Three: pack_size = 16 / sizeof(scalar_t) makes each thread move a 16-byte uint4, the widest single load the memory system supports. Wide, aligned, coalesced loads are how you reach 3.35 TB/s instead of a fraction of it.

SGLang's Triton version of the same operation makes the launch configuration explicit in Python:

@triton.jit
def merge_state_kernel(
    output,  # [NUM_TOKENS, NUM_HEADS, HEAD_SIZE] v_merged
    output_lse,  # [NUM_TOKENS, NUM_HEADS] s_merged
    prefix_output,  # [NUM_TOKENS, NUM_HEADS, HEAD_SIZE] v_a
    prefix_lse,  # [NUM_TOKENS, NUM_HEADS] s_a
    suffix_output,  # [NUM_TOKENS, NUM_HEADS, HEAD_SIZE] v_b
    suffix_lse,  # [NUM_TOKENS, NUM_HEADS] s_b
    HEAD_SIZE: tl.constexpr,
    PADDED_HEAD_SIZE: tl.constexpr,
    OUTPUT_LSE: tl.constexpr,
):
    token_idx = tl.program_id(0)
    num_tokens = tl.num_programs(0)
    head_idx = tl.program_id(1)
    num_heads = tl.num_programs(1)

    p_lse = tl.load(prefix_lse + token_idx * num_heads + head_idx)
    s_lse = tl.load(suffix_lse + token_idx * num_heads + head_idx)
# ...
    num_tokens = output.shape[0]
    num_query_heads = output.shape[1]
    head_size = output.shape[2]
    padded_head_size = triton.next_power_of_2(head_size)

    merge_state_kernel[(num_tokens, num_query_heads)](

The difference is abstraction level. In CUDA you name a thread and hand-derive its token, head and pack from a flat index. In Triton you name a program — roughly a thread block — with tl.program_id(0) and tl.program_id(1), and the compiler decides how many threads and warps back it and how to vectorise the loads. The 2-D grid [(num_tokens, num_query_heads)] makes vLLM's hand index arithmetic disappear, and head_mask = head_arange < HEAD_SIZE replaces the early-return guard, which is how Triton supports head sizes that are not powers of two.

Neither is better. Triton is far faster to write and read, which is why SGLang's Triton attention backend exists; hand-written CUDA still wins where you need exact control of shared-memory staging and tensor-core fragment layout — a tradeoff §8.3 covers properly. Note that SGLang treats Triton as the fallback: at python/sglang/srt/layers/attention/merge_state.py:L34-L46 it calls the compiled merge_state_v2 when the dtype and head dim allow, and drops to Triton otherwise.

Llama-3-8B decoding at batch 32 with split-KV attention: merge_attn_states runs on an output of shape [32, 32, 128] in bf16 — 32 tokens (one per sequence), 32 query heads, head_size 128.

1. merge_attn_states(...) at csrc/libtorch_stable/attention/merge_attn_states.cu:L337-L362 checks dtypes, then dispatches on prefix_output.scalar_type() into merge_attn_states_launcher<__nv_bfloat16>.
2. In the launcher, pack_size = 16 / sizeof(bf16) = 8, so threads_per_head = 128 / 8 = 16. Each thread will move exactly one 16-byte pack.
3. total_threads = 32 × 32 × 16 = 16,384.
4. dim3 block(128); dim3 grid((16384 + 127) / 128) = 128.
5. The <<<grid, block, 0, stream>>> in the LAUNCH_MERGE_ATTN_STATES macro (csrc/libtorch_stable/attention/merge_attn_states.cu:L221-L237) enqueues 128 blocks of 128 threads, 0 bytes of dynamic shared memory, on the current stream.
6. The grid can occupy at most 128 SMs simultaneously, without a promised one-block-per-SM placement. Each block has 4 warps; an SM holding exactly one block would have 6.25% occupancy.
7. Thread 5 of block 0: global_idx = 5, token_head_idx = 5/16 = 0, pack_idx = 5 % 16 = 5, so token_idx = 0, head_idx = 0. It loads one uint4 from prefix_output and one from suffix_output at element offset 5 × 8 into head 0 of token 0, rescales by the softmax weights derived from the two log-sum-exps, and stores one uint4.

The last box is the punchline. This kernel moves 0.79 MB, so it has about 0.24 µs of ideal compulsory-traffic HBM time (arithmetic: 786,432 bytes at 3.35 TB/s), while the launch costs at least a microsecond of CPU-side driver work and possibly seven. The launch is several times more expensive than the work. Repeat that across the ~400 launches in a decode step and you have derived why CUDA graphs are not an optimisation but a requirement.

Shared memory above the default limit requires opt-in

CUDA gives a kernel 48 KB of dynamic shared memory by default even though Hopper allows 228 KB. Ask for more without opting in and the launch fails with cudaErrorInvalidValue — "invalid argument" — usually reported at some unrelated later synchronisation point. The opt-in is a call vLLM wraps at csrc/cuda_compat.h:L70-L76 as VLLM_DevFuncAttribute_SET_MaxDynamicSharedMemorySize, which expands to cudaFuncSetAttribute(FUNC, cudaFuncAttributeMaxDynamicSharedMemorySize, VAL). The companion query lives at csrc/cuda_utils.h:L29-L31:

int64_t get_device_attribute(int64_t attribute, int64_t device_id);

int64_t get_max_shared_memory_per_block_device_attribute(int64_t device_id);

If an attention backend works on A100 and dies on another card, check this first — A100 allows 164 KB, Hopper 228 KB, consumer cards far less.

Errors surface late, at the wrong line

Launches are asynchronous. A failed launch, or a kernel writing out of bounds, typically raises at the next synchronising call — a .cpu() hundreds of lines later, in innocent code. Set CUDA_LAUNCH_BLOCKING=1 to get the real line, then turn it off: it also destroys the launch pipelining above and will change your timings by an order of magnitude.

Hardcoding 32 for the warp size

The reduction at csrc/libtorch_stable/attention/attention_utils.cuh:L38-L47 is parameterised on THREAD_GROUP_SIZE, not a literal 32, and csrc/cuda_compat.h:L7-L38 returns 64 on AMD __GFX9__. A butterfly reduction with a hardcoded 32 does not crash on a 64-wide warp; it silently reduces half the lanes and produces wrong attention scores. Silent numerical corruption on one vendor's hardware is the worst class of bug here.

Chasing occupancy

The common self-inflicted wound is tuning until Nsight Compute reports high occupancy and finding the kernel slower. Raising occupancy usually means cutting registers or shared memory per thread, which means spilling to L1 or re-reading from L2 — trading a stall you could hide for traffic you cannot.

Measuring a launch benchmark and calling it a decode benchmark

Inter-token latency with CUDA graphs disabled largely measures your CPU, your Python and your driver. Numbers with --enforce-eager are not comparable to numbers without it — §10.3.

1. Read your own card's numbers and check them against the whitepaper. Every figure in Figure 1 is queryable:

python3 -c "import torch; p = torch.cuda.get_device_properties(0); print(p)"

You are looking for multi_processor_count (132 on H100 SXM), max_threads_per_multi_processor (2048), regs_per_multiprocessor (65536), shared_memory_per_multiprocessor and warp_size. Attribute names vary slightly across PyTorch versions, which is why the command prints the whole object.

2. Count the kernels in one decode step. Run a short generation under the PyTorch profiler with --enforce-eager, then group by kernel name and count. The number of CUDA kernel invocations per decode iteration (not distinct names) is the number this chapter has been reasoning about — and it is the number CUDA graphs collapses to one replay.

vllm serve meta-llama/Meta-Llama-3-8B-Instruct --enforce-eager --max-model-len 4096
# then, in a second shell, drive one request and compare ITL against the same
# server started WITHOUT --enforce-eager. The flag is defined at
# vllm/engine/arg_utils.py:L880 and defaults to False in vllm/config/model.py:L241.

3. Predict before you look. Which moves more when you drop --enforce-eager: time-to-first-token, or inter-token latency? Write your answer down before running it. The reasoning is in this chapter; the confirmation is §8.1.

4. Read one kernel end to end. csrc/libtorch_stable/attention/merge_attn_states.cu is 362 lines and the most approachable real CUDA in either repository: one kernel, one launcher, no templated tile machinery. Explain every line of its index arithmetic and you can read §3.3.

• NVIDIA H100 Tensor Core GPU Architecture whitepaper (v1.02). Table 3 has SM counts, cache sizes and peak rates; Table 4 has the compute-capability 9.0 limits — 32 threads/warp, 64 warps/SM, 2048 threads/SM, 32 blocks/SM, 65,536 registers/SM, shared memory up to 228 KB. Every capacity number here came from those two tables.
• NVIDIA H100 datasheet — the shipping SXM numbers (80 GB, 3.35 TB/s, 1,979 BF16 tensor TFLOPS with sparsity, 67 FP32 TFLOPS), which differ slightly from the whitepaper's preliminary figures.
• Alan Gray, "Getting Started with CUDA Graphs", NVIDIA Developer Blog — the source of every launch-overhead number here.
• Luo et al., "Benchmarking and Dissecting the Nvidia Hopper GPU Architecture", IPDPS 2024. Table IV is the memory-latency table; Table V the per-clock bandwidths and the 4.23× L2-versus-global ratio; Table VI the SASS shapes wgmma compiles to.
• CUDA C++ Programming Guide, §"Hardware Implementation", §"Technical Specifications per Compute Capability", and the shared-memory cudaFuncSetAttribute opt-in.
• vLLM PR #16173 — "[Kernel] support merge_attn_states CUDA kernel, 3x speedup", where the kernel traced in §5 landed; vLLM PR #47361 — "Delete PagedAttention", which removed the last hand-written attention kernel from csrc/.
• Triton tutorials — the vector-add and fused-softmax lessons. Thirty minutes there makes every Triton kernel in SGLang readable, and §8.3 assumes you have done them.