Skip to content
KernelIndex
Search⌘K

submission 66690

Saint of the Famished · python · License unknown

Use it

Vendorable · source mirrored · license unknownView source →

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

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

Reported · How evidence levels are derived →

Source and license

sourceavailable
revision digestsha256:36cd101ca19af9b8d47d6b0b526bb15ae12a391b971148c8f700c3a173e98688
license declaredunknown
license concludedunknown
authorsSaint of the Famished
imported2026-08-15

Kernel source

submission.py41 lines
#!POPCORN leaderboard vectorsum_v2
#!POPCORN gpus B200

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


@triton.jit
def sum_kernel(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)
    block_sum = tl.sum(x, axis=0)
    tl.store(partial_sums_ptr + pid, block_sum)


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

    BLOCK_SIZE = 4096

    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)

    if partial_sums.numel() > 16384:
        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)
        partial_sums = next_level

    result = partial_sums.sum()
    return result
scrolls · 41 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 66686.

⋯ 22 unchanged lines
input, output = data
n_elements = input.numel()
- BLOCK_SIZE = 1024
+ BLOCK_SIZE = 4096
+
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)
+ if partial_sums.numel() > 16384:
+ 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)
+ partial_sums = next_level
+
result = partial_sums.sum()
return result

Best evidence level for this revision: reported

JSON