AI-Generated Kernel Engineering

Suppose an AI coding agent generates a CUDA extension that clears its visible smoke test and reports an eye-popping 1.4x speedup on an H100 GPU. The pull request looks clean, the pull request comments celebrate the win, and the kernel merges into main. Two days later in staging, inference worker threads hang indefinitely, customer prompts with 1009 tokens trigger silent NaN outputs, and uncoalesced memory reads drop real memory bandwidth across the cluster from 3.35 TB/s down to 120 GB/s.

During post-mortem debugging, engineers discover the root cause: the model generated a single specialized branch for the exact public test shape [2, 4096]. Every non-standard row length triggered an out-of-bounds shared memory read, and a missing barrier created a race condition that passed the unit test only because thread warps ran in lockstep on an idle GPU.

Generated GPU code isn't an implementation; it's an untrusted proposal. A language model predicts token sequences based on open-source patterns rather than executing hardware instructions or reasoning about memory controllers, thread divergence, and cache lines. Promoting generated code into a production library requires undeniable empirical evidence: the candidate must preserve mathematical semantics across its full domain, execute safely without memory corruption, beat a comparable baseline fairly, and roll back instantly when production workloads disagree with lab benchmarks.

The GPU Runtime Execution Lab made host submission, device streams, static storage, fallback paths, and timing scopes explicit. Generated-kernel engineering takes ownership one layer lower: prove that the candidate symbol behind that dispatch computes the correct operator safely before the runtime ever replays or promotes it.

One realistic operator carries this process throughout the chapter: weighted root mean square normalization (RMSNorm), weighted_rmsnorm(x, weight, eps). It rescales each row across its last dimension and applies a learned feature weight. For an example row [3.0, 4.0], the mean square is $(9 + 16) / 2 = 12.5$. Divide each entry by $\sqrt{12.5 + \epsilon}$, then multiply by its corresponding weight. The positive scalar $\epsilon$ ensures an all-zero row never divides by zero.

For row index $i$, feature index $j$, and row width $H$:

$$y_{i,j} = x_{i,j}\, w_j\left(\frac{1}{H}\sum_{k=1}^{H}x_{i,k}^{2}+\epsilon\right)^{-1/2}$$

The reference implementation is pure PyTorch. A candidate may use CUDA, Triton, or compiler code. The promotion gates stay identical. The executable examples below run on CPU: they test evaluator logic, proving how to audit a candidate before giving it device access.

AI kernel generation dynamics and common failure modes

LLMs generate CUDA and Triton by pattern-matching code patterns from open-source repositories like FlashAttention and PyTorch. They can reproduce idiomatic launch grids, shared memory declarations, and warp-reduction syntax effortlessly. Yet that high syntactic fluency hides lethal semantic blind spots: models struggle with spatial coordinate transforms, thread-block boundary conditions, and hardware memory hierarchies.

Three failure modes recur across generated GPU kernels:

1. Subtle race conditions and missing barriers: In reduction operators like RMSNorm, threads compute partial sums and exchange them through shared memory (smem). LLMs regularly omit __syncthreads() in CUDA or block-level barriers in Triton between write and read stages, assuming warp-synchronous execution or implicit ordering. On a small test tensor where warps run in lockstep across a single Streaming Multiprocessor (SM), the test passes reliably. In production, with 132 SMs executing asynchronous warps under dynamic clock frequency scaling, the kernel suffers Read-After-Write (RAW) data hazards, corrupting outputs non-deterministically.
2. Out-of-bounds shared memory indexing: Block tile sizes (such as 128 or 256 threads) rarely divide real-world tensor dimensions cleanly. When a sequence length or hidden dimension is prime (such as $H = 1009$), edge tiles must mask out inactive threads. LLMs often guard global memory reads with if (idx < H), but forget to mask the shared memory store or fail to clear reduction padding. Uninitialized shared memory floats get summed into the reduction tree, quietly poisoning the denominator.
3. Uncoalesced memory access passing small tests by luck: Models frequently confuse column-major and row-major layout indexing, swapping threadIdx.x and threadIdx.y or indexing across rows instead of along contiguous row elements. On a tiny smoke test (such as $2 \times 4$ floats), the entire tensor fits inside a single 128-byte cache line; the test reports a false speedup because there's no memory latency penalty. Under production batching (such as batch 64, sequence length 4096, hidden dim 8192), non-consecutive addresses force the memory controller to issue 32 distinct 32-byte sector transactions instead of one coalesced 128-byte transaction. Effective memory bandwidth plummets from 3.35 TB/s down to 120 GB/s.

Treating generated code as untrusted build input protects your cluster from these hidden defects.

Write the operator contract before the prompt

Prompt engineering can't repair an underspecified operation. A code generator needs both the exact mathematical formula and the library contract surrounding it. PyTorch custom operators declare mutation, aliasing, and schemas via torch.library.custom_op, while torch.library.opcheck verifies registration, fake-tensor integration, and compilation mechanics rather than numerical correctness.^{[1]Reference 1The Custom Operators Manualhttps://docs.pytorch.org/docs/main/library.html} Treat those as distinct obligations.

Here is the contract for the running operator:

The dispatch predicate is an active gate in the contract. Suppose the first generated kernel supports CUDA, contiguous x and weight, H divisible by 128, and all three dtypes. The wrapper must verify those conditions, check the target GPU architecture, and confirm that backward autograd isn't needed; any other valid call routes to the reference. Invalid public inputs must still raise the declared library error. A fast-path restriction is acceptable; an unannounced semantic limitation isn't.

A contiguous last dimension alone doesn't justify flattening rows: sliced outer dimensions leave gaps between them. Either compute row addresses from strides or require the entire input tensor to be contiguous. Active autograd requires both torch.is_grad_enabled() and at least one input tensor with requires_grad=True.

Use a prompt specification that exposes obligations without leaking private evaluation fixtures:

Exact evaluation shapes and adversarial values stay outside the model context. Structural rules remain visible so the generator implements the intended program rather than guessing benchmark inputs.

What must a reviewer be able to state before generation begins?

Answer

A reviewer must state the accepted shapes, dtypes, strides, devices, aliasing, error behavior, numerical policy, and fallback predicate without reading the candidate. If an axis is unclear, the contract isn't frozen.

Isolate generated code and the toolchain

Generated kernel source is untrusted code. Compiling it invokes preprocessors, linkers, build scripts, dynamic linkers, and shell hooks. Run the compiler and generated binaries in a disposable sandbox with a pinned toolchain image, no credentials, no network access, tight CPU and memory limits, a wall-clock timeout, and a single writable build directory. Expose only public smoke fixtures there. An external orchestrator calls the model provider without exposing API keys to the build worker. Allow-list compiler commands and artifact types. Keep the evaluator, hidden fixtures, promotion registry, and production codebase completely outside the sandbox.

Device execution requires a second security boundary. Running untrusted binaries on a shared GPU worker risks memory snooping, kernel hangs, and driver crashes that destabilize co-located workloads. Execute tests on an isolated, short-lived GPU worker with no production traffic. Expose only evaluator-allocated input and output buffers, kill the worker process on timeout or illegal memory access, and reset or quarantine the device before reuse. A host container doesn't prevent arbitrary device code from executing illegal instructions or exhausting GPU resources. The compiler sandbox limits host-side build capabilities; the disposable GPU worker isolates hardware faults.

These boundaries mirror secure code-generation agent architectures: model output receives only the capabilities required for its task, avoiding ambient authority over the repository or secrets. Record every tool call, compiler invocation, exit code, and artifact hash. Reject attempts to modify the harness, reference, timing harness, or test suites.

Catch a candidate that only handles the visible shape

A minimal evaluator demonstrates incomplete domain coverage immediately. The CPU example below inspects a deliberately flawed candidate that computes the operator for the public contiguous 2 x 4 float32 shape, but returns zeros for any other input.

Scaling and negation preserve tensor shape, so the candidate's hardcoded branch still passes them. The new shape and transposed view expose the missing domain. Shape specialization is acceptable when a wrapper routes unsupported calls to a verified fallback. Returning zeros for unhandled inputs violates the operator contract.

The autonomous evaluation harness: six sequential gates

An autonomous evaluation harness subjects candidate GPU code to a battery of deterministic gates. Skipping early gates to run benchmarks wastes GPU compute on broken code. The pipeline enforces six sequential checks:

1. Reference oracle

The reference oracle defines functional ground truth. Written in pure PyTorch, it enforces the formal operator schema, checks input metadata, and produces the baseline output tensor. It operates with full precision accumulation (FP32) to prevent precision loss.

2. Compilation check and static inspection

Compiling with nvcc -O3 and target architecture flags (such as -arch=sm_90a for NVIDIA H100) produces diagnostic logs from ptxas, the PTX optimizing assembler.^{[2]Reference 2NVIDIA CUDA Compiler Driver NVCChttps://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/} Static inspection extracts three critical metrics:

• Register allocation and spills: Inspect the ptxas output string (e.g., ptxas info : Used 64 registers, 0 bytes cmem, 1024 bytes spill stores, 1024 bytes spill loads). If a thread requires more registers than the hardware allows per warp, ptxas spills variables into local memory (backed by high-latency DRAM), cratering execution speed.
• Shared memory allocation: High shared memory usage per thread block reduces active warp occupancy per SM. For instance, requesting 64 KB of shared memory on an SM with a 99 KB capacity limits concurrency to one block per SM.
• Symbol verification: Tools like cuobjdump and nvdisasm confirm that the compiled .cubin contains the declared symbol rather than falling back to host PyTorch routines or embedding precomputed constant tables.^{[3]Reference 3CUDA Binary Utilitieshttps://docs.nvidia.com/cuda/cuda-binary-utilities/}

3. Multi-scale numerical verification

A candidate that passes on powers of two can fail on boundary conditions. Test families must probe diverse dimensions and memory alignments:

• Powers of two: Shapes like $H \in \{128, 512, 4096\}$ test aligned, coalesced memory access paths.
• Prime row lengths: Odd and prime dimensions like $H \in \{1009, 4099\}$ test edge-tile boundary masks and thread reduction cleanup loops.
• Non-contiguous views and ragged strides: Transposed matrices (x.t()), strided slices, and non-zero storage offsets test pointer arithmetic.
• Adversarial floating-point inputs: Zeros, subnormals, numbers near overflow ($10^{19}$ squaring to infinity in FP32), and exact NaN or infinity placements verify compliance with IEEE 754 standards.

Componentwise numerical acceptance requires a declared formula:

$$|c - r| \le \mathrm{atol} + \mathrm{rtol} \cdot |r|$$

PyTorch's torch.testing.assert_close implements this rule and checks that NaNs match at identical positions.^{[4]Reference 4torch.allclosehttps://docs.pytorch.org/docs/stable/generated/torch.allclose} Choose tolerances based on operator mathematics and reduction length. Reporting maximum absolute error, maximum relative error, and failing element counts prevents an aggregate pass from hiding localized divergence.

The script below audits input preservation, memory aliasing, and non-finite policies:

Checking dense tensor values, input mutation, and pointer aliasing catches bugs before device timing runs.

4. Compute Sanitizer gating

Numerical tests don't reveal undefined behavior that survived by luck. NVIDIA Compute Sanitizer isolates execution hazards into four dedicated tools:^{[5]Reference 5Compute Sanitizerhttps://docs.nvidia.com/compute-sanitizer/ComputeSanitizer/}

• memcheck: Detects out-of-bounds memory accesses, misaligned addresses, and device heap errors.
• racecheck: Intercepts shared memory hazards, including Read-After-Write (RAW), Write-After-Read (WAR), and Write-After-Write (WAW) data races.
• initcheck: Flags reads from uninitialized global or shared device memory.
• synccheck: Identifies divergent synchronization barriers, such as threads within a block reaching different __syncthreads() calls, which deadlocks hardware warps.

Invoke the sanitizer with an explicit non-zero exit code:

Any report of a hazard or misaligned access halts evaluation immediately.

5. Clean CUDA event benchmarking

Accurate GPU timing requires strict methodology. Python host timers (time.time()) measure kernel launch enqueue time rather than device completion.^{[6]Reference 6CUDA semanticshttps://docs.pytorch.org/docs/stable/notes/cuda.html} High-integrity benchmarking enforces five rules:^{[7]Reference 7PyTorch Benchmarkhttps://docs.pytorch.org/tutorials/recipes/recipes/benchmark.html}

• CUDA events on the work stream: Place cudaEventRecord before and after the kernel on the active stream, synchronizing on the end event before reading elapsed time.
• Warmup passes: Run 20 to 50 warmup iterations to bring the GPU core and memory clocks out of idle power states (P-states).
• Interleaved A-B-B-A execution: Alternate between baseline ($A$) and candidate ($B$) executions to cancel out thermal throttling and dynamic frequency variations over time.
• Rotated buffer cache busting: If a benchmark runs the same 16 MB tensor in a tight loop on an H100 GPU (which has 50 MB of L2 cache), the tensor stays resident in L2 cache. The benchmark reports an unrealistic 12 TB/s cache speed instead of measuring HBM memory bandwidth. Cycle through a ring of separate memory buffers to force real global memory traffic.
• Statistical reporting: Record full timing distributions (median, P90, P99) instead of cherry-picking the minimum latency.

6. Roofline sanity check

Compare measured kernel throughput against the theoretical hardware ceiling using the Roofline model.^{[8]Reference 8Nsight Compute Profiling Guidehttps://docs.nvidia.com/nsight-compute/ProfilingGuide/}

For weighted RMSNorm, calculate the arithmetic intensity ($I$):

• Operations: Sum of squares ($H$ multiplies, $H-1$ adds), mean (1 divide), rsqrt (1 op), normalization ($H$ multiplies), and weighting ($H$ multiplies). Total: roughly $4H$ floating-point operations.
• Memory traffic: Read input $x$ ($H$ elements), read weight $w$ ($H$ elements), and write output $y$ ($H$ elements). For FP16 ($2$ bytes/element), total memory traffic is $3 \cdot H \cdot 2 = 6H$ bytes.
• Arithmetic intensity: $I = \frac{4H \text{ FLOPs}}{6H \text{ Bytes}} \approx 0.67 \text{ FLOPs/Byte}$.

Because an H100 SXM5 GPU features an arithmetic intensity ridge point around 150 FLOPs/Byte, RMSNorm is deeply memory bandwidth bound. Achieved bandwidth can't exceed the physical HBM3 limit of 3.35 TB/s:

$$\text{Bandwidth}_{\text{achieved}} = \frac{6 \cdot N \cdot H}{\text{Latency}_{\text{seconds}}}$$

If a benchmark reports an effective bandwidth of 5.2 TB/s on an H100, the kernel isn't fast; it's skipping computation, returning early, or dropping memory writes.

The refinement feedback loop: hardware-guided iteration

Autonomous generation agents (such as KernelAgent) don't rely on blind one-shot generation; they analyze diagnostic traces to iteratively optimize kernel code.^{[9]Reference 9KernelAgent: Hardware-Guided GPU Kernel Optimization via Multi-Agent Orchestrationhttps://pytorch.org/blog/kernelagent-hardware-guided-gpu-kernel-optimization-via-multi-agent-orchestration/} The harness closes the loop by feeding structured telemetry back to the model:

The feedback payload provides three categories of actionable signals:

1. Compiler diagnostics: Capture ptxas warnings, register spill counts, and shared memory footprints. If local memory spills occur, the model responds by reducing unroll factors, hoisting variables, or tiling registers.
2. Sanitizer traces: When racecheck or memcheck fails, extract the instruction pointer, memory address, thread coordinates (threadIdx=(x,y,z)), and block coordinates (blockIdx=(x,y,z)). The agent receives the exact instruction line where a barrier is missing or an address is miscalculated.
3. Warp stall counters: Telemetry from NVIDIA Nsight Compute (NCU) reveals runtime bottlenecks via hardware performance counters:^{[8]Reference 8Nsight Compute Profiling Guidehttps://docs.nvidia.com/nsight-compute/ProfilingGuide/}
   • smsp__warp_issue_stalled_long_scoreboard_pct: Warps waiting for global DRAM loads. Feedback prompts the model to vectorize memory loads using float4 or implement asynchronous copies (cp.async).
   • smsp__warp_issue_stalled_short_scoreboard_pct: Warps stalled on shared memory or MIO operations, signaling shared memory bank conflicts. The model pads shared memory buffers (such as __shared__ float smem[128 + 1]).
   • smsp__warp_issue_stalled_barrier_pct: Warps stalled waiting at __syncthreads(), pointing to thread divergence or poor load balancing. The model switches to warp shuffle intrinsics (__shfl_down_sync) across 32-thread warps.

To prevent prompt poisoning and benchmark memorization, enforce three strict boundaries: sanitize logs to remove raw tensors from hidden fixtures, cap the refinement loop to a maximum of 5 iterations, and hold out an untouched validation suite that runs only once after candidate generation freezes.

The promotion evidence ladder: from lab to production infrastructure

Passing laboratory unit tests is necessary, but it doesn't prove that a kernel will run reliably under messy production traffic. The promotion evidence ladder establishes four progressive validation tiers:

Tier 1: Lab verification

The candidate clears the automated harness: the operator contract passes, ptxas shows zero local memory spills, multi-scale tests verify prime lengths and adversarial floats, Compute Sanitizer reports zero hazards, and benchmark throughput aligns with the Roofline model.

Tier 2: Shadow execution

Deploy the compiled candidate to a production inference node in shadow mode. The inference runtime duplicates real customer requests: the trusted reference kernel computes the user-facing response, while the candidate kernel runs concurrently on an asynchronous stream with zero impact on user latency.

An asynchronous worker compares candidate outputs against reference tensors. Shadow execution audits three critical metrics under production traffic:

• Numerical consistency: Verify that real input distributions (such as varying context lengths and sparsity) don't trigger unexpected precision divergence or NaNs.
• Host and device memory leaks: Monitor virtual memory metrics (VMM) and CUDA caching allocators over millions of invocations.
• JIT compilation cardinality: Ensure the runtime doesn't compile a new binary for every novel sequence length, which exhausts host RAM and GPU memory.

Tier 3: Canary routing

Route a small fraction of live traffic (e.g., 1% ramping to 5%) directly to the candidate kernel behind a dynamic feature flag. Execution is restricted to tested GPU architectures via dispatch predicates.

During canary routing, automated tripwires monitor three failure indicators:

• Tail latency regressions: A sudden spike in P99 or P99.9 latency indicates warp serialization or cache thrashing.
• Driver and hardware errors: Any GPU XID error, CUDA illegal address trap, or hardware engine reset trips the circuit breaker immediately.
• Downstream logit drift: Sample output token distributions to ensure downstream generation remains identical.

If any tripwire triggers, the feature flag flips within milliseconds, routing all traffic back to the trusted reference without dropping requests.

Tier 4: Production gate

Full promotion commits the candidate to the core operator registry. The kernel includes an automated circuit breaker, a permanent reference fallback path, and a cryptographically signed promotion receipt.

Read benchmark scores within their boundary

Research benchmarks measure model capabilities across standardized harnesses, but they don't certify production readiness:

Aggregate benchmark scores can mask localized regressions. A candidate might achieve a 1.5x average speedup across common shapes while silently producing incorrect results on strided views. Always evaluate candidates against your application's exact operator contract and production distribution.

Preserve a promotion receipt

Every promoted kernel requires an immutable, machine-readable promotion receipt that records its complete provenance:

A human reviewer inspects the candidate code, verifies that dispatch predicates match the library specification, reviews licensing, and confirms that the rollback target is tested before signing the receipt.

Promotion rubric

Promotion uses hard gates before evaluating performance scores:

No speedup justifies bypassing a safety or contract gate. Once all hard gates pass, teams can rank candidates by latency, maintainability, and resource utilization.

Mastery check

Evaluation rubric

• Foundational: Distinguish between generated source code, benchmark pass rates, and production promotion criteria.
• Intermediate: Design a multi-scale hidden test suite, Compute Sanitizer verification plan, and Roofline sanity check for a GPU operator.
• Advanced: Implement a closed-loop refinement agent using NCU warp stall counters and manage canary rollouts with sub-millisecond fallbacks.

Diagnose a misleading pass