submission 776241 · KernelIndex

Use it

Vendorable · source mirrored · license unknownView source →

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

submission.py

curl "https://kernelindex.com/api/v1/implementations/kernelbot-vectorsum-v2-776241?include=source"

interfacepython

Compatibility

measured onNVIDIA A100

declared hardwareNVIDIA A100

architecturessm_80

dtypesfp32

Benchmark evidence

1 measurement across 1 GPU, fastest first.

Operation / workload

Hardware

Latency

Rank

Observed

Vector sum reductionsuite of 6 cases

NVIDIA A100

142.9µs

#26 of 96

2026-04-20

Run →

Reported · How evidence levels are derived →

Source and license

sourceavailable

revision digestsha256:a0b7268f7bdc8046338710e8fc3025be76f1cc68b96981ccb6e7edb08985c889

license declaredunknown

license concludedunknown

authorsx3C49

imported2026-08-15

Techniques

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

num-warps = 16num_warps=16, # 512 threads/block → 4 blocks/SM → 64 warps/SM = 100% occupancy

Kernel source

submission.py63 lines

import torch
import triton
import triton.language as tl
from task import input_t, output_t

GRID = 512
BLOCK = 4096


@triton.jit
def _reduce_pass1(
    data_ptr, partial_ptr, N,
    BLOCK: tl.constexpr,
    GRID: tl.constexpr,
):
    pid = tl.program_id(0)
    acc = tl.zeros((BLOCK,), dtype=tl.float64)
    start = pid * BLOCK
    stride = GRID * BLOCK

    # ── hot path: full tiles, no mask ──────────────────────────────────────
    while start + BLOCK <= N:
        acc += tl.load(data_ptr + start + tl.arange(0, BLOCK)).to(tl.float64)
        start += stride

    # ── tail: partial tile (not reached for any benchmark size) ────────────
    if start < N:
        offs = start + tl.arange(0, BLOCK)
        acc += tl.load(data_ptr + offs, mask=offs < N, other=0.0).to(tl.float64)

    tl.store(partial_ptr + pid, tl.sum(acc, axis=0))


@triton.jit
def _reduce_pass2(partial_ptr, out_ptr, GRID: tl.constexpr):
    s = tl.sum(tl.load(partial_ptr + tl.arange(0, GRID)), axis=0)
    tl.store(out_ptr, s.to(tl.float32))


_scratch: torch.Tensor | None = None


def custom_kernel(data: input_t) -> output_t:
    global _scratch
    input_tensor, output_tensor = data
    N = input_tensor.numel()

    if _scratch is None:
        _scratch = torch.empty(GRID, device="cuda", dtype=torch.float64)

    _reduce_pass1[(GRID,)](
        input_tensor, _scratch, N,
        BLOCK=BLOCK,
        GRID=GRID,
        num_warps=16,   # 512 threads/block → 4 blocks/SM → 64 warps/SM = 100% occupancy
    )
    _reduce_pass2[(1,)](
        _scratch, output_tensor,
        GRID=GRID,
        num_warps=4,
    )

    return output_tensor.view([])

scrolls · 63 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 775031.

  ⋯ 7 unchanged lines
   
   
  @triton.jit
- def _pass1(data_ptr, partial_ptr, N,
-             BLOCK_SIZE: tl.constexpr, GRID: tl.constexpr):
+ def _reduce_pass1(
+     data_ptr, partial_ptr, N,
+     BLOCK: tl.constexpr,
+     GRID: tl.constexpr,
+ ):
      pid = tl.program_id(0)
-     acc = tl.zeros((BLOCK_SIZE,), dtype=tl.float64)
-     block_start = pid * BLOCK_SIZE
-     step = GRID * BLOCK_SIZE
-     while block_start < N:
-         offs = block_start + tl.arange(0, BLOCK_SIZE)
-         mask = offs < N
-         vals = tl.load(data_ptr + offs, mask=mask, other=0.0)
-         acc += vals.to(tl.float64)
-         block_start += step
-     partial = tl.sum(acc, axis=0)
-     tl.store(partial_ptr + pid, partial)
+     acc = tl.zeros((BLOCK,), dtype=tl.float64)
+     start = pid * BLOCK
+     stride = GRID * BLOCK
   
+     # ── hot path: full tiles, no mask ──────────────────────────────────────
+     while start + BLOCK <= N:
+         acc += tl.load(data_ptr + start + tl.arange(0, BLOCK)).to(tl.float64)
+         start += stride
   
+     # ── tail: partial tile (not reached for any benchmark size) ────────────
+     if start < N:
+         offs = start + tl.arange(0, BLOCK)
+         acc += tl.load(data_ptr + offs, mask=offs < N, other=0.0).to(tl.float64)
+  
+     tl.store(partial_ptr + pid, tl.sum(acc, axis=0))
+  
+  
  @triton.jit
- def _pass2(partial_ptr, output_ptr, N_PARTIALS: tl.constexpr):
-     offs = tl.arange(0, N_PARTIALS)
-     vals = tl.load(partial_ptr + offs)
-     total = tl.sum(vals, axis=0)
-     tl.store(output_ptr, total.to(tl.float32))
+ def _reduce_pass2(partial_ptr, out_ptr, GRID: tl.constexpr):
+     s = tl.sum(tl.load(partial_ptr + tl.arange(0, GRID)), axis=0)
+     tl.store(out_ptr, s.to(tl.float32))
   
   
  _scratch: torch.Tensor | None = None
  ⋯ 7 unchanged lines
      if _scratch is None:
          _scratch = torch.empty(GRID, device="cuda", dtype=torch.float64)
   
-     _pass1[(GRID,)](
-         input_tensor,
-         _scratch,
-         N,
-         BLOCK_SIZE=BLOCK,
+     _reduce_pass1[(GRID,)](
+         input_tensor, _scratch, N,
+         BLOCK=BLOCK,
          GRID=GRID,
-         num_warps=8,
+         num_warps=16,   # 512 threads/block → 4 blocks/SM → 64 warps/SM = 100% occupancy
      )
-  
-     _pass2[(1,)](
-         _scratch,
-         output_tensor,
-         GRID,
+     _reduce_pass2[(1,)](
+         _scratch, output_tensor,
+         GRID=GRID,
          num_warps=4,
      )

scrolls · 79 diff lines total

Best evidence level for this revision: reported

JSON