Pith. sign in

REVIEW 3 major objections 5 minor 40 references

Splitting an LLM kernel-writing task into verified subkernels yields more correct, faster GPU code than generating the whole kernel at once, while consuming up to 35 percent fewer LLM tokens.

Reviewed by Pith at T0; open to challenge. T0 means a machine referee read the full paper against a public rubric. the ladder, T0–T4 →

Challenge this review Re-run · record.json Download PDF Read on arXiv ↗

T0 review · deepseek-v4-flash

2026-08-02 10:33 UTC pith:UM43MBLI

load-bearing objection Solid, useful Triton-kernel generation work with a real novelty claim, but the unreported TotalFuse fallback undermines the modularity attribution; worth refereeing after artifact release and added diagnostics. the 3 major comments →

arxiv 2607.20501 v1 pith:UM43MBLI submitted 2026-06-20 cs.AI cs.MA

MKEvolve: A Modular Multi-Agent Framework for Kernel Code Generation

Jason Yoo , Rajarshi Saha , Shaowei Zhu , Tao Yu , Wei Tang , Youngsuk Park This is my paper
classification cs.AI cs.MA
keywords LLM code generationGPU kernelsTritonmodular decompositionbeam searchmulti-agent frameworkkernel correctnesscompositional verification
verification ladder T0 review T1 audit T2 compute T3 formal T4 reserved

The pith

A machine-rendered reading of the paper's core claim, the machinery that carries it, and where it could break.

This paper argues that the right unit for LLM-driven kernel synthesis is the subkernel, not the whole kernel. It introduces a framework that repeatedly decomposes a PyTorch module into submodules, independently improves a Triton kernel for each via beam search, and assembles the verified pieces into a top-level program. Across 150 GPU-kernel tasks spanning operator sequences and full models, this modular regime raises correctness and speedups over end-to-end beam search and parallel scaling, on both a frontier model and an open-weight model, with 15-35 percent fewer tokens. The paper also shows the composed kernels are swappable and transferable: adapting a full model to a sibling architecture costs about four LLM calls. If the local-verification premise holds, modularity becomes an organizing principle for inference-time kernel generation.

Core claim

The central claim, stated in the experimental section, is that MKEvolve, through its modular subkernel synthesis strategy, outperforms end-to-end baselines in correctness and speed on both the simpler operator-sequence tasks and the harder full-model tasks, while using fewer LLM tokens. The framework co-evolves two things in lockstep: the decomposition of the target module (via LLM agents that split failing subproblems and optionally fuse succeeded ones) and the per-subkernel Triton implementations (via independent beam search scored by speedup over the reference compiler, with zero reward for incorrect kernels). The end product is a top-level module that programmatically chains subkernels,

What carries the argument

The load-bearing mechanism is the co-evolution loop: an LLM decomposition agent first partitions the module into subproblems and a top-level orchestrator; every few iterations a topology-refinement step splits subproblems whose kernels still fail and, if enabled, fuses subproblems whose kernels all pass; a budget allocator then distributes LLM calls among subproblems, prioritizing failing or new ones; and per-subproblem beam search produces the best verified subkernel. Composition is programmatic, not LLM-generated, which is what makes the pieces independently verifiable, swappable, and reusable.

Load-bearing premise

The claim leans on the assumption that a kernel assembled from subkernels, each verified independently at atol/rtol=1e-4, is itself correct end-to-end at the same tolerance — a local-to-global transfer the paper invokes 'under certain assumptions' without a proven theorem, while its own appendix shows fixed end-to-end thresholds can miss bugs that per-stage checks catch.

What would settle it

Run the same pipeline on tasks into which a known subtle bug is injected (for example, a BF16 accumulator inside an attention QK^T product) and test whether every subkernel passes its local 1e-4 check while the composed kernel's output differs from the reference by more than 1e-4. A cheaper audit: recompute all reported correctness scores with the end-to-end tolerance tightened to 1e-5 and check whether MKEvolve's lead over end-to-end beam search shrinks or reverses.

Watch this falsifier — get emailed when new claim-graph text bears on it.

If this is right

  • Higher correctness and speed with fewer tokens: on 150 GPU-kernel tasks, the modular pipeline beats end-to-end beam search on correctness and most speed thresholds with 15-35 percent fewer LLM tokens, across two different base models.
  • Interpretable failures: because each subkernel is verified separately, a wrong or slow piece can be localized to a specific subproblem instead of debugging the whole kernel.
  • Cheap adaptation: replacing one verified subkernel produces working kernels for sibling architectures with only a handful of LLM calls, as shown with three pooling variants of one full model.
  • Complementary to retraining: the approach works purely at inference time using LLM APIs, so it can be layered on top of post-trained models without weight access.
  • Swapping in simpler implementations for lagging subkernels lifts end-to-end speedups further, and post-hoc fusion recovers some fusion gains at the cost of tokens.

Where Pith is reading between the lines

These are editorial extensions of the paper, not claims the author makes directly.

  • If per-subkernel verification is sound, the framework turns kernel synthesis into a compounding library: verified subkernels (attention blocks, fused linear+activation, pooling) could be reused across models, amortizing LLM cost across projects — a consequence the paper only gestures at via the transfer example.
  • The token savings suggest that search effort concentrates where it matters; a testable prediction is that MKEvolve's advantage over monolithic beam search grows with task size and with the number of independent stages, since the monolithic baseline must re-touch all of a large kernel each iteration.
  • The appendix's feasibility-gap analysis implies a recipe beyond what the paper runs: using per-stage calibrated tolerances (max plus 3 sigma from correct runs) instead of a fixed 1e-4 should make lower-precision kernel generation viable; the paper stops short of using that in its main experiments.
  • The decomposition inherits a hidden sensitivity: the LLM must preserve the original module's random-initialization order when constructing submodules, otherwise weights differ and the functional-equivalence check fails — an engineering constraint that could be relaxed by capturing and replaying RNG states.

Editorial analysis

A structured set of objections, weighed in public.

Desk editor's note, referee report, simulated authors' rebuttal, and a circularity audit.

Referee Report

3 major / 5 minor

Summary. The paper presents MKEvolve, an agentic framework that decomposes a PyTorch module into subproblems, synthesizes and beam-searches Triton kernels for each subproblem with LLM agents, and programmatically composes the resulting subkernels. The decomposition is periodically revised by split/fuse agents, and correctness is checked locally and end-to-end with a TritonRL-based cheating detector. Experiments on 100 KernelBench L2 and 50 L3 problems with two LLMs report higher Correct and Fast@p fractions than Parallel Scaling, Beam Search, and a KernelFalcon baseline, while using 15–35% fewer tokens than Beam Search. Appendices add a second seed, an ablation, a transferability study, and a bug-detection case study.

Significance. If the results hold, the paper makes a useful contribution: it provides evidence that structuring kernel generation around independently verified subkernels can be more token-efficient and more interpretable than monolithic beam search, and it evaluates on a standard external benchmark with two different LLMs. Strengths include the controlled equal-kernel-budget comparison against Beam Search and Parallel Scaling, a second seed in Appendix C, and a concrete appendix (D) showing where fixed tolerance checking fails. The central quantitative claim is plausible, but two load-bearing gaps—the unreported TotalFuse fallback and the unproven global-error statement—need to be closed before the modularity attribution can be accepted.

major comments (3)
  1. [Algorithm 2, lines 8–10] This is the load-bearing issue. The paper's headline claim in Section 4 is that MKEvolve 'through its modular subkernel synthesis strategy' outperforms baselines. However, Algorithm 2 silently replaces P,S with TotalFuse(P)—a monolithic end-to-end search—any time local and global correctness checks disagree. The manuscript never reports how often this trigger fires over the 150 KernelBench tasks or for either LLM. If it fires frequently, especially on L3, the reported Correct/Fast/token numbers may be produced by the fallback plus decomposition overhead, not by modular synthesis. Please report the trigger rate per benchmark and model, and report the main metrics separately for tasks that actually used the modular path; if the fallback is frequent, re-state the contribution accordingly.
  2. [§3 last paragraph; Appendix E; Appendix D Tables 12–13] The compositional correctness guarantee is asserted but not established. Section 3 says 'under certain assumptions, the global model error can be bounded by the local errors of individual subkernels' and points to Appendix E. Appendix E only states that a bound is 'possible' via probabilistic backward error analysis; no theorem, explicit assumptions, or Lipschitz constants are given. Appendix D itself shows that fixed per-stage or end-to-end tolerances can both reject correct outputs and miss bugs in composed operators, and two injected BF16 bugs are undetectable by either method. Thus local subkernel verification at atol/rtol=1e-4 does not, by itself, guarantee end-to-end FP32 correctness. Please either provide a concrete theorem/bound for the FP32 setting used in the main experiments or weaken the wording from 'ensuring correctness' to 'empirical evidence on KernelBench'.
  3. [Tables 1–2 and Appendix C] The empirical comparisons are reported as single-run proportions with no confidence intervals or error bars. Appendix C provides only one additional seed and only for Beam Search, MKEvolve, and MKEvolve (Swap). Given that several headline differences are small (e.g., Table 1 L2 Fast@1: 0.36 vs 0.49; Table 2 L3 Fast@1: 0.04 vs 0.06), sampling noise could affect the ordering. Please provide per-problem results with bootstrap CIs, or at minimum report all seeds and per-problem speedups. This is necessary to support the numeric precision of the 'up to 35% token reduction' and speedup claims.
minor comments (5)
  1. [§4, Setup] The setup lists a metric 'the proportion of which it produces a kernel that is faster than all other baselines,' but this metric is not reported in Tables 1–2. Clarify how this is computed or remove it from the setup.
  2. [Figure 4] The win-rate heatmaps are visually cluttered and the 'vs Beam Search / vs MKEvolve' labels are ambiguous. Consider using a clearer matrix format with confidence intervals or counts.
  3. [Appendix A, KernelFalcon Setup] The KernelFalcon baseline is modified substantially (FP32 conversion, retry logic, serialized subproblem refinement, and external-LLM conversion of the wrapper). These modifications could affect fairness. Please state this limitation in the main text and, if possible, provide the exact patch used.
  4. [Appendix D, Tables 7 and 8] There are formatting typos in the numeric entries, e.g., 'Matmul512 2 5×10−5–7×10−5' in Table 7. Please clean up the table formatting.
  5. [Appendix E] The title 'Local Smoothness Implies Global Smoothness' promises a theorem that is not actually provided. Consider renaming the appendix to reflect the heuristic argument, or include a precise statement with assumptions.

Circularity Check

0 steps flagged

No significant circularity: MKEvolve's headline correctness/speed/token claims are measured against the external KernelBench suite and torch.compile; no fitted quantity is renamed as a prediction.

full rationale

The paper's central empirical claim is an engineering comparison, not a derived prediction. Correctness, Fast@p, and token counts are all measured directly against KernelBench's external CUDA-stream pipeline and torch.compile baselines, with the same prompts and evaluators applied to MKEvolve and the beam-search/parallel-scaling baselines. No parameter is fitted to the reported outcomes and then 'predicted' back. The one self-citation is the TritonRL cheating detector (Woo et al., 2026), reused from prior work by two co-authors; it is applied uniformly to all methods and is not the source of the modularity or token-efficiency claim, so it is a minor, non-load-bearing reuse rather than a circular step. Two robustness gaps flagged in the manuscript are real but are not circularity: Algorithm 2's TotalFuse fallback (lines 8-10) can silently revert to monolithic search, and Appendix E promises only 'under certain assumptions' a global error bound without stating a theorem; Appendix D even shows fixed local thresholds can miss bugs. These weaken the causal attribution 'through its modular subkernel synthesis strategy' and the strength of the correctness guarantee, but they do not make any prediction equal to its input by construction. The empirical wins are plausible and externally benchmarked, so the circularity score is 0.

Axiom & Free-Parameter Ledger

5 free parameters · 5 axioms · 1 invented entities

The central empirical claim rests on benchmark-domain assumptions and algorithm hyperparameters, not on fitted physical constants. The most fragile assumptions are the sufficiency of per-subkernel 1e-4 tolerance for global correctness and the reliability of the adopted cheating detector.

free parameters (5)
  • Refinement period L = 2
    Algorithm 2 re-decomposes every L iterations; chosen by the authors with no sensitivity analysis reported.
  • Outer iterations T = 5
    Algorithm 1 stopping point; chosen by the authors; no ablation on T except a single-split variant in Appendix F.
  • Per-iteration LLM call budget = 32
    Budget per subproblem evolution iteration; chosen to be comparable with the 160-kernel baseline budget.
  • Beam width and expansion size = 4 and 1
    Beam search hyperparameters shared by MKEvolve and the beam-search baseline; no tuning study.
  • Swap threshold tau = 0.9
    MKEvolve (Swap) replaces subkernels below 0.9x torch.compile speedup with PyTorch implementations; this only affects the Swap variant.
axioms (5)
  • domain assumption KernelBench L2/L3 and its atol/rtol=1e-4 correctness criterion are a meaningful proxy for kernel-generation usefulness.
    All headline correctness numbers are measured against this benchmark criterion (Section 4).
  • domain assumption The TritonRL-based cheating detector reliably identifies kernels that secretly call PyTorch operations.
    Cheating detection is adopted from Woo et al. (2026) and used as a filter for all methods (Section 3, Appendix A).
  • domain assumption Error accumulation across composed subkernels is mild enough that per-subkernel 1e-4 checks imply end-to-end correctness for FP32.
    Appendix E only sketches a probabilistic bound, and Appendix D shows threshold dilemmas for BF16, so this premise is not established for all datatypes.
  • domain assumption Current LLMs can produce correct decompositions and Triton kernels for these benchmarks.
    The whole method is an LLM agent pipeline; if the base LLM is too weak, decomposition and subkernel synthesis fail.
  • domain assumption torch.compile is a fair speedup baseline for the generated kernels.
    All speedups are reported relative to torch.compile on an A100, following KernelBench conventions (Section 4).
invented entities (1)
  • None no independent evidence
    purpose: No new physical, mathematical, or theoretical entities are introduced.
    The 'subproblems', 'subkernels', and 'agents' are software artifacts, not new postulated entities with independent falsifiable handles.

pith-pipeline@v1.3.0-alltime-deepseek · 36180 in / 13782 out tokens · 133657 ms · 2026-08-02T10:33:38.298193+00:00 · methodology

0 comments
read the original abstract

Despite rapid progress in LLM-based code generation, writing correct and performant kernels for hardware accelerators remains a key bottleneck in scaling modern ML workloads. We present MKEvolve (Modular Kernel Evolve), a framework that iteratively co-evolves a modular decomposition of complex PyTorch modules and the LLM-generated kernel for each submodule, refining the decomposition by splitting and fusing across iterations while independently improving each subkernel via LLM-driven beam search. The resulting kernels are programmatic compositions of independently verified subkernels, making them configurable (subkernel implementations are swappable), interpretable (errors and speedups are traceable to specific subkernels), and readily adaptable to related model architectures. Experiments with Triton on KernelBench L2 and L3, spanning multi-operator sequences and full model architectures, show that MKEvolve improves both correctness and speedup over end-to-end direct synthesis baselines while reducing LLM token usage by up to 35%.

Figures

Figures reproduced from arXiv: 2607.20501 by Jason Yoo, Rajarshi Saha, Shaowei Zhu, Tao Yu, Wei Tang, Youngsuk Park.

Figure 1
Figure 1. Figure 1: MKEVOLVE kernel synthesis on the KernelBench L3 ConvolutionalVisionTransformer task for 5 outer loop iterations. Left: Subproblem structure evolution during refinement, with checkmarks and crosses respectively denoting correct and invalid subkernels. MKEVOLVE iteratively decomposes the subproblem structure until all subproblems admit correct subkernels, then considers subproblem fusion. Right: Subkernel (L… view at source ↗
Figure 2
Figure 2. Figure 2: MKEvolve Algorithm 1 Visualization. problem decomposition itself, resulting in more efficient kernel implementations. Finally, MKEVOLVE relies on a programmatic evaluation pipeline with cheating detection and strict correctness criteria, instead of LLM-generated tests, improving robustness and reliability. 3. Methodology Algorithm 1 outlines the core MKEVOLVE workflow. MKE￾VOLVE produces a Python codebase … view at source ↗
Figure 3
Figure 3. Figure 3: Transferability of MKEvolve-generated Triton kernels across MobileNetV1 variants. Top: The base kernel invokes four subkernels sequentially, achieving 1.8× speedup over torch.compile. Middle & Bottom: Replacing AvgPool2D with MaxPool2D or LPPool2D yields adapted kernels achieving 1.7× and 1.8× speedups, requiring only 4 LLM calls to produce the new pooling subkernels. 63.0% 67.5% 58.5% 68.0% 58.5% 52.5% vs… view at source ↗
Figure 4
Figure 4. Figure 4: Win-rate (proportion of tasks for which the row method achieved higher speedup than the column method) heatmaps com￾paring various MKEvolve variants and beam search, using both Claude 4.5 Opus and gpt-oss as backbone LLM to solve L2 and L3 tasks sampled from KernelBench. Further Analysis [PITH_FULL_IMAGE:figures/full_fig_p007_4.png] view at source ↗
Figure 5
Figure 5. Figure 5: Claude 4.5 Opus KernelBench L2 experiment improvement over time plot for all baselines. We note that MKEvolve and its variants achieve the best final performance on all metrics while consuming fewer tokens per outer iteration. 10 [PITH_FULL_IMAGE:figures/full_fig_p010_5.png] view at source ↗
Figure 6
Figure 6. Figure 6: Claude 4.5 Opus KernelBench L3 experiment improvement over time plot for all baselines. We note that MKEvolve and its variants achieve the best final performance on all but one metric while consuming fewer tokens per outer iteration. C. Example Second Seed Results METHOD Correct (↑) Fast0.5 (↑) Fast1 (↑) Fast2 (↑) # Tokens Beam Search 0.95 0.66 0.37 0.03 2.0 × 106 MKEvolve 0.99 0.75 0.49 0.04 1.7 × 106 MKE… view at source ↗

discussion (0)

Sign in with ORCID, Apple, or X to comment. Anyone can read and Pith papers without signing in.

Reference graph

Works this paper leans on

40 extracted references · 2 linked inside Pith

  1. [1]

    **Functional Equivalence**:`ModelNew( *get_inputs())`must produce the same output as`Model( *get_inputs())`from the original problem file (numerical equivalence check) ,→ ,→

  2. [2]

    **Subkernel Standalone Execution**: Each subkernel in`subkernel_problems/` must execute successfully with its own`get_inputs()`and`get_init_inputs()` (or`get_input_configs()`): ,→ ,→ -`model = Model( *get_init_inputs())`must instantiate without error -`output = model( *get_inputs())`must execute without error

  3. [3]

    error profile

    URL https://arxiv.org/abs/2509.0 7506. Wen, Z., Zhang, Y ., Li, Z., Liu, Z., Xie, L., and Zhang, T. MultiKernelBench: A Multi-Platform Benchmark for Kernel Generation, 2025. URL https://arxiv.or g/abs/2507.17773. Woo, J., Zhu, S., Nie, A., Jia, Z., Wang, Y ., and Park, Y . Tritonrl: Training llms to think and code triton without cheating, 2026. URL https:...

    Pith/arXiv arXiv 2025

  4. [4]

    **Init Function Consistency**: The`get_init_inputs()`in each subkernel file must return arguments that`ModelNew`uses to initialize the corresponding `Model`module ,→ ,→

  5. [5]

    **Input Function Coverage**: If`ModelNew`nn.Module at `decomposed_problem.py`contain multiple instantiation of subkernel PyTorch modules in`subkernel_problems`, the corresponding subkernel PyTorch module file must have the function`get_input_configs()`that returns a list of dicts with functions that return`get_init_inputs()`and`get_inputs()` shapes presen...

  6. [6]

    **Shape Consistency**: The`get_inputs()`in each subkernel file must return tensors with shapes matching what`ModelNew.forward()`will actually pass to that submodule ,→ ,→

  7. [7]

    Assume that the modules will be in PyTorch`eval`mode (this is highly relevant for BatchNorm, etc),→ {% if assert_single_launch %}

  8. [8]

    M", "N",

    The`ModelNew`module must only launch exactly ONE fused Triton kernel (multiple launches of the same kernel allowed) that performs all meaningful computation of the module ,→ ,→ {% endif %} {% if show_example -%} Example solution format for Triton persistent matmul kernel (wrapper can use PyTorch, but kernel must purely use Triton):,→ ```python import torc...

  9. [9]

    reduction requires global synchronization

    **Cheating Detection**: The only PyTorch modules initialized at `decomposed_problem.py`should be modules from`subkernel_problems/`,→ ## FUSION OBJECTIVE 30 MKEvolve: A Modular Multi-Agent Framework for Kernel Code Generation Maximize fusion opportunities by grouping operations that can be efficiently executed in a single kernel. Only split into separate s...

  10. [10]

    Analyze the forward pass as a dataflow graph

  11. [11]

    Identify operations requiring global synchronization (fusion blockers)

  12. [12]

    Group all other operations between blockers into single submodules

  13. [13]

    Each submodule = one kernel opportunity

  14. [14]

    get_init_inputs

    Ensure exact output semantics, dtype behavior, and broadcasting rules are preserved,→ ## SUBKERNEL INPUT FUNCTION REQUIREMENTS **CRITICAL DECISION RULE:** For each subkernel file, count how many times`ModelNew.__init__`instantiates that subkernel's Model class:,→ - **1 instantiation** →define`get_inputs()`and`get_init_inputs()` 31 MKEvolve: A Modular Mult...

  15. [15]

    ,→ ,→ ,→

    **One file per operation pattern**: If the same fused operation pattern (e.g., Linear+ReLU, Conv2d+ReLU+MaxPool) appears multiple times in the original model, create ONE subkernel file and instantiate it multiple times in `ModelNew.__init__`with different parameters. ,→ ,→ ,→

  16. [16]

    **Parameterize by dimensions**: The subkernel`Model.__init__`should accept dimension arguments (e.g.,`in_features`,`out_features`,`in_channels`, `out_channels`) that configure the operation for different shapes. ,→ ,→

  17. [17]

    Linear+ReLU

    **Identify operation patterns, not instances**: When analyzing the forward pass, group by operation TYPE (e.g., "Linear+ReLU", "Conv2d+ReLU"), not by layer index or specific dimension values. ,→ ,→

  18. [18]

    **get_inputs() shows ONE example shape**: The`get_inputs()`function should return tensors for one valid input configuration.,→

  19. [19]

    PYTORCH TO KERNEL PROBLEM

    **get_input_configs() shows MULTIPLE example shapes**: If the same fused operation pattern appears multiple times in the original model, the file corresponding to the pattern must have`get_input_configs()`function that returns a list of dictionaries with functions`get_inputs`and `get_init_inputs`for all shapes ,→ ,→ ,→ ,→ 32 MKEvolve: A Modular Multi-Agen...

  20. [20]

    PyTorch module named`Model`that extends`torch.nn.Module`

  21. [21]

    Function that returns a sample argument for the module's`__init__`method `get_init_inputs`,→

  22. [22]

    Function that returns a sample argument for the module's`forward`method `get_inputs`,→ {{ problem_descriptor }} Your task is to produce a PyTorch module named`ModelNew`that is a Triton implementation of the`Model`module satisfying the following requirements.,→

  23. [23]

    The`ModelNew`module contains`__init__`and`forward`methods with the same input and output signatures as the`Model`module,→

  24. [24]

    The`ModelNew`module's`forward`method must - Correctly and efficiently implement the`Model`module's`forward`method in Triton,→ - Invoke a separately written Triton kernel function decorated with @triton.jit - The Triton kernel can be named anything (e.g., _kernel) - The Triton kernel function must explicitly accept all necessary tensors as function inputs ...

  25. [25]

    If the`Model`'s`__init__`method implicitly initializes parameters/tensors, explicitly initialize these values in the exact same order in the`__init__` method ,→ ,→ - Example 1:`Model`contains`nn.Linear`/`nn.Conv3D`and its weights and bias values need to be explicitly passed to the Triton kernel function,→ ```python # This is how nn.Linear initializes its ...

  26. [26]

    The`ModelNew`module must not cheat by initializing torch.nn module or trying to replicate`Model`'s forward pass using PyTorch operations or shortcuts,→

  27. [27]

    Do not repeat the raw original`Model`code anywhere, even in the comments

  28. [28]

    Keep logic in the`forward`method and the kernel minimal; speed is important and the entire`forward`method execution time is measured,→

  29. [31]

    Correctness: Are the outputs of`Model`and`ModelNew`'s`forward`methods sufficiently close?,→

  30. [32]

    Speedup: How fast is the`ModelNew`'s Triton-based`forward`method compared to`Model`'s PyTorch-based`forward`method?,→ ## General Triton Guidelines (Helpful):

  31. [33]

    KERNEL STRUCTURE: - Use @triton.jit decorator for kernel functions - Use tl.constexpr for compile-time constants (BLOCK_SIZE, etc.) - Include proper type hints and launch metadata when needed

  32. [34]

    MEMORY ACCESS PATTERNS: - Use tl.load and tl.store with proper masking 37 MKEvolve: A Modular Multi-Agent Framework for Kernel Code Generation - Coalesce memory accesses for optimal performance - Use tensor descriptors for advanced memory operations (TMA) - Handle boundary conditions with masks

  33. [35]

    INDEXING AND GRID: - Use tl.program_id(axis) for block indices - Calculate offsets: pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE) - Use tl.cdiv for ceiling division - Always mask for out-of-bounds protection

  34. [36]

    OPTIMIZATION TECHNIQUES: - Use @triton.autotune for automatic configuration selection (Autotuned parameters must be declared as tl.constexpr meta-parameters and must not be passed as runtime kernel arguments) ,→ ,→ - Choose appropriate BLOCK_SIZE (powers of 2: 64, 128, 256, 512, 1024) - Leverage tensor cores with tl.dot for matrix operations - Use warp sp...

  35. [37]

    COMMON PATTERNS: a) Elementwise operations: Load -> Compute -> Store b) Reductions: Use tl.reduce with proper axis (can also use tl.sum(), tl.max(), tl.min(), etc),→ c) Matrix multiplication: Use tl.dot with accumulator d) Softmax: Online normalization for numerical stability e) Fused operations: Combine multiple ops in single kernel and document the fuse...

  36. [38]

    ADVANCED FEATURES: - Persistent kernels for better SM utilization - Tensor Memory Accelerator (TMA) descriptors - Multi-stage pipelines with num_stages - Warp specialization with warp_specialize parameter

  37. [39]

    RUNTIME CONSTRAINTS: - Wrappers: validate/allocate/launch only; no math 38 MKEvolve: A Modular Multi-Agent Framework for Kernel Code Generation - All compute runs in Triton kernels; no torch.nn, torch.nn.functional (e.g., F.*), or other PyTorch compute ops, including general tensor-tensor math like torch.matmul/mm/bmm/einsum or their Tensor method forms, ...

  38. [40]

    TRITON MATH DOCS (from`triton.language`module or`tl`): - The following is a comprehensive list of math related`tl`ops; if a math op in`tl`is not in the following list, it does not exist,→ - LinAlg Ops: tl.dot, tl.dot_scaled - Math Ops: tl.abs, tl.cdiv, tl.ceil, tl.clamp, tl.cos, tl.div_rn, tl.erf, tl.exp, tl.exp2, tl.fdiv, tl.floor, tl.fma, tl.log, tl.log...

  39. [2024]

    URL https://arxiv.org/abs/2312.1 4820. Chen, T., Ye, Z., Xu, B., Ye, Z., Liu, T., Hassani, A., Chen, T., Kerr, A., Wu, H., Xu, Y ., Chen, Y .-J., Chen, H., Kane, A., Krashinsky, R., Liu, M.-Y ., Grover, V ., Ceze, L., Bringmann, R., Tran, J., Liu, W., Xie, F., Lightstone, M., and Shi, H. Avo: Agentic variation operators for autonomous evolutionary search,...

    arXiv 2026

  40. [2025]

    Liao, G., Qin, H., Wang, Y ., Golden, A., Kuchnik, M., Yetim, Y ., Ang, J

    URL https://arxiv.org/abs/2507.0 5687. Liao, G., Qin, H., Wang, Y ., Golden, A., Kuchnik, M., Yetim, Y ., Ang, J. J., Fu, C., He, Y ., Hsia, S., Jiang, Z., Li, D., Pashkevich, U., Puvvada, V ., Shi, F., Steiner, M., Xiao, R., Yan, N., Yu, X., Fang, Z., Levenstein, R., Ho, K., Zhu, H., Hammond, A., Li, R., Mathews, A., Gondkar, K., Zainul-Abedin, A., Singh...

    Pith/arXiv arXiv 2026