This directory contains pretuned Helion kernels, benchmark shape sweeps, and checked-in AOT heuristic files. They are meant to be useful copy/paste starting points for common kernel patterns while also being runnable examples for people who want to quickly try Helion.

The checked-in heuristics let these kernels run immediately without online autotuning. Each entry lists the NVIDIA architecture it supports (currently H100, B200, or both), and Helion picks the matching file at runtime. Treat the files as kernel recipes: copy the kernel and its local _helion_aot_* heuristic into your code, then retune when your target shapes or hardware differ materially from the included sweep.

Each kernel module has a main() that benchmarks against one or more named reference baselines.

pretuned_kernels/
├── README.md
├── vector_add/
│   ├── vector_add.py                          # the kernel + main()
│   ├── _helion_aot_vector_add_cuda_sm100.py   # B200 heuristic
│   └── _helion_aot_vector_add_cuda_sm90.py    # H100 heuristic
├── softmax/
├── layer_norm/
├── rms_norm/
├── cross_entropy/
├── rope/
├── scaled_mm/
├── scale_mm_cute/                    # B200 CuTe (tcgen05) rowwise FP8 GEMM
├── grouped_gemm/                     # B200 grouped FP16 GEMM vs CuTeDSL
├── grouped_gemm_deepgemm/            # B200 grouped BF16 vs DeepGEMM + cuBLAS
├── nvfp4_gemv/                       # B200 Triton NVFP4 (W4A4 / W4A16) decode GEMV
├── nvfp4_gemv_cute/                  # B200 Helion CuTe NVFP4 decode GEMV
├── projection_rotary/                # B200 CuTe projection + rotary fragment epilogue
├── interleaved_swiglu/               # B200 CuTe projection + compact SwiGLU epilogue
├── silu_mul_fp8/                     # ported from vLLM (vllm/kernels/helion/ops)
├── dynamic_per_token_scaled_fp8_quant/
├── per_token_group_fp8_quant/
├── rms_norm_dynamic_per_token_quant/
├── rms_norm_per_block_quant/
├── silu_and_mul_per_block_quant/
├── fused_qk_norm_rope/
├── causal_conv1d/                    # TPU/Pallas fixed-config decode kernel
├── gdn_decode/                       # TPU/Pallas fixed-config recurrent decode
└── megakernels/
    ├── kda_decode/                   # Kimi-K3 KDA decode sublayer
    ├── qwen3_8b/                     # Qwen3-8B FP8 and BF16 decode layers
    ├── qwen3_30b_a3b_moe/            # Qwen3-30B-A3B MoE block
    ├── llama32_1b_decode_layer_bf16/ # Llama-3.2-1B BF16 decode layer
    ├── gemma4_a4b_moe/               # Gemma 4 26B-A4B MoE
    ├── gpt_oss_moe/                  # gpt-oss MoE
    ├── flash_mla/                    # DeepSeek-V3 MLA decode, one request
    ├── deepseek_v3_moe_nvfp4/        # DeepSeek-V3 NVFP4 MoE
    ├── deepseek_v3_moe_nvfp4_tp/     # DeepSeek-V3 NVFP4 MoE + all-reduce + RMSNorm, TP4
    ├── deepseek_v3_attention_nvfp4_tp/
    └── muse_spark/

Each kernel ships with one heuristic file per supported compute capability. At runtime Helion picks the file matching the current GPU.

Most kernels additionally benchmark against torch.compile of the listed PyTorch baseline (a speedup-comparison baseline only -- correctness is checked against the eager reference). The grouped-GEMM entries instead use the named CUDA references in the table. The headline speedup is Helion vs the fastest available baseline, and the per-kernel dropdown reports every baseline.

grouped_gemm compares Helion with the same pinned CUTLASS kernel. Required device pointer tables are initialized before graph capture, and both implementations are checked against the same pointer-target outputs before timing.

grouped_gemm_deepgemm is separate because it follows DeepGEMM's BF16 packed-A layout and official shape suite rather than the CUTLASS FP16 heterogeneous-M/N/K suite. Its AOT flow searches only valid 4--7-stage worklist_nm schedules and checks them against a PyTorch reference. It also reports cuBLAS on each logical group while excluding aligned padding from the common correctness contract. Both grouped benchmarks replay pre-captured graphs, clear L2 before every measurement, warm the GPU before each case, and rotate/reverse implementation order to reduce clock and ordering bias.

The model-region entries are fixed-capacity B200 examples. KDA decode ships ten Kimi-K3 physical envelopes (B1/B2/B4/B8/B16 by TP8/H12 or TP16/H6). The perf gate benchmarks H12 against vLLM's production CUDA fusion plus identical projections; correctness also covers vLLM's H6 fallback. Decode sequence length is absent from the signature, while mutable convolution/recurrent caches and state-slot indices remain runtime data. A source-matched seven-launch PDL graph separately measures the advantage from persistent cross-root scheduling. Each Helion definition contains the complete computation in one generated Triton kernel; the checked-in source does not import or assemble code from the probe directory. Their benchmark baselines require vLLM and correctness is checked before timing. Qwen executes the compiled Qwen3DecoderLayer with vLLM's auto policy for attention and FP8 linears. Gemma likewise leaves MoE dispatch on auto; its printed baseline name records the backend available in the installed vLLM/FlashInfer build. The Qwen check compares final numerical results against that production layer. When vLLM selects DeepGEMM, it also rounds dynamic activation scales to UE8M0; the Helion kernel retains its FP32 activation scales, so the comparison is not instruction-for-instruction.

The model-region benchmarks measure each captured implementation sequentially, with a fresh thermal warmup and cold-L2 samples, so one implementation's graph does not perturb the other's measurement.

The kernels ported from vLLM (vllm/kernels/helion/ops) benchmark each fused Helion kernel under CUDA graphs against a torch-native (unfused, eager) reference; silu_mul_fp8 ships an sm90 heuristic only, the rest ship both sm90 and sm100.

causal_conv1d and gdn_decode are Pallas kernels with checked-in fixed configurations rather than CUDA AOT heuristic files. Their module-level main() functions run real-TPU JAX-export correctness checks. They are not registered in the CUDA-only aggregate runner. The implementations and baselines follow the Apache-licensed vllm-project/tpu-inference state-cache contracts; gdn_decode currently targets H=2 because larger recurrent-state tiles exceed Helion's current aligned indirect-DMA VMEM plan.

Use this directory as a collection of pretuned kernels and runnable examples. For production code, copy the relevant kernel pattern into the application. If the shapes or target hardware differ from the included sweep, generate and commit an AOT heuristic for the application's target shapes and hardware.

For AOT pretuned kernels, Helion's runtime looks for _helion_aot_<kernel>_cuda_sm<NN>.py next to the kernel source file. Helion looks for AOT heuristic files for the current compute capability first, then falls back to older compatible CUDA/ROCm capabilities. For example, on sm120, if no sm120 heuristic exists, an sm100 heuristic file can be used.

Each kernel module has a main() that benchmarks the Helion kernel against its named reference baselines across the included shape set:

cd pretuned_kernels/softmax
python softmax.py

The external grouped-GEMM references need explicit pinned checkouts:

export HELION_CUTLASS_GROUPED_GEMM_SOURCE=/path/to/CUTLASS/.../grouped_gemm.py
python -m pretuned_kernels.grouped_gemm.grouped_gemm

export HELION_DEEPGEMM_ROOT=/path/to/built/DeepGEMM
python -m pretuned_kernels.grouped_gemm_deepgemm.grouped_gemm_deepgemm

The modules report the required commits when either variable is missing; the nightly workflow fetches, builds, and verifies those revisions automatically.

These kernels ship pretuned heuristics for specific GPU architectures. To add another GPU, run the AOT autotune workflow on that hardware against the kernel; the runner emits a new _helion_aot_<kernel>_<device>_<cc>.py next to the kernel source, which you commit alongside the existing one(s). Helion picks the right one at runtime based on the running GPU's compute capability (with fallback to older compatible capabilities, e.g. sm120 → sm100).

For a kernel whose main() benchmarks under CUDA graphs (use_cudagraph() returns True, e.g. scaled_mm and the vLLM-ported ops), set HELION_BENCHMARK_CUDAGRAPH=1 for the autotune run so the autotuner benchmarks candidate configs the same way — matching the deployment/benchmark timing regime.

See the Ahead-of-Time (AOT) Heuristic Tuning section of docs/deployment_autotuning.md for the end-to-end workflow, runner CLI, generated artifacts, and runtime fallback rules — including a worked "Pretuning a kernel for new hardware" walkthrough.