Skip to content
KernelIndex
Search⌘K
Sign in

submission 66970

Saint of the Famished · python · License unknown
Kernel source · 47 lines ↓

Use it

Vendorable · source mirrored · license unknownView source →

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

submission.py
curl "https://kernelindex.com/api/v1/implementations/kernelbot-vectorsum-v2-66970?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
47.5µs
#18 of 88
2025-11-05
Run →

Reported · How evidence levels are derived →

Source and license

sourceavailable
revision digestsha256:340b78f302f3d01770690f4be83e4fcf042c253fee6aa657b12d995cf99834d8
license declaredunknown
license concludedunknown
authorsSaint of the Famished
imported2026-08-15

Kernel source

submission.py47 lines
#!POPCORN leaderboard vectorsum_v2

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


@triton.jit
def sum_kernel(
    x_ptr,
    n_elements,
    BLOCK_SIZE: tl.constexpr,
):
    pid = tl.program_id(0)
    block_start = pid * BLOCK_SIZE
    offsets = block_start + tl.arange(0, BLOCK_SIZE)
    mask = offsets < n_elements
    x = tl.load(x_ptr + offsets, mask=mask, other=0.0, eviction_policy="evict_first")
    block_sum = tl.sum(x, axis=0)
    tl.store(x_ptr + pid, block_sum)


def custom_kernel(data: input_t) -> output_t:
    input, output = data
    n_elements = input.numel()

    if n_elements >= 10_000_000:
        BLOCK_SIZE = 8192
    elif n_elements >= 1_000_000:
        BLOCK_SIZE = 4096
    elif n_elements >= 100_000:
        BLOCK_SIZE = 2048
    elif n_elements >= 10_000:
        BLOCK_SIZE = 1024
    else:
        BLOCK_SIZE = 512

    grid = lambda meta: (triton.cdiv(n_elements, meta["BLOCK_SIZE"]),)

    n_blocks = triton.cdiv(n_elements, BLOCK_SIZE)
    grid = (n_blocks,)

    sum_kernel[grid](input, n_elements, BLOCK_SIZE=BLOCK_SIZE)

    return input[:n_blocks].sum()
scrolls · 47 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 66966.

  ⋯ 8 unchanged lines
  @triton.jit
  def sum_kernel(
      x_ptr,
-     partial_sums_ptr,
      n_elements,
      BLOCK_SIZE: tl.constexpr,
  ):
  ⋯ 3 unchanged lines
      mask = offsets < n_elements
      x = tl.load(x_ptr + offsets, mask=mask, other=0.0, eviction_policy="evict_first")
      block_sum = tl.sum(x, axis=0)
-     tl.store(partial_sums_ptr + pid, block_sum)
+     tl.store(x_ptr + pid, block_sum)
   
   
  def custom_kernel(data: input_t) -> output_t:
  ⋯ 11 unchanged lines
      else:
          BLOCK_SIZE = 512
   
-     n_blocks = triton.cdiv(n_elements, BLOCK_SIZE)
-     partial_sums = torch.empty(n_blocks, device=input.device, dtype=input.dtype)
+     grid = lambda meta: (triton.cdiv(n_elements, meta["BLOCK_SIZE"]),)
   
+     n_blocks = triton.cdiv(n_elements, BLOCK_SIZE)
      grid = (n_blocks,)
-     sum_kernel[grid](input, partial_sums, n_elements, BLOCK_SIZE=BLOCK_SIZE)
   
-     return partial_sums.sum()
+     sum_kernel[grid](input, n_elements, BLOCK_SIZE=BLOCK_SIZE)
+  
+     return input[:n_blocks].sum()
scrolls · 33 diff lines total

Best evidence level for this revision: reported

JSON