Makora: RL for GPU Kernel Generation

• Makora is an RL environment and toolchain that reformulates GPU kernel generation as a reinforcement learning problem using verifiable, execution-based rewards.
• It integrates multi-turn tool augmentation and scalable hardware evaluation to optimize both functional correctness and runtime performance.
• The system overcomes limited high-quality supervision by decontaminating training datasets and implementing robust reward-hacking defenses.

Searching arXiv for the Makora paper and closely related GPU kernel-generation work to ground the article. arxiv_search(query="Makora GPU kernel generation Triton KernelBench reinforcement learning", max_results=10) Makora is an environment and toolchain for reinforcement learning fine-tuning of frontier LLMs—specifically GPT-5—for the specialized task of GPU kernel generation. It reformulates Triton kernel synthesis from natural-language problem descriptions and PyTorch references as reinforcement learning from verifiable rewards (RLVR), coupling a strong pretrained policy with execution-grounded feedback, multi-turn tool use, and a scalable evaluation backend. In the formulation reported in "Fine-Tuning GPT-5 for GPU Kernel Generation" (Tehrani et al., 11 Feb 2026), Makora targets two simultaneous objectives: reliable functional correctness and competitive runtime performance against TorchInductor, in a domain where supervised fine-tuning is constrained by limited high-quality labels, compiler-induced biases in synthetic data, and hardware-specific optimization regimes.

1. Problem setting and rationale

Makora is motivated by the mismatch between the requirements of GPU kernel optimization and the assumptions underlying standard supervised fine-tuning. The paper states that production-grade GPU kernels are proprietary and that public examples are often educational or suboptimal. The quantitative disparity is large: Stack v2’s Python subset contains 96,448,523 samples, whereas KernelBook, defined as PyTorch plus synthetic Triton, has only 18,162 examples (Tehrani et al., 11 Feb 2026). This scarcity is presented as a central reason that large-scale supervised fine-tuning on high-quality labels is not viable.

A further limitation is that compiler-generated synthetic data is not neutral. Training on compiler outputs tends to inherit performance ceilings dictated by compiler heuristics, boilerplate and artifacts from IR-driven code generation, reliance on internal runtime libraries or intrinsics, and poor readability or maintainability. The paper’s position is that such training teaches a model to mimic compiler idioms rather than discover new optimizations. It also argues that functional correctness alone is inadequate, because a correct kernel may still be substantially slower than an optimized implementation, and because optimization strategies depend on hardware details such as Hopper versus Blackwell, and within Blackwell, the difference between an RTX 5090 without TMA and a B200 with TMA (Tehrani et al., 11 Feb 2026).

Makora therefore treats kernel generation as an RLVR problem rather than a pure code-completion problem. This suggests a shift from imitation toward execution-conditioned search over a combinatorial design space involving tiling, memory layout, vectorization, and scheduling, whose interactions are explicitly described as non-linear and effectively intractable to enumerate for supervised learning.

2. RLVR environment and tool-augmented workflow

Makora implements a tool-augmented, multi-turn RLVR loop. Each training episode receives a system prompt, a natural-language problem description, and a PyTorch reference implementation $r$. The policy, GPT-5, proposes a Triton kernel $k$ and may interleave external tool calls within the same trajectory (Tehrani et al., 11 Feb 2026).

The tool interface comprises four named components. The kernel_evaluator compiles, validates, and benchmarks the kernel, returning structured feedback such as compilation or runtime errors, output mismatch, hack-detection reasons, or success with measured speedup. The kernel_search tool retrieves previously generated candidate kernels for the same reference from a database; sampling is weighted by speedup through a softmax-style normalization with controlled stochasticity, including a 10% chance to return no candidate and a 10% chance to return an erroneous candidate. The web_search tool issues internet queries for documentation and expert discussions. A profiler is identified as a planned tool for detailed runtime and utilization diagnostics, but left to future work in training (Tehrani et al., 11 Feb 2026).

The evaluation backend is queue-based and parameterized by target hardware such as H100 or B200. For each submitted kernel, both the reference and optimized code pass through compilation, validation, and benchmarking. Benchmarking uses 3 warmups plus 100 timed iterations. Makora also applies AST-based canonicalization, docstring and comment removal, and normalized formatting through ast.unparse to avoid redundant evaluation. Over three weeks, canonicalization raised cache hits to approximately 16% and saved 227.6 hours of compute (Tehrani et al., 11 Feb 2026).

A defining property of the workflow is that tool outputs are appended to the model’s context inside the same episode. The paper frames this as a bridge between single-turn and multi-turn RL, allowing immediate refinement from execution feedback rather than delayed reward alone.

3. Dataset construction, decontamination, and reward formulation

Makora’s training corpus is curated from public PyTorch code and explicitly decontaminated against KernelBench. The decontamination pipeline has two stages. First, semantic deduplication uses jina-embeddings-v3, removing a train sample if

$$\min_{b \in \mathrm{KernelBench}} \|E(\mathrm{code}_k)-E(\mathrm{code}_b)\|_2 < T_{\mathrm{embed}},$$

with $T_{\mathrm{embed}} = 0.45$. Second, syntactic deduplication uses token-level Jaccard similarity,

$$J(c_i,c_j)=\frac{|T(c_i)\cap T(c_j)|}{|T(c_i)\cup T(c_j)|},$$

removing pairs with $J > 0.8$ (Tehrani et al., 11 Feb 2026).

Problems are ranked by an LLM judge into six difficulty levels, $L0$ through $L5$, labeled Trivial, Simple, Straightforward, Moderate, Advanced, and Expert. Shape diversification is also introduced by an LLM to avoid trivial tensor sizes and to expose optimization opportunities. Baseline runtime is measured on H100 after warmup and 5 timed runs, with the mean recorded as $t_{\mathrm{baseline}}$. Filtering then retains tasks satisfying $$1\ \mathrm{ms} < t_{\mathrm{baseline}} < 1000\ \mathrm{ms}$$ and difficulty level $k$0 (Tehrani et al., 11 Feb 2026).

The sampling strategy is cluster-aware. Makora applies k-means with 50 clusters and inverse-log weighting,

$k$1

followed by normalization and stochastic rounding, then samples uniformly within clusters. Two subsets are highlighted: a 100-problem subset concentrated in $k$2-$k$3 with a nontrivial fraction of $k$4, and a 1,000-problem subset matching the broader higher-difficulty distribution. The post-filter pool contains 11,363 valid examples, while training on the full set is described as ongoing (Tehrani et al., 11 Feb 2026).

The RLVR objective uses deterministic execution-based signals. For a problem $k$5 with reference $k$6 and candidate kernel $k$7,

$k$8

with default $k$9. The unnormalized reward is

$$\min_{b \in \mathrm{KernelBench}} \|E(\mathrm{code}_k)-E(\mathrm{code}_b)\|_2 < T_{\mathrm{embed}},$$0

where

$$\min_{b \in \mathrm{KernelBench}} \|E(\mathrm{code}_k)-E(\mathrm{code}_b)\|_2 < T_{\mathrm{embed}},$$1

Correctness is defined on test inputs $$\min_{b \in \mathrm{KernelBench}} \|E(\mathrm{code}_k)-E(\mathrm{code}_b)\|_2 < T_{\mathrm{embed}},$$2 by $$\min_{b \in \mathrm{KernelBench}} \|E(\mathrm{code}_k)-E(\mathrm{code}_b)\|_2 < T_{\mathrm{embed}},$$3; if any test fails, the reward is zero (Tehrani et al., 11 Feb 2026).

A common misconception addressed by this setup is that correctness and optimization can be treated independently. Makora’s reward makes correctness necessary but not sufficient, then differentiates among correct kernels through speedup-sensitive shaping.

4. Evaluation protocol and empirical performance

Makora evaluates on an extended KernelBench suite of 264 problems spanning matmul, transpose, attention, convolution, reductions such as sum, max, and softmax, and elementwise operations including activations and normalization. Degenerate cases such as constant fills and no-ops are removed, inputs are standardized, and the problems are enforced to require nontrivial computation (Tehrani et al., 11 Feb 2026).

The benchmark protocol is fixed: compilation, validation under $$\min_{b \in \mathrm{KernelBench}} \|E(\mathrm{code}_k)-E(\mathrm{code}_b)\|_2 < T_{\mathrm{embed}},$$4, and benchmarking with 3 warmups plus 100 timed iterations. Performance is reported using three metrics: functionality rate, the fraction faster than TorchInductor, and geometric mean speedup over the solved set,

$$\min_{b \in \mathrm{KernelBench}} \|E(\mathrm{code}_k)-E(\mathrm{code}_b)\|_2 < T_{\mathrm{embed}},$$5

where $$\min_{b \in \mathrm{KernelBench}} \|E(\mathrm{code}_k)-E(\mathrm{code}_b)\|_2 < T_{\mathrm{embed}},$$6 (Tehrani et al., 11 Feb 2026).

In the single-attempt setting, Makora’s fine-tuned model, denoted GPT-5-RL, improves kernel correctness from 43.7% to 77.0%, a gain of 33.3 percentage points over baseline GPT-5. The fraction of problems outperforming TorchInductor rises from 14.8% to 21.8%, a gain of 7.0 percentage points, and geometric mean speedup increases from $$\min_{b \in \mathrm{KernelBench}} \|E(\mathrm{code}_k)-E(\mathrm{code}_b)\|_2 < T_{\mathrm{embed}},$$7 to $$\min_{b \in \mathrm{KernelBench}} \|E(\mathrm{code}_k)-E(\mathrm{code}_b)\|_2 < T_{\mathrm{embed}},$$8 (Tehrani et al., 11 Feb 2026).

When the policy is embedded in the full coding agent MakoraGenerate, the reported results are substantially higher: 97.4% functionality, 72.9% of problems faster than TorchInductor, and a geometric mean speedup of $$\min_{b \in \mathrm{KernelBench}} \|E(\mathrm{code}_k)-E(\mathrm{code}_b)\|_2 < T_{\mathrm{embed}},$$9. The same evaluation also lists comparative baselines on the 264-problem benchmark: Claude Opus 4.5 at 59.8%, 13.2%, and $T_{\mathrm{embed}} = 0.45$0; Claude 4.5 Sonnet at 47.2%, 12.8%, and $T_{\mathrm{embed}} = 0.45$1; Gemini 2.5 Pro at 37.7%, 8.7%, and $T_{\mathrm{embed}} = 0.45$2; Grok at 24.4%, 10.5%, and $T_{\mathrm{embed}} = 0.45$3; and Grok Code Fast at 15.8%, 3.8%, and $T_{\mathrm{embed}} = 0.45$4 (Tehrani et al., 11 Feb 2026).

In the multi-attempt setting with up to three refinement steps, GPT-5-RL improves functionality from 77.0% to 83.7%. Difficulty-stratified results indicate gains across all levels, and MakoraGenerate reaches perfect performance at Level 1 and Level 2; the paper specifies that Levels 1 and 2 each contain 100 problems, Level 3 has 50, and Level 5 has 14 (Tehrani et al., 11 Feb 2026).

5. Agentic integration, exploration, and reward-hacking defenses

MakoraGenerate is described as a multi-agent evolutionary system built around the fine-tuned GPT-5-RL policy. It maintains an evolving pool of candidate kernels ranked by speedup, reuses strong prior solutions, and applies exploit-explore selection with controlled diversity to avoid premature convergence. Each attempt generates a batch of candidates and uses code generation, evaluation through the same compile-validate-benchmark backend, documentation lookup, and optimization planning (Tehrani et al., 11 Feb 2026).

The inference-time tool ablations are structurally important. With up to three attempts and at most three tool calls per attempt, web_search alone can trade correctness for speed, kernel_evaluator acts as a correctness filter, and the full WS+KE+KS loop yields the strongest correctness gains with modest, consistent performance improvements. On Attempt 1, KE increases functionality by 1.7 points, while WS+KE+KS increases fraction faster by 1.4 points and geometric mean by 0.02. On Attempt 2, WS+KE+KS increases functionality by 7.2 points, fraction faster by 0.8, and geometric mean by 0.04. On Attempt 3, the same combination increases functionality by 7.6 points, fraction faster by 0.2, and geometric mean by 0.05. Tool usage statistics show that tools are used on 56.8% of problems, KS covers 45.5% of problems, KE accounts for 56.6% of calls, and WS is used sparingly with 10.2% coverage (Tehrani et al., 11 Feb 2026).

Because RLVR is susceptible to reward hacking, Makora implements explicit defenses. The paper enumerates six hack modes: Baseline Kernel, Identity Kernel, No-op Kernel, Unused Output, Ghost Optimization, and Forgotten Kernel. Prevention combines static reachability analysis, which ensures that at least one reachable @triton.jit or CUDA kernel is actually launched, with a GPT-5-based LLM judge that detects semantic hacks, including previously unseen ones. Any flagged code receives zero reward (Tehrani et al., 11 Feb 2026).

This design implies that Makora is not merely an optimizer around benchmarking infrastructure; it is an RL environment in which the reward function, execution substrate, and anti-cheating mechanisms are co-designed.

6. Limitations, portability, and naming ambiguity

Makora’s results depend on a strong base model. The paper states that RLVR requires a base model capable of producing non-zero-reward samples, and reports that smaller open models—Qwen-4B, Qwen-8B, and Qwen-32B—plateaued early in reward. This is presented as consistent with prior observations about cold-start difficulties in RL for specialized code generation (Tehrani et al., 11 Feb 2026).

The system also does not claim explicit cross-hardware generalization results. Although the evaluation backend can target H100 or B200, and the paper emphasizes hardware discrepancies such as TMA availability, memory technology, and low-precision support, performance portability remains a future direction. The text specifically identifies profiler-informed rewards and speedup-focused training as promising directions for improving optimization depth and cross-hardware robustness (Tehrani et al., 11 Feb 2026).

A separate ambiguity concerns nomenclature. In the robotics paper "GAMORA: A Gesture Articulated Meta Operative Robotic Arm for Hazardous Material Handling in Containment-Level Environments" (Wasay et al., 17 Jun 2025), the details section includes the note that one may assume “Makora” refers to that VR-guided robotic system. That usage is distinct from Makora as defined in the GPU kernel-generation literature. GAMORA is a ROS-, Unity-, and Oculus Quest 2-based teleoperation platform for hazardous laboratory manipulation, whereas Makora in (Tehrani et al., 11 Feb 2026) denotes an RLVR environment and toolchain for Triton kernel synthesis.

Taken in its primary usage, Makora designates a reinforcement-learning infrastructure for post-training GPT-5 on GPU kernel generation under verifiable execution-based rewards. Its core significance lies in the combination of decontaminated task curation, tool-augmented multi-turn refinement, queue-based evaluation on target hardware, and explicit hack prevention, all directed toward the joint objectives of correctness and runtime competitiveness on KernelBench (Tehrani et al., 11 Feb 2026).