Kernel Forge: An Agent Harness for LLM-Based CUDA Kernel Generation and Optimisation

Teaser: A July 2026 paper introduces Kernel Forge — an agent harness that generates CUDA kernels from natural language descriptions, then iteratively compiles, profiles, and optimises them through a feedback loop of automated benchmarking and LLM-driven refinement. This article examines the architecture, the agent loop, and what it means for GPU programming accessibility.

Introduction

CUDA kernel programming remains one of the steepest learning curves in software engineering. Writing a correct kernel is achievable with some parallel programming experience; writing an efficient one — one that saturates memory bandwidth, minimises warp divergence, and exploits shared memory — requires years of GPU architecture knowledge and hands-on tuning experience.

Kernel Forge, introduced in a July 2026 paper from ETH Zurich, proposes an LLM agent harness that automates this expertise gap. Given a natural language description of a computation, Kernel Forge:

1. Generates an initial CUDA kernel
2. Compiles it with nvcc and checks for errors
3. Profiles it on the target GPU
4. Analyses performance metrics against hardware roofline models
5. Iteratively refines the kernel until performance converges

The result: across 20 common benchmark operations (vector addition, matrix multiplication, convolution, reduction, scan, stencil, sorting), Kernel Forge produces kernels that achieve 80% of hand-optimised library performance (cuBLAS, CUTLASS) within 5–7 refinement iterations.

The Agent Harness Architecture

Kernel Forge is structured as a feedback-driven agent loop:

┌─────────────────────────────────────────────────────────────┐
│                    Natural Language Prompt                   │
│  "Write a CUDA kernel that performs a 2D convolution with   │
│   a 5x5 filter on a 4096x4096 input"                        │
└─────────────────────┬───────────────────────────────────────┘
                      │
┌─────────────────────▼───────────────────────────────────────┐
│  1. LLM Generator Agent                                     │
│     • Parses specification → kernel skeleton                │
│     • Selects template (shared memory, tiling, etc.)        │
│     • Generates initial .cu file                            │
└─────────────────────┬───────────────────────────────────────┘
                      │
┌─────────────────────▼───────────────────────────────────────┐
│  2. Compilation Agent                                       │
│     • Invokes nvcc with target architecture flags           │
│     • Parses compilation errors → structured feedback       │
│     • If errors: return to Generator with error context      │
└─────────────────────┬───────────────────────────────────────┘
                      │
┌─────────────────────▼───────────────────────────────────────┐
│  3. Profiling Agent                                         │
│     • Launches kernel with representative input sizes       │
│     • Captures: runtime, memory throughput, SM occupancy    │
│     • Computes roofline metrics (FLOP/s, util %)           │
└─────────────────────┬───────────────────────────────────────┘
                      │
┌─────────────────────▼───────────────────────────────────────┐
│  4. Analysis Agent                                          │
│     • Compares against roofline model                       │
│     • Identifies bottlenecks (compute-bound vs mem-bound)   │
│     • Suggests optimisation strategies                      │
│     • Annotates kernel source with bottleneck regions       │
└─────────────────────┬───────────────────────────────────────┘
                      │
                      ▼
            ┌──────────────────┐
            │ Converged?       │
            │ (Δ perf < 5%)    │──── Iterate (max 10 rounds)
            └──────┬───────────┘
                   │ No
                   │ (back to Generator with analysis feedback)
                   │
            ┌──────▼──────────────────────────────────────────┐
            │  5. Optimiser Agent                              │
            │     • Applies targeted transformations:          │
            │       - Tiling / shared memory allocation        │
            │       - Loop unrolling / coalescing              │
            │       - Warp-level reduction                     │
            │       - Register pressure balancing             │
            │     • Regenerates kernel with optimisations      │
            └──────────────────────────────────────────────────┘
                   │
                   ▼
            ┌──────────────────┐
            │ Final kernel     │ ← When converged or max iterations reached
            │ + optimisation   │
            │ report           │
            └──────────────────┘

The feedback loop is the core innovation. Rather than asking the LLM to write a perfect kernel in one shot (which rarely succeeds for complex operations), Kernel Forge treats optimisation as an empirical process: try, measure, analyse, refine.

The Feedback Loop in Detail

Round 1: Initial Generation

Given the prompt "2D convolution with 5×5 filter on 4096×4094 float input", the Generator produces a naive kernel:

__global__ void conv2d(const float* input, float* output,
                       const float* filter, int width, int height) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;

    float sum = 0.0f;
    int half = 5 / 2;

    for (int fy = -half; fy <= half; fy++) {
        for (int fx = -half; fx <= half; fx++) {
            int ix = x + fx;
            int iy = y + fy;
            if (ix >= 0 && ix < width && iy >= 0 && iy < height) {
                sum += input[iy * width + ix] * filter[(fy + half) * 5 + (fx + half)];
            }
        }
    }
    output[y * width + x] = sum;
}

Round 1: Profile Results

Kernel: conv2d_naive
Runtime: 184.2 ms
Memory throughput: 127 GB/s (22% of H100 peak)
SM occupancy: 34%
Main bottleneck: Global memory bandwidth (no shared memory usage)

Round 2: Analysis → Optimisation

The Analysis agent identifies that each input pixel is loaded 25 times (once per filter tap). The Optimiser generates a tiled version with shared memory:

__global__ void conv2d_tiled(const float* input, float* output,
                              const float* filter, int width, int height) {
    __shared__ float tile[TILE_SIZE + 4][TILE_SIZE + 4];  // Halo region

    int tx = threadIdx.x, ty = threadIdx.y;
    int x = blockIdx.x * TILE_SIZE + tx - 2;  // Halo offset
    int y = blockIdx.y * TILE_SIZE + ty - 2;

    // Cooperative load with halo
    if (x >= 0 && x < width && y >= 0 && y < height) {
        tile[ty][tx] = input[y * width + x];
    } else {
        tile[ty][tx] = 0.0f;
    }
    __syncthreads();

    // Compute only for interior region
    if (tx >= 2 && tx < TILE_SIZE + 2 && ty >= 2 && ty < TILE_SIZE + 2
        && x < width && y < height) {
        float sum = 0.0f;
        #pragma unroll
        for (int fy = 0; fy < 5; fy++) {
            #pragma unroll
            for (int fx = 0; fx < 5; fx++) {
                sum += tile[ty + fy - 2][tx + fx - 2] * filter[fy * 5 + fx];
            }
        }
        output[y * width + x] = sum;
    }
}

Round 2: Profile Results

Kernel: conv2d_tiled
Runtime: 47.3 ms
Memory throughput: 493 GB/s (85% of H100 peak)
SM occupancy: 72%
Improvement: 3.9× over naive
Remaining bottleneck: Boundary condition branches causing some warp divergence

The loop continues until convergence, typically reaching 5–7 rounds for complex kernels.

Empirical Results

Kernel Forge was evaluated on an NVIDIA H100 GPU across 20 benchmarks spanning four categories:

Category	Operations	Hand-Optimised Performance	Kernel Forge	% of Hand-Opt
Memory-bound	copy, add, scale, triad, axpy, gather, scatter	cuBLAS / hand-written	96–102%	98%
Compute-bound (dense)	sgemm, dgemm, batch matmul	cuBLAS	74–83%	78%
Stencil / convolution	1D stencil, 2D conv 3×3/5×5/7×7, separable conv	CUTLASS / hand-written	71–88%	81%
Reduction / scan	sum, max, min, prefix sum, argmax	CUB / Thrust	68–79%	74%
Sorting	radix sort (32-bit), merge sort	CUB	52–61%	57%

The weakest category is sorting, where the agent consistently struggles with the complex warp-level primitives that make CUB's radix sort state-of-the-art. The paper notes this as a known limitation — sorting requires algorithmic innovation more than parameter tuning.

Iteration Behaviour

The paper tracked performance across iterations for representative kernels:

graph LR
    subgraph conv2d_5x5[2D Convolution 5×5]
        direction LR
        I1[1: 184ms] --> I2[2: 47ms]
        I2 --> I3[3: 31ms]
        I3 --> I4[4: 28ms]
        I4 --> I5[5: 27ms ✓]
    end

    subgraph sgemm[SGEMM 2048×2048]
        direction LR
        S1[1: 342ms] --> S2[2: 98ms]
        S2 --> S3[3: 52ms]
        S3 --> S4[4: 44ms]
        S4 --> S5[5: 41ms]
        S5 --> S6[6: 39ms ✓]
    end

    subgraph reduce[Reduction sum 16M]
        direction LR
        R1[1: 12.4ms] --> R2[2: 4.1ms]
        R2 --> R3[3: 2.8ms]
        R3 --> R4[4: 2.4ms ✓]
    end

    classDef fast fill:#54A24B,stroke:#3a7a35,color:#fff
    classDef slow fill:#E45756,stroke:#b33d3d,color:#fff
    classDef mid fill:#F58518,stroke:#b35a0e,color:#fff
    class I1,S1,R1 slow
    class I2,I3,S2,S3,R2 mid
    class I4,I5,S4,S5,S6,R3,R4 fast

Convergence typically occurs within 5 iterations. The largest gains always come in the first 2–3 rounds, where obvious performance bugs (no shared memory, no tiling, no vectorisation) are fixed.

What Makes a Good CUDA Kernel Generator?

The paper analysed which LLM capabilities correlate most strongly with Kernel Forge performance:

Capability	Impact on Final Kernel Performance	Notes
GPU architecture knowledge	High	Understanding of H100 SM layout, memory hierarchy, warp scheduling
PTX/assembly literacy	Medium	Helps interpret profiler output (achieved occupancy vs theoretical)
Roofline analysis understanding	High	Determines whether to optimise compute or memory path
Shared memory tiling	Very high	Single biggest performance lever across all benchmarks
Warp-level primitives	Medium	Needed for reductions and scans; current weakest area
Error message parsing	Medium	The Compilation agent must distinguish real errors from false positives

The Generator agent uses Claude 4 Sonnet and GPT-5.6 Sol as backend LLMs; the paper reports no significant difference between them for the Generator role, but the Analysis agent benefits from GPT-5.6 Sol's longer context window (256K tokens) when processing profiler output.

Practical Applications

Custom Kernel Generation for ML

Kernel Forge is most immediately useful for ML researchers who need custom fused kernels for novel operations:

"I need a fused kernel that does a group-query attention with ALiBi positional encoding, with variable group sizes per head."

Instead of hand-writing a CUDA kernel or waiting for a library update, Kernel Forge can generate and optimise a working kernel in minutes.

Automated Kernel Porting

Kernel Forge can take an existing CUDA kernel (or an OpenCL kernel) and re-optimise it for a different GPU architecture — e.g., A100 → H100 or H100 → B200 — by changing the target architecture flag and letting the profiling loop tune parameters for the new hardware.

Educational Use

For students learning CUDA, Kernel Forge's optimisation reports are a teaching tool in themselves. Each iteration shows exactly what change was made and why, with before/after profiling data:

Round 2: Added shared memory tiling (16×16 tiles with 2-element halo). Reduced global memory loads from 25× per output element to 1×. Performance improved 3.9×.

Limitations

Limitation	Details
Sorting performance gap	52–61% of CUB. The agent cannot replicate hand-tuned warp-level bitonic sort primitives.
Single-GPU only	No multi-GPU or multi-node kernel generation.
Fixed precision	FP32 kernels only. FP16/FP8/INT8 tensor core utilisation not yet supported.
nvcc dependency	Requires a local CUDA toolchain with nvcc and nsys/nsight-compute.
Convergence guarantee	Not guaranteed. 4.3% of runs diverged (performance worsened) and required rollback.

Conclusion

Kernel Forge demonstrates that LLM-based code generation, when coupled with a tight empirical feedback loop of compile → profile → analyse → refine, can produce CUDA kernels approaching hand-optimised performance for a wide range of operations. The key insight is not that LLMs write perfect kernels — they don't — but that they can navigate the optimisation search space more efficiently than either a human starting from scratch or an autotuner starting from random configurations.

For the GPU programming community, Kernel Forge suggests a future where the programmer describes what computation they want, and the agent system handles how to map it efficiently onto the hardware. The 80% figure is not a ceiling — it is a baseline that will improve with better LLMs, richer profiling feedback, and larger optimisation repertoires.

The paper and code are available on arXiv and GitHub (July 2026).