submission 66696 · KernelIndex

Use it

Vendorable · source mirrored · license unknownView source →

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

submission.py

curl "https://kernelindex.com/api/v1/implementations/kernelbot-vectorsum-v2-66696?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.5µs

#42 of 88

2025-11-05

Run →

Reported · How evidence levels are derived →

Source and license

sourceavailable

revision digestsha256:972cde549126d76500a01d3478c9a0d43ac64a485fad25e1e2a7af3e16830471

license declaredunknown

license concludedunknown

authorsSaint of the Famished

imported2026-08-15

Kernel source

submission.py77 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_optimized(
    x_ptr,
    partial_sums_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(partial_sums_ptr + pid, block_sum)


@triton.jit
def sum_kernel_atomic(
    x_ptr,
    output_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.atomic_add(output_ptr, block_sum)


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

    BLOCK_SIZE = 8192

    n_blocks = triton.cdiv(n_elements, BLOCK_SIZE)
    partial_sums = torch.empty(n_blocks, device=input.device, dtype=input.dtype)

    grid = (n_blocks,)
    sum_kernel_optimized[grid](input, partial_sums, n_elements, BLOCK_SIZE=BLOCK_SIZE)

    while partial_sums.numel() > BLOCK_SIZE:
        n_elements = partial_sums.numel()
        n_blocks = triton.cdiv(n_elements, BLOCK_SIZE)
        next_level = torch.empty(n_blocks, device=input.device, dtype=input.dtype)

        grid = (n_blocks,)
        sum_kernel_optimized[grid](partial_sums, next_level, n_elements, BLOCK_SIZE=BLOCK_SIZE)
        partial_sums = next_level

    if partial_sums.numel() <= 128:
        final_output = torch.zeros(1, device=input.device, dtype=input.dtype)
        n_elements = partial_sums.numel()
        sum_kernel_atomic[(1,)](partial_sums, final_output, n_elements, BLOCK_SIZE=BLOCK_SIZE)
        return final_output[0]
    else:
        n_elements = partial_sums.numel()
        n_blocks = triton.cdiv(n_elements, BLOCK_SIZE)
        next_level = torch.empty(n_blocks, device=input.device, dtype=input.dtype)

        grid = (n_blocks,)
        sum_kernel_optimized[grid](partial_sums, next_level, n_elements, BLOCK_SIZE=BLOCK_SIZE)
        return next_level.sum()

scrolls · 77 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 66694.

  ⋯ 6 unchanged lines
   
   
  @triton.jit
- def sum_kernel(x_ptr, partial_sums_ptr, n_elements, BLOCK_SIZE: tl.constexpr):
+ def sum_kernel_optimized(
+     x_ptr,
+     partial_sums_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)
+     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)
   
   
+ @triton.jit
+ def sum_kernel_atomic(
+     x_ptr,
+     output_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.atomic_add(output_ptr, block_sum)
+  
+  
  def custom_kernel(data: input_t) -> output_t:
      input, output = data
      n_elements = input.numel()
   
-     BLOCK_SIZE = 4096
+     BLOCK_SIZE = 8192
   
      n_blocks = triton.cdiv(n_elements, BLOCK_SIZE)
      partial_sums = torch.empty(n_blocks, device=input.device, dtype=input.dtype)
-     sum_kernel[(n_blocks,)](input, partial_sums, n_elements, BLOCK_SIZE=BLOCK_SIZE)
   
-     # Keep reducing on GPU while there are many elements left
-     # Stop when small enough that CPU reduction is faster
+     grid = (n_blocks,)
+     sum_kernel_optimized[grid](input, partial_sums, n_elements, BLOCK_SIZE=BLOCK_SIZE)
+  
      while partial_sums.numel() > BLOCK_SIZE:
          n_elements = partial_sums.numel()
          n_blocks = triton.cdiv(n_elements, BLOCK_SIZE)
          next_level = torch.empty(n_blocks, device=input.device, dtype=input.dtype)
-         sum_kernel[(n_blocks,)](partial_sums, next_level, n_elements, BLOCK_SIZE=BLOCK_SIZE)
+  
+         grid = (n_blocks,)
+         sum_kernel_optimized[grid](partial_sums, next_level, n_elements, BLOCK_SIZE=BLOCK_SIZE)
          partial_sums = next_level
   
-     # Final small sum on CPU is faster than launching another tiny kernel
-     return partial_sums.sum()
+     if partial_sums.numel() <= 128:
+         final_output = torch.zeros(1, device=input.device, dtype=input.dtype)
+         n_elements = partial_sums.numel()
+         sum_kernel_atomic[(1,)](partial_sums, final_output, n_elements, BLOCK_SIZE=BLOCK_SIZE)
+         return final_output[0]
+     else:
+         n_elements = partial_sums.numel()
+         n_blocks = triton.cdiv(n_elements, BLOCK_SIZE)
+         next_level = torch.empty(n_blocks, device=input.device, dtype=input.dtype)
+  
+         grid = (n_blocks,)
+         sum_kernel_optimized[grid](partial_sums, next_level, n_elements, BLOCK_SIZE=BLOCK_SIZE)
+         return next_level.sum()

scrolls · 81 diff lines total

Best evidence level for this revision: reported

JSON