From a4ac5c49059b6206d73a631ca0af508f8d0829f7 Mon Sep 17 00:00:00 2001 From: Will Feng Date: Fri, 13 Mar 2026 15:38:52 -0700 Subject: [PATCH 1/8] docs: add TileIR backend usage guide to helion-hackathon.md Documents ENABLE_TILE=0 vs ENABLE_TILE=1 and the TileIR compilation pipeline available via nvtriton on B200 instances. Covers how to enable TileIR with Helion (ENABLE_TILE=1 + HELION_BACKEND=tileir), the different tunables (num_ctas/occupancy vs num_warps/maxnreg), and how to hardcode TileIR configs in submissions. Co-Authored-By: Claude Opus 4.6 (1M context) --- docs/helion-hackathon.md | 60 ++++++++++++++++++++++++++++++++++++++++ 1 file changed, 60 insertions(+) diff --git a/docs/helion-hackathon.md b/docs/helion-hackathon.md index 5986ad0..bb43286 100644 --- a/docs/helion-hackathon.md +++ b/docs/helion-hackathon.md @@ -250,6 +250,66 @@ This is the approach you should use for KernelBot submissions — a fixed config 3. **Hardcode both the config and ACF path** in your submission file 4. **Verify** the ACF path (`/opt/booster_pack/...`) exists on B200 — it does on all hackathon instances +## Using the TileIR Backend + +The B200 instances ship with **nvtriton**, NVIDIA's extended Triton compiler that includes a **TileIR** backend. TileIR is an alternative compilation pipeline that bypasses LLVM and compiles directly to CUBIN via NVIDIA's `tileiras` compiler. It can produce faster kernels on Blackwell (B200) GPUs. + +### ENABLE_TILE=0 vs ENABLE_TILE=1 + +| | `ENABLE_TILE=0` (default) | `ENABLE_TILE=1` | +|---|---|---| +| **Pipeline** | Triton IR → LLVM IR → PTX → CUBIN | Triton IR → TileIR → CUBIN (via `tileiras`) | +| **Tunables** | `num_warps`, `num_stages` (1-8), `maxnreg` | `num_ctas` (1-2), `num_stages` (1-10), `occupancy` (1-8) | +| **Helion backend** | `triton` (default) | `tileir` | + +### Enabling TileIR with Helion + +Set both environment variables **before** importing Triton or Helion: + +```bash +export ENABLE_TILE=1 +export HELION_BACKEND=tileir +``` + +Then write your kernel as usual. Helion automatically adjusts the autotuner search space for TileIR — it searches over `num_ctas` and `occupancy` instead of `num_warps` and `maxnreg`: + +```python +# With TileIR, the autotuner explores tileir-specific tunables +@helion.kernel(static_shapes=True) +def my_kernel(...): + ... +``` + +### Hardcoding a TileIR config in your submission + +After autotuning with TileIR, hardcode the config as usual. TileIR configs use `num_ctas` and `occupancy` instead of `num_warps`: + +```python +@helion.kernel(config=helion.Config( + block_sizes=[64, 64], + num_ctas=1, + num_stages=5, + occupancy=4, + # ... rest of your tuned config +)) +def my_kernel(...): + ... +``` + +**Important:** Your submission still needs `ENABLE_TILE=1` and `HELION_BACKEND=tileir` at runtime. Add them at the top of your submission file: + +```python +import os +os.environ["ENABLE_TILE"] = "1" +os.environ["HELION_BACKEND"] = "tileir" + +import helion # must be imported after setting env vars +``` + +### Which backend should I use? + +Try both. Autotune your kernel with the default Triton backend (`ENABLE_TILE=0`) and with TileIR (`ENABLE_TILE=1`), then submit whichever gives better benchmark numbers. TileIR can be significantly faster on B200 for some kernel patterns, but the default backend may win for others. + ## Tips - **Iterate locally first.** Use your Nebius B200 to develop and autotune. Only submit to KernelBot once you have a hardcoded config that works. From 8bbd20dd393371d8ea36bc87264d7d08d30e3dc7 Mon Sep 17 00:00:00 2001 From: Will Feng Date: Fri, 13 Mar 2026 15:46:44 -0700 Subject: [PATCH 2/8] docs: restructure ACF + TileIR as optional performance knobs Group both sections under a single "Optional: Extra Performance Knobs" heading to emphasize neither is required. Streamline both into step 1 (autotune) / step 2 (hardcode) format. Add a "Which combination" section showing all 4 options to try. Co-Authored-By: Claude Opus 4.6 (1M context) --- docs/helion-hackathon.md | 83 +++++++++++++++------------------------- 1 file changed, 31 insertions(+), 52 deletions(-) diff --git a/docs/helion-hackathon.md b/docs/helion-hackathon.md index bb43286..78809c9 100644 --- a/docs/helion-hackathon.md +++ b/docs/helion-hackathon.md @@ -193,22 +193,24 @@ popcorn submit causal_conv1d_py/submission.py --gpu B200_Nebius --leaderboard ca This returns GPU throughput, pipe utilization, and warp stall metrics, plus a downloadable `.ncu-rep` trace file you can open in the Nsight Compute GUI. See [profiling.md](profiling.md) for details on interpreting the output. -## Using ACF Files (Booster Pack) +## Optional: Extra Performance Knobs -Each B200 instance comes with pre-tuned **PTXAS Advanced Controls Files (ACFs)** at `/opt/booster_pack/`. These are low-level NVIDIA PTX assembler configurations that can improve kernel performance beyond what Helion's standard autotuner finds. +The sections below describe two **optional** techniques that can squeeze extra performance out of your kernels. Neither is required — you can place on the leaderboard without them. Try them after you have a working kernel with a tuned config. + +### ACF Files (Booster Pack) + +Each B200 instance has pre-tuned **PTXAS Advanced Controls Files (ACFs)** at `/opt/booster_pack/`. ACFs are low-level NVIDIA PTX assembler configurations that can improve performance beyond what Helion's standard autotuner finds. Available files: ``` /opt/booster_pack/ -├── causal_conv_0.acf ... causal_conv_2.acf -├── chunk_fwd_h_0.acf ... chunk_fwd_h_1.acf -├── chunk_fwd_o_0.acf ... chunk_fwd_o_6.acf -├── fp8_group_quant_0.acf ... fp8_group_quant_6.acf -└── recompute_w_u_fwd_0.acf ... recompute_w_u_fwd_4.acf +├── causal_conv_*.acf (3 files) +├── chunk_fwd_h_*.acf (2 files) +├── chunk_fwd_o_*.acf (7 files) +├── fp8_group_quant_*.acf (7 files) +└── recompute_w_u_fwd_*.acf (5 files) ``` -### Using ACFs during autotuning - -Pass `autotune_search_acf` to the `@helion.kernel` decorator. Helion treats each ACF as another tunable parameter — every config candidate gets tried with each ACF file (plus the default `-O3` baseline): +**Step 1: Autotune with ACFs.** Pass `autotune_search_acf` to include ACFs in the search space. Helion tries each ACF file (plus the default `-O3` baseline) as another tunable parameter: ```python from pathlib import Path @@ -223,11 +225,9 @@ def my_kernel(...): ... ``` -> **Important:** `autotune_search_acf` only takes effect when the autotuner actually runs. If you set `autotune_effort="none"` or provide a fixed `config=`, the ACF list is ignored. - -### Hardcoding an ACF in your submission +> **Note:** `autotune_search_acf` only takes effect when the autotuner actually runs. It is ignored with `autotune_effort="none"` or a fixed `config=`. -After autotuning finds the best ACF, include it in your hardcoded config via `advanced_controls_file`. The autotuner prints the winning ACF path — copy it into your `Config`: +**Step 2: Hardcode the best ACF in your submission.** After autotuning, look for the `advanced_controls_file` field in the best config and copy it: ```python @helion.kernel(config=helion.Config( @@ -241,20 +241,9 @@ def my_kernel(...): ... ``` -This is the approach you should use for KernelBot submissions — a fixed config with a fixed ACF, no autotuning at runtime. - -### Recommended workflow - -1. **Autotune with ACFs locally** on your B200, using the matching `*_*.acf` files for your problem -2. **Check the best config output** — look for the `advanced_controls_file` field -3. **Hardcode both the config and ACF path** in your submission file -4. **Verify** the ACF path (`/opt/booster_pack/...`) exists on B200 — it does on all hackathon instances - -## Using the TileIR Backend - -The B200 instances ship with **nvtriton**, NVIDIA's extended Triton compiler that includes a **TileIR** backend. TileIR is an alternative compilation pipeline that bypasses LLVM and compiles directly to CUBIN via NVIDIA's `tileiras` compiler. It can produce faster kernels on Blackwell (B200) GPUs. +### TileIR Backend -### ENABLE_TILE=0 vs ENABLE_TILE=1 +The B200 instances also ship with **nvtriton**, NVIDIA's extended Triton compiler that includes a **TileIR** backend — an alternative compilation pipeline that bypasses LLVM and compiles directly to CUBIN via NVIDIA's `tileiras` compiler. | | `ENABLE_TILE=0` (default) | `ENABLE_TILE=1` | |---|---|---| @@ -262,29 +251,25 @@ The B200 instances ship with **nvtriton**, NVIDIA's extended Triton compiler tha | **Tunables** | `num_warps`, `num_stages` (1-8), `maxnreg` | `num_ctas` (1-2), `num_stages` (1-10), `occupancy` (1-8) | | **Helion backend** | `triton` (default) | `tileir` | -### Enabling TileIR with Helion - -Set both environment variables **before** importing Triton or Helion: +**Step 1: Enable TileIR.** Set both environment variables **before** importing Triton or Helion: ```bash export ENABLE_TILE=1 export HELION_BACKEND=tileir ``` -Then write your kernel as usual. Helion automatically adjusts the autotuner search space for TileIR — it searches over `num_ctas` and `occupancy` instead of `num_warps` and `maxnreg`: +Then autotune as usual. Helion automatically adjusts the search space for TileIR — it searches over `num_ctas` and `occupancy` instead of `num_warps` and `maxnreg`. -```python -# With TileIR, the autotuner explores tileir-specific tunables -@helion.kernel(static_shapes=True) -def my_kernel(...): - ... -``` +**Step 2: Hardcode the TileIR config in your submission.** TileIR configs use `num_ctas` and `occupancy` instead of `num_warps`. Your submission also needs the env vars set before imports: -### Hardcoding a TileIR config in your submission +```python +import os +os.environ["ENABLE_TILE"] = "1" +os.environ["HELION_BACKEND"] = "tileir" -After autotuning with TileIR, hardcode the config as usual. TileIR configs use `num_ctas` and `occupancy` instead of `num_warps`: +import helion # must be imported after setting env vars +import helion.language as hl -```python @helion.kernel(config=helion.Config( block_sizes=[64, 64], num_ctas=1, @@ -296,19 +281,13 @@ def my_kernel(...): ... ``` -**Important:** Your submission still needs `ENABLE_TILE=1` and `HELION_BACKEND=tileir` at runtime. Add them at the top of your submission file: - -```python -import os -os.environ["ENABLE_TILE"] = "1" -os.environ["HELION_BACKEND"] = "tileir" - -import helion # must be imported after setting env vars -``` - -### Which backend should I use? +### Which combination should I use? -Try both. Autotune your kernel with the default Triton backend (`ENABLE_TILE=0`) and with TileIR (`ENABLE_TILE=1`), then submit whichever gives better benchmark numbers. TileIR can be significantly faster on B200 for some kernel patterns, but the default backend may win for others. +Try all combinations and submit whichever gives the best benchmark numbers: +- Default Triton backend without ACF +- Default Triton backend with ACF +- TileIR backend without ACF +- TileIR backend with ACF ## Tips From 4882ec0a6d7be4d4aa69932bdd5db7eb207ab3a2 Mon Sep 17 00:00:00 2001 From: Will Feng Date: Fri, 13 Mar 2026 15:48:42 -0700 Subject: [PATCH 3/8] docs: remove "(Booster Pack)" from ACF heading Co-Authored-By: Claude Opus 4.6 (1M context) --- docs/helion-hackathon.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/docs/helion-hackathon.md b/docs/helion-hackathon.md index 78809c9..35df7e1 100644 --- a/docs/helion-hackathon.md +++ b/docs/helion-hackathon.md @@ -197,7 +197,7 @@ This returns GPU throughput, pipe utilization, and warp stall metrics, plus a do The sections below describe two **optional** techniques that can squeeze extra performance out of your kernels. Neither is required — you can place on the leaderboard without them. Try them after you have a working kernel with a tuned config. -### ACF Files (Booster Pack) +### ACF Files Each B200 instance has pre-tuned **PTXAS Advanced Controls Files (ACFs)** at `/opt/booster_pack/`. ACFs are low-level NVIDIA PTX assembler configurations that can improve performance beyond what Helion's standard autotuner finds. Available files: From f10d2356372c30b2eb297b8dc407ef2cc31b4a1c Mon Sep 17 00:00:00 2001 From: Will Feng Date: Fri, 13 Mar 2026 15:49:32 -0700 Subject: [PATCH 4/8] docs: consolidate TileIR env var instructions MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Remove duplicate bash export block — the Python os.environ in the code example is sufficient for both local autotuning and submissions. Co-Authored-By: Claude Opus 4.6 (1M context) --- docs/helion-hackathon.md | 11 ++--------- 1 file changed, 2 insertions(+), 9 deletions(-) diff --git a/docs/helion-hackathon.md b/docs/helion-hackathon.md index 35df7e1..6b760dd 100644 --- a/docs/helion-hackathon.md +++ b/docs/helion-hackathon.md @@ -251,16 +251,9 @@ The B200 instances also ship with **nvtriton**, NVIDIA's extended Triton compile | **Tunables** | `num_warps`, `num_stages` (1-8), `maxnreg` | `num_ctas` (1-2), `num_stages` (1-10), `occupancy` (1-8) | | **Helion backend** | `triton` (default) | `tileir` | -**Step 1: Enable TileIR.** Set both environment variables **before** importing Triton or Helion: +**Step 1: Enable TileIR and autotune.** Set the env vars before importing Helion, then autotune as usual. Helion automatically adjusts the search space — it searches over `num_ctas` and `occupancy` instead of `num_warps` and `maxnreg`. -```bash -export ENABLE_TILE=1 -export HELION_BACKEND=tileir -``` - -Then autotune as usual. Helion automatically adjusts the search space for TileIR — it searches over `num_ctas` and `occupancy` instead of `num_warps` and `maxnreg`. - -**Step 2: Hardcode the TileIR config in your submission.** TileIR configs use `num_ctas` and `occupancy` instead of `num_warps`. Your submission also needs the env vars set before imports: +**Step 2: Hardcode the TileIR config in your submission.** TileIR configs use `num_ctas` and `occupancy` instead of `num_warps`. The env vars must be set before imports: ```python import os From e4cb531e4415399726e9726f9d2244451d16367b Mon Sep 17 00:00:00 2001 From: Will Feng Date: Fri, 13 Mar 2026 15:50:22 -0700 Subject: [PATCH 5/8] docs: clarify TileIR tunables come from autotuner output Co-Authored-By: Claude Opus 4.6 (1M context) --- docs/helion-hackathon.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/docs/helion-hackathon.md b/docs/helion-hackathon.md index 6b760dd..650015f 100644 --- a/docs/helion-hackathon.md +++ b/docs/helion-hackathon.md @@ -253,7 +253,7 @@ The B200 instances also ship with **nvtriton**, NVIDIA's extended Triton compile **Step 1: Enable TileIR and autotune.** Set the env vars before importing Helion, then autotune as usual. Helion automatically adjusts the search space — it searches over `num_ctas` and `occupancy` instead of `num_warps` and `maxnreg`. -**Step 2: Hardcode the TileIR config in your submission.** TileIR configs use `num_ctas` and `occupancy` instead of `num_warps`. The env vars must be set before imports: +**Step 2: Hardcode the TileIR config in your submission.** Copy the best config from the autotuner output (it will include TileIR-specific fields like `num_ctas` and `occupancy`). The env vars must be set before imports: ```python import os From 115daeb1d675648e1217be81bc4593534c0bb1de Mon Sep 17 00:00:00 2001 From: Will Feng Date: Fri, 13 Mar 2026 15:56:05 -0700 Subject: [PATCH 6/8] docs: shorten "Which should I use?" section Co-Authored-By: Claude Opus 4.6 (1M context) --- docs/helion-hackathon.md | 8 ++------ 1 file changed, 2 insertions(+), 6 deletions(-) diff --git a/docs/helion-hackathon.md b/docs/helion-hackathon.md index 650015f..cb40481 100644 --- a/docs/helion-hackathon.md +++ b/docs/helion-hackathon.md @@ -274,13 +274,9 @@ def my_kernel(...): ... ``` -### Which combination should I use? +### Which should I use? -Try all combinations and submit whichever gives the best benchmark numbers: -- Default Triton backend without ACF -- Default Triton backend with ACF -- TileIR backend without ACF -- TileIR backend with ACF +Try both backends and with/without ACFs, then submit whichever gives the best benchmark numbers. ## Tips From 2b1bef67182a187b9d79676d380459989964ff10 Mon Sep 17 00:00:00 2001 From: Will Feng Date: Fri, 13 Mar 2026 15:56:47 -0700 Subject: [PATCH 7/8] docs: be explicit about ENABLE_TILE=0 vs ENABLE_TILE=1 Co-Authored-By: Claude Opus 4.6 (1M context) --- docs/helion-hackathon.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/docs/helion-hackathon.md b/docs/helion-hackathon.md index cb40481..0c72abe 100644 --- a/docs/helion-hackathon.md +++ b/docs/helion-hackathon.md @@ -276,7 +276,7 @@ def my_kernel(...): ### Which should I use? -Try both backends and with/without ACFs, then submit whichever gives the best benchmark numbers. +Try both `ENABLE_TILE=0` and `ENABLE_TILE=1`, with and without ACFs, then submit whichever gives the best benchmark numbers. ## Tips From 1a127c1be45f4fa14192ad55e22755c0a64efb91 Mon Sep 17 00:00:00 2001 From: Will Feng Date: Fri, 13 Mar 2026 16:02:17 -0700 Subject: [PATCH 8/8] docs: simplify TileIR comparison table to just backend names Co-Authored-By: Claude Opus 4.6 (1M context) --- docs/helion-hackathon.md | 6 ++---- 1 file changed, 2 insertions(+), 4 deletions(-) diff --git a/docs/helion-hackathon.md b/docs/helion-hackathon.md index 0c72abe..c71230b 100644 --- a/docs/helion-hackathon.md +++ b/docs/helion-hackathon.md @@ -247,11 +247,9 @@ The B200 instances also ship with **nvtriton**, NVIDIA's extended Triton compile | | `ENABLE_TILE=0` (default) | `ENABLE_TILE=1` | |---|---|---| -| **Pipeline** | Triton IR → LLVM IR → PTX → CUBIN | Triton IR → TileIR → CUBIN (via `tileiras`) | -| **Tunables** | `num_warps`, `num_stages` (1-8), `maxnreg` | `num_ctas` (1-2), `num_stages` (1-10), `occupancy` (1-8) | -| **Helion backend** | `triton` (default) | `tileir` | +| **Helion backend** | `triton` | `tileir` | -**Step 1: Enable TileIR and autotune.** Set the env vars before importing Helion, then autotune as usual. Helion automatically adjusts the search space — it searches over `num_ctas` and `occupancy` instead of `num_warps` and `maxnreg`. +**Step 1: Enable TileIR and autotune.** Set the env vars before importing Helion, then autotune as usual. Helion automatically adjusts the search space for the TileIR backend. **Step 2: Hardcode the TileIR config in your submission.** Copy the best config from the autotuner output (it will include TileIR-specific fields like `num_ctas` and `occupancy`). The env vars must be set before imports: