Reasoning Models (o1/R1), GRPO & Triton Kernels

Advanced curriculum with interactive signal rubrics, mathematical proofs, and candidate checklists.

Reasoning Models (o1/R1), GRPO/DPO Alignment & Triton Kernels

Covers QA-46 through QA-50. Master Test-Time Compute scaling, Process Reward Models (PRMs), Group Relative Policy Optimization (GRPO), and writing custom fused Triton GPU kernels.

QA-46 • Frontier Modeling

Test-Time Compute Scaling & Reasoning Models (OpenAI o1 / DeepSeek R1)

Track: Senior AI Researcher / Applied ML • Difficulty: Hard • Focus: MCTS & PRM Verification

▼

The Question: "Explain the paradigm shift from Pre-training Scaling Laws to Test-Time Compute Scaling. How do models like OpenAI o1 and DeepSeek R1 trade inference FLOPs for accuracy, and what is the role of Process Reward Models (PRMs) vs Outcome Reward Models (ORMs)?"

Whiteboard Explanation:

Pre-training Saturation: Chinchilla scaling laws govern pre-training tokens and parameters. However, internet-scale text is finite, and training larger models yields diminishing returns.
Test-Time Compute Scaling: By allocating more computation (tokens and rollouts) at inference time, reasoning performance continues to scale logarithmically without changing model weights!
PRMs vs ORMs:
- ORM (Outcome Reward Model): Scores only the final answer (0 or 1). Vulnerable to false positives (getting right answer through wrong steps).
- PRM (Process Reward Model): Scores every individual reasoning step. Enables Monte Carlo Tree Search (MCTS) or Best-of-N beam search to prune hallucinated branches before the final answer is generated.

🎯 Interviewer Signal Rubric Checklist: Stated that inference compute trades latency/cost for accuracy dynamically. Differentiated step-level PRMs from final-output ORMs. Mentioned self-correction backtracking and tree search rollout pruning.

QA-47 • Post-Training Alignment

DPO vs PPO vs GRPO: The Modern Alignment Revolution

Track: AI Engineer / Applied ML • Difficulty: Hard • Focus: Reference-Free Optimization

▼

The Question: "Why is standard RLHF with PPO notorious for instability, and how do Direct Preference Optimization (DPO) and Group Relative Policy Optimization (GRPO) eliminate the Critic network?"

The PPO Nightmare: PPO requires loading 4 separate models into GPU VRAM simultaneously: Actor (policy), Critic (value network), Reference Model (KL guard), and Reward Model. It is memory-prohibitive and sensitive to hyperparameters.
DPO (Direct Preference Optimization): Analytically solves for the optimal policy, directly re-parameterizing the reward function into the policy loss:
```
Loss_DPO = -E[log(sigmoid(beta * log(pi(y_w|x)/ref(y_w|x)) - beta * log(pi(y_l|x)/ref(y_l|x))))]
```
Eliminates the Critic and Reward models completely—trains directly on pairwise preferences like supervised fine-tuning!
GRPO (Group Relative Policy Optimization): Used by DeepSeekMath and R1. Samples a group of outputs `{o_1, o_2, ..., o_G}` for each prompt and calculates rewards relative to the group mean, eliminating the baseline value model while maintaining reinforcement learning exploration.

🎯 Interviewer Signal Rubric Checklist: Explained why PPO requires 4 concurrent models in GPU memory. Wrote out or explained the Bradley-Terry preference probability in DPO. Explained how GRPO replaces the Critic network using group-relative baseline normalization.

QA-48 • GPU Kernel Optimization

Writing Custom Fused Kernels in OpenAI Triton vs CUDA C++

Track: AI Systems / Performance Engineer • Difficulty: Hard • Focus: Memory-Fused RMSNorm & Softmax

▼

The Question: "In PyTorch, chaining `x = norm(x); x = gelu(x)` launches two separate GPU kernels that write intermediate activations back to high-latency HBM. How do you write a fused kernel in OpenAI Triton to keep data in SRAM?"

The Kernel Fusion Advantage:

import triton
import triton.language as tl

@triton.jit
def fused_rmsnorm_kernel(X_ptr, Y_ptr, W_ptr, stride, N, eps, BLOCK_SIZE: tl.constexpr):
    row_idx = tl.program_id(0)
    col_offsets = tl.arange(0, BLOCK_SIZE)
    mask = col_offsets < N
    
    # Load input row into on-chip SRAM registers
    x = tl.load(X_ptr + row_idx * stride + col_offsets, mask=mask, other=0.0)
    
    # Compute RMS variance locally in registers
    variance = tl.sum(x * x, axis=0) / N
    rsqrt = 1.0 / tl.sqrt(variance + eps)
    
    # Load scale weights and normalize without touching HBM
    w = tl.load(W_ptr + col_offsets, mask=mask)
    y = x * rsqrt * w
    
    # Write final output directly back to DRAM once
    tl.store(Y_ptr + row_idx * stride + col_offsets, y, mask=mask)

Result: Halves memory traffic to HBM, cutting kernel execution time by up to 60%.

🎯 Interviewer Signal Rubric Checklist: Explained kernel launch overhead and intermediate HBM read/write roundtrips. Identified register/SRAM accumulation as the primary performance lever. Contrasted Triton's automatic block tiling and memory coalescing vs manual CUDA thread index math.

QA-49 • Precision Formats & Quantization

FP8 (E4M3 vs E5M2) vs INT4 vs AWQ Kernel Execution

Track: AI Systems / Infrastructure • Difficulty: Hard • Focus: Hopper Tensor Cores & Outlier Scaling

▼

The Question: "Compare FP8 (E4M3 vs E5M2 formats on NVIDIA Ada Lovelace / Hopper GPUs) against weight-only quantization (AWQ/GPTQ INT4) and FP16. In what serving bottlenecks does FP8 outperform INT4, and how does per-tensor vs per-block scaling prevent activation outlier degradation?"

Whiteboard Breakdown:

E4M3 vs E5M2 Precision:
- E4M3 (1 sign, 4 exponent, 3 mantissa): Higher precision, narrower dynamic range ($[-448, 448]$). Optimal for inference forward passes and activations where numerical fidelity matters.
- E5M2 (1 sign, 5 exponent, 2 mantissa): Wider dynamic range matching FP16 exponents ($[-57344, 57344]$), but lower precision. Essential for backpropagation gradients where loss gradients span orders of magnitude.
The INT4 vs FP8 Execution Bottleneck:
- INT4 (AWQ/GPTQ): Weights are stored in 4-bit, but NVIDIA Tensor Cores do NOT support mixed INT4-weight/FP16-activation GEMM natively. The GPU kernel must dequantize INT4 weights back into FP16 in registers on every compute cycle—wasting valuable arithmetic instructions!
- FP8: Hopper H100 and Ada GPUs possess native FP8 Tensor Cores. Both weights and activations remain in 8-bit, achieving 2x raw TFLOPS compared to FP16 with ZERO register dequantization overhead.
Activation Outlier Handling: Per-tensor scaling fails when 0.1% of activation channels have large magnitudes (e.g., >100). Modern serving uses per-block scaling (Microscaling MXFP8) where vectors are partitioned into 32-element tiles with localized scaling factors.

🎯 Interviewer Signal Rubric Checklist: Correctly distinguished E4M3 (inference forward) vs E5M2 (gradient backward) bit allocations. Identified that INT4 requires on-the-fly register dequantization whereas Hopper has native FP8 GEMM. Formulated block-wise / channel-wise scaling to preserve activation outlier fidelity.

QA-50 • Frontier Inference Architecture

Speculative Decoding & Multi-Token Prediction (MTP) Engine

Track: Senior AI Architect / Systems • Difficulty: Hard • Focus: EAGLE Tree Attention & Rollback Throttling

▼

The Question: "Design a high-throughput Multi-Token Prediction (MTP) or Speculative Verification Engine for ultra-low latency code generation. How does tree-based verification (Medusa/EAGLE) differ from draft-model speculation, and how do you prevent speculation rollback bubbles?"

Architecture & Verification Flow:

Multi-Token Prediction Trunk: Instead of predicting only token $t+1$, DeepSeek-V3 and advanced architectures add $K$ lightweight prediction heads on top of the shared transformer representations. Head $k$ predicts token $t+k+1$.
Tree-Based Verification (Medusa/EAGLE): Linear draft speculative decoding produces a single sequence of $N$ tokens. If the base model rejects token 2, tokens 3 to $N$ are wasted. Tree Attention generates a tree of candidate tokens (e.g. top-2 at depth 1, top-2 at depth 2) and verifies all branches in parallel with a single base model forward pass using a custom 2D tree attention mask.
Rollback Throttling & Bubble Avoidance: If the empirical acceptance rate falls below a threshold $ au < 0.35$ (e.g. during creative writing or high sampling temperature), the engine dynamically throttles the speculation window from $K=5$ down to $K=1$, preventing wasted memory bandwidth.

🎯 Interviewer Signal Rubric Checklist: Detailed parallel auxiliary heads on top of the shared representation trunk. Contrasted linear draft rejection bubbles vs 2D tree-attention multi-branch verification. Outlined dynamic speculation depth throttling based on real-time token acceptance rate.