Advanced GPU Optimization: How to tech an LLM with CUDA and ROCm? - Part 5 (Final Part) A developer's final part of a CUDA/ROCm optimization series details implementing Mixture of Experts (MoE) inference, including a custom top-k gating kernel, expert parallelism with All-to-All communication, and CPU offloading for large models. The post provides code for a sparse softmax router and describes using NCCL/RCCL for token dispatch, enabling trillion-parameter models to run on multi-GPU setups. Welcome back, you absolute madman You finished Part 4, implemented Flash Attention, and squeezed FP8 out of your silicon. But the industry doesn't stop at dense Transformers. In 2024/2025, every major model Grok, Mixtral, Gemini uses Mixture of Experts MoE to scale to trillions of parameters without exploding compute costs. Furthermore, if you actually try to run these monsters on a single node, you will hit the VRAM wall instantly. That is where CPU Offloading and Zero-Inference come to the rescue. In this fifth and I swear final part, we will: Prerequisites: Parts 1-4, a multi-GPU setup 4+ is ideal , and a realization that training is only half the battle—serving is the other half. A dense 1.5 Trillion parameter model would require 3TB of VRAM just for weights. MoE solves this by having hundreds of "expert" FFNs, but only activating 2 of them per token. The Math: y = Σ Softmax Router x i Expert i x for the top-k experts usually k=2 . 1.1 The Router Kernel Top-k Gating First, we need a kernel that takes the token embeddings x and computes the routing scores, then selects the top-2 indices and their weights. js global void moe router kernel const float x, float gate logits, int expert indices, float expert weights, int num tokens, int num experts { int tid = blockIdx.x blockDim.x + threadIdx.x; if tid = num tokens return; // Compute logits for this token dot product with router weights // Assume router weights are in shared memory or loaded via L2 cache float max1 = -INFINITY, max2 = -INFINITY; int idx1 = -1, idx2 = -1; for int e = 0; e < num experts; ++e { float score = 0.0f; for int d = 0; d < D; ++d { score += x tid D + d router weights e D + d ; } gate logits tid num experts + e = score; // Track top-2 manual reduction if score max1 { max2 = max1; idx2 = idx1; max1 = score; idx1 = e; } else if score max2 { max2 = score; idx2 = e; } } // Apply Softmax only to the top-2 scores sparse softmax float denom = expf max1 - max1 + expf max2 - max1 ; // numerically stable expert indices tid 2 = idx1; expert indices tid 2 + 1 = idx2; expert weights tid 2 = 1.0f / denom; expert weights tid 2 + 1 = expf max2 - max1 / denom; } 1.2 Expert Parallelism with All-to-All Communication In dense models, we used All-Reduce. In MoE, we use All-to-All because different GPUs hold different experts. Tokens must be sent to the GPU that owns their assigned expert. Implementation using NCCL/RCCL: // Step 1: Token dispatch send tokens to expert owners // We build a send buffer per rank based on the routing decisions. ncclGroupStart ; for int r = 0; r < world size; ++r { // Send tokens assigned to expert owned by rank 'r' if send counts r 0 { COMM SEND send buffer r , send counts r D, COMM FLOAT, r, comm, stream ; } // Recv tokens for local experts COMM RECV recv buffer r , recv counts r D, COMM FLOAT, r, comm, stream ; } ncclGroupEnd ; // Step 2: Run local experts FFN on the received tokens locally. run local experts kernel<<<... recv buffer, local expert weights, ... ; // Step 3: All-to-All again to return the processed tokens to their original GPUs. This ensures we scale to thousands of experts across a cluster with minimal communication overhead. When you have 70B parameters, even ZeRO-3 might not fit if you only have 40GB VRAM. ZeRO-Offload moves the optimizer states momentum, variance and sometimes the gradients to CPU RAM DDR . The trick is overlapping the GPU computation with PCIe transfers. During the backward pass, we asynchronously copy gradients to CPU, compute the AdamW update on the CPU using a separate thread , and copy the updated FP32 weights back to the GPU just in time for the next forward pass. C++ implementation using pinned memory and streams: // Allocate pinned memory on CPU for offloaded states float cpu momentum; float cpu variance; hipHostMalloc &cpu momentum, num params sizeof float , hipHostMallocDefault ; // During the training step: void offloaded train step { // 1. Forward/Backward on GPU FP16 - dW stays in GPU memory temporarily. backward pass ... ; // 2. Asynchronous D2H copy of gradients for the offloaded partition // while GPU continues computing the next layer. hipMemcpyAsync cpu gradients, d gradients, partition size, hipMemcpyDeviceToHost, data stream ; // 3. CPU thread computes: m = beta1 m + 1-beta1 grad, etc. // This runs on a std::async thread to not block the main loop cpu adam update cpu momentum, cpu variance, cpu gradients, partition size ; // 4. Asynchronous H2D copy of updated weights back to GPU hipMemcpyAsync d model weights, cpu weights, partition size, hipMemcpyHostToDevice, data stream ; // 5. Synchronize streams at the end. } Why this works: PCIe Gen 5.0 can do ~32 GB/s. If we overlap this with the 1-2 seconds it takes to compute a backward pass on a large model, the offloading overhead becomes nearly invisible. If you profile a real LLM, you will see that kernel launch overhead the CPU time spent telling the GPU to do things is surprisingly high—often 10-20 microseconds per kernel. A GPT-3 forward pass has ~1,000 kernels. That's 20ms wasted just on launching. CUDA Graphs and HIP Graphs record a sequence of kernel launches and replay them with a single API call. This is crucial for inference. Implementation: hipGraph t graph; hipGraphExec t instance; // 1. Capture the graph in a stream hipStreamBeginCapture stream, hipStreamCaptureModeGlobal ; // Launch all kernels for the forward pass layernorm kernel<<<..., stream ... ; matmul kernel<<<..., stream ... ; flash attention<<<..., stream ... ; // ... everything ... hipStreamEndCapture stream, &graph ; // 2. Instantiate it this compiles it down to a single executable hipGraphInstantiate &instance, graph, NULL, NULL, 0 ; // 3. Replay it every iteration launches all kernels in ~1 microsecond for int i = 0; i < 1000; ++i { // Update input pointers if needed using memcpy or host-side updates hipGraphLaunch instance, stream ; hipStreamSynchronize stream ; } For static shapes fixed sequence length, batch size , this gives a massive 10-15% end-to-end speedup. During autoregressive generation e.g., ChatGPT , we compute the Key K and Value V for every token and store them to avoid recomputing. This is the KV Cache. 4.1 Pre-allocated KV Cache Instead of dynamically allocating memory per token, we pre-allocate a contiguous buffer. // Shape: batch, num heads, max seq len, d head float kv cache; hipMalloc &kv cache, batch num heads max seq len d head 2 sizeof float ; // During decoding, we write the current token's K and V into the 'pos' slot. global void append kv kernel float cache, const float K, const float V, int batch, int head, int pos, int d head { // Write K cache offset + pos d head + idx = K idx ; // Write V stored contiguously after K cache offset + max seq len d head + pos d head + idx = V idx ; } 4.2 PagedAttention vLLM Style If you have multiple sequences of different lengths, contiguous KV caches lead to massive fragmentation memory is wasted because you allocate max seq len for everyone . PagedAttention virtualizes the KV cache into "pages" blocks in GPU memory, similar to OS virtual memory. We implement a simple block table: php struct BlockTable { int block ids; // maps logical block - physical block address int num blocks; }; // Instead of pos, we compute physical address: physical addr = block id block size + offset in block global void paged attention kernel ..., int block table, int block size { int block id = block table logical block ; int physical pos = block id block size + offset; // ... load K/V from physical pos ... } This completely eliminates memory waste and allows you to serve 3x more concurrent users on the same hardware. Finally, here is the loop for a production-grade inference server: void serve requests std::vector