submission 553231 · KernelIndex

Use it

Vendorable · source mirrored · license unknownView source →

No package. Vendor the mirrored source: 82 lines, June 9 Researcher Reciprocity License v1.0.

submission.py

curl "https://kernelindex.com/api/v1/implementations/kernelbot-gated-deltanet-recompute-w-u-553231?include=source"

interfacepython

Compatibility

measured onNVIDIA B200

declared hardwareNVIDIA B200

architecturessm_100

dtypesfp32

Benchmark evidence

1 measurement across 1 GPU, fastest first.

Operation / workload

Hardware

Latency

Rank

Observed

Gated DeltaNet recompute_w_usuite of 3 cases

NVIDIA B200

153.6µs

#19 of 28

2026-03-14

Run →

Reported · How evidence levels are derived →

Source and license

sourceavailable

revision digestsha256:699fb3855ad18265ff1e5cd9b6b6a603f736d32b247d74cfa9b76a8501f74959

license declaredunknown

license concludedunknown

authorsAyush10

imported2026-08-15

Techniques

Extracted from the mirrored source by pattern, never inferred. Each row cites its line.

num-warps = 4(64, 64): helion.Config(num_warps=4, num_stages=3, l2_groupings=[4]),

stages = 3(64, 64): helion.Config(num_warps=4, num_stages=3, l2_groupings=[4]),

Kernel source

submission.py82 lines

#!POPCORN leaderboard gated_deltanet_recompute_w_u
#!POPCORN gpu B200_Nebius
# Team: KernalForge
# Fix: Group configs by (K,V) to reduce JIT compilations from 10 to 4
# Remove static_shapes so B,T,H are runtime values — fits within 420s timeout
from task import input_t, output_t

import torch
import helion
import helion.language as hl


# 4 unique (K,V) pairs across all test+benchmark shapes:
#   (64,64)   — 7 shapes
#   (64,128)  — 1 shape
#   (100,100) — 1 shape
#   (128,128) — 2 shapes
# Each group compiles ONE Triton kernel. 4 compilations × ~60s = ~240s < 420s timeout.
KV_CONFIGS: dict[tuple[int, int], helion.Config] = {
    (64, 64): helion.Config(num_warps=4, num_stages=3, l2_groupings=[4]),
    (64, 128): helion.Config(num_warps=8, num_stages=3),
    (100, 100): helion.Config(num_warps=8, num_stages=4, l2_groupings=[4]),
    (128, 128): helion.Config(num_warps=8, num_stages=4, l2_groupings=[8]),
}


def _make_kernel(config: helion.Config):
    @helion.kernel(dot_precision="ieee", config=config)
    def kernel(
        k: torch.Tensor,
        v: torch.Tensor,
        beta: torch.Tensor,
        A: torch.Tensor,
        beta_g: torch.Tensor,
    ) -> tuple[torch.Tensor, torch.Tensor]:
        B, T, H, K = k.shape
        V = v.shape[-1]
        C = hl.specialize(A.shape[-1])
        K = hl.specialize(K)
        V = hl.specialize(V)

        w_out = torch.empty_like(k)
        u_out = torch.empty_like(v)
        BH = B * H

        for flat_bh, rt in hl.tile([BH, T], block_size=[1, C]):
            b_idx = flat_bh.begin // H
            h_idx = flat_bh.begin % H

            a_chunk = A[b_idx, rt, h_idx, :].to(torch.float32)
            beta_chunk = beta[b_idx, rt, h_idx].to(torch.float32)
            beta_g_chunk = beta_g[b_idx, rt, h_idx].to(torch.float32)
            rhs_k = k[b_idx, rt, h_idx, :].to(torch.float32) * beta_g_chunk[:, None]
            rhs_v = v[b_idx, rt, h_idx, :].to(torch.float32) * beta_chunk[:, None]

            w_out[b_idx, rt, h_idx, :] = hl.dot(
                a_chunk,
                rhs_k,
                out_dtype=torch.float32,
            ).to(w_out.dtype)
            u_out[b_idx, rt, h_idx, :] = hl.dot(
                a_chunk,
                rhs_v,
                out_dtype=torch.float32,
            ).to(u_out.dtype)

        return w_out, u_out

    return kernel


_KERNELS = {kv: _make_kernel(cfg) for kv, cfg in KV_CONFIGS.items()}


def custom_kernel(data: input_t) -> output_t:
    k, v, beta, A, g = data
    K = k.shape[-1]
    V = v.shape[-1]
    beta_g = beta * torch.exp(g)
    kernel = _KERNELS[(K, V)]
    return kernel(k, v, beta, A, beta_g)

scrolls · 82 lines total

Source code from GPU Mode and the KernelBot dataset · June 9 Researcher Reciprocity License v1.0

Changes from previous submission

Against this author's previous submission submission 553133.

  #!POPCORN leaderboard gated_deltanet_recompute_w_u
  #!POPCORN gpu B200_Nebius
- # Team: Kernal Forge
- # Precompute beta_g = beta * exp(g) on host to eliminate exp() inside kernel
- from task import input_t, output_t
- 
- import torch
- import helion
- import helion.language as hl
- 
- 
- SHAPE_CONFIGS: dict[tuple[int, int, int, int, int], helion.Config] = {
-     (1, 64, 2, 64, 64): helion.Config(num_warps=4, num_stages=2),
-     (2, 128, 4, 64, 64): helion.Config(num_warps=4, num_stages=3),
-     (1, 256, 4, 64, 128): helion.Config(num_warps=8, num_stages=3),
-     (1, 64, 1, 64, 64): helion.Config(num_warps=4, num_stages=2),
-     (2, 512, 3, 64, 64): helion.Config(num_warps=4, num_stages=3, l2_groupings=[4]),
-     (2, 1024, 3, 64, 64): helion.Config(num_warps=4, num_stages=4, l2_groupings=[4]),
-     (3, 1024, 4, 100, 100): helion.Config(num_warps=8, num_stages=4, l2_groupings=[4]),
-     (4, 1024, 4, 128, 128): helion.Config(num_warps=8, num_stages=4, l2_groupings=[8]),
-     (2, 1536, 4, 128, 128): helion.Config(num_warps=8, num_stages=5, l2_groupings=[8]),
-     (4, 2048, 8, 64, 64): helion.Config(num_warps=8, num_stages=4, l2_groupings=[8]),
+ # Team: KernalForge
+ # Fix: Group configs by (K,V) to reduce JIT compilations from 10 to 4
+ # Remove static_shapes so B,T,H are runtime values — fits within 420s timeout
+ from task import input_t, output_t
+  
+ import torch
+ import helion
+ import helion.language as hl
+  
+  
+ # 4 unique (K,V) pairs across all test+benchmark shapes:
+ #   (64,64)   — 7 shapes
+ #   (64,128)  — 1 shape
+ #   (100,100) — 1 shape
+ #   (128,128) — 2 shapes
+ # Each group compiles ONE Triton kernel. 4 compilations × ~60s = ~240s < 420s timeout.
+ KV_CONFIGS: dict[tuple[int, int], helion.Config] = {
+     (64, 64): helion.Config(num_warps=4, num_stages=3, l2_groupings=[4]),
+     (64, 128): helion.Config(num_warps=8, num_stages=3),
+     (100, 100): helion.Config(num_warps=8, num_stages=4, l2_groupings=[4]),
+     (128, 128): helion.Config(num_warps=8, num_stages=4, l2_groupings=[8]),
  }
   
   
  def _make_kernel(config: helion.Config):
-     @helion.kernel(static_shapes=True, dot_precision="ieee", config=config)
-     def kernel(
-         k: torch.Tensor,
-         v: torch.Tensor,
+     @helion.kernel(dot_precision="ieee", config=config)
+     def kernel(
+         k: torch.Tensor,
+         v: torch.Tensor,
          beta: torch.Tensor,
          A: torch.Tensor,
          beta_g: torch.Tensor,
      ) -> tuple[torch.Tensor, torch.Tensor]:
          B, T, H, K = k.shape
          V = v.shape[-1]
-         C = 64
+         C = hl.specialize(A.shape[-1])
+         K = hl.specialize(K)
+         V = hl.specialize(V)
   
          w_out = torch.empty_like(k)
          u_out = torch.empty_like(v)
  ⋯ 25 unchanged lines
      return kernel
   
   
- _KERNEL_CACHE: dict[tuple[int, int, int, int, int], callable] = {}
+ _KERNELS = {kv: _make_kernel(cfg) for kv, cfg in KV_CONFIGS.items()}
   
   
- def _get_kernel(shape: tuple[int, int, int, int, int]):
-     kernel = _KERNEL_CACHE.get(shape)
-     if kernel is None:
-         kernel = _make_kernel(SHAPE_CONFIGS[shape])
-         _KERNEL_CACHE[shape] = kernel
-     return kernel
-  
-  
  def custom_kernel(data: input_t) -> output_t:
      k, v, beta, A, g = data
-     B, T, H, K = k.shape
+     K = k.shape[-1]
      V = v.shape[-1]
-     # Precompute beta * exp(g) on device before kernel launch
      beta_g = beta * torch.exp(g)
-     kernel = _get_kernel((B, T, H, K, V))
+     kernel = _KERNELS[(K, V)]
      return kernel(k, v, beta, A, beta_g)

scrolls · 94 diff lines total

Best evidence level for this revision: reported

JSON