Skip to content
KernelIndex
Search⌘K
Sign in

submission 780485

x3C49 · python · License unknown
Kernel source · 49 lines ↓

Use it

Vendorable · source mirrored · license unknownView source →

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

submission.py
curl "https://kernelindex.com/api/v1/implementations/kernelbot-vectorsum-v2-780485?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
Vector sum reductionsuite of 6 cases
NVIDIA B200
53.7µs
#44 of 88
2026-05-01
Run →

Reported · How evidence levels are derived →

Source and license

sourceavailable
revision digestsha256:32d8498c741b96b025e1c66c30badabc996dec103cee9ddde53a2a814d78ede1
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,

Kernel source

submission.py49 lines
import torch
import triton
import triton.language as tl
from task import input_t, output_t
GRID = 512
BLOCK = 8192

# vectorsum_v2 Kernel Optimization for A100 GPU leaderboard

@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.float32)
    start = pid * BLOCK
    stride = GRID * BLOCK
    while start + BLOCK <= N:
        acc += tl.load(data_ptr + start + tl.arange(0, BLOCK)).to(tl.float32)
        start += stride
    if start < N:
        offs = start + tl.arange(0, BLOCK)
        acc += tl.load(data_ptr + offs, mask=offs < N, other=0.0).to(tl.float32)
    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.float32)
    _reduce_pass1[(GRID,)](
        input_tensor, _scratch, N,
        BLOCK=BLOCK,
        GRID=GRID,
        num_warps=16,
    )
    _reduce_pass2[(1,)](
        _scratch, output_tensor,
        GRID=GRID,
        num_warps=4,
    )
    return output_tensor.view([])
scrolls · 49 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 779751.

  ⋯ 2 unchanged lines
  import triton.language as tl
  from task import input_t, output_t
  GRID = 512
- BLOCK = 8192          # Doubled block size → loop iterations halved (25 → 13)
+ BLOCK = 8192
+  
+ # vectorsum_v2 Kernel Optimization for A100 GPU leaderboard
+  
  @triton.jit
  def _reduce_pass1(
      data_ptr, partial_ptr, N,
  ⋯ 1 unchanged lines
      GRID: tl.constexpr,
  ):
      pid = tl.program_id(0)
-     acc = tl.zeros((BLOCK,), dtype=tl.float32)   # FP32 halves register use
+     acc = tl.zeros((BLOCK,), dtype=tl.float32)
      start = pid * BLOCK
      stride = GRID * BLOCK
      while start + BLOCK <= N:

Best evidence level for this revision: reported

JSON