submission 510381 · KernelIndex

Use it

Vendorable · source mirrored · license unknownView source →

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

submission_atomic_chunked_v3_w4.py

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

interfacepython

Compatibility

measured onNVIDIA L4

declared hardwareNVIDIA L4

architecturessm_89

dtypesfp32

Benchmark evidence

1 measurement across 1 GPU, fastest first.

Operation / workload

Hardware

Latency

Rank

Observed

Vector sum reductionsuite of 6 cases

NVIDIA L4

944.5µs

#9 of 26

2026-02-27

Run →

Reported · How evidence levels are derived →

Source and license

sourceavailable

revision digestsha256:3fc0cb11766d181e1ea6257e6ddad5484239f8b96183f2fae5061cf6c1f4b1d6

license declaredunknown

license concludedunknown

authorsiharryli

imported2026-08-15

Techniques

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

num-warps = 1_zero_scalar[(1,)](out, num_warps=1, num_stages=1)

stages = 1_zero_scalar[(1,)](out, num_warps=1, num_stages=1)

Kernel source

submission_atomic_chunked_v3_w4.py64 lines


import triton
import triton.language as tl

from task import input_t, output_t


@triton.jit
def _zero_scalar(out_ptr):
    tl.store(out_ptr, 0.0)


@triton.jit
def _sum_atomic_chunked(
    x_ptr,
    out_ptr,
    n_elements,
    BLOCK: tl.constexpr,
    ITERS: tl.constexpr,
):
    pid = tl.program_id(0)
    base = pid * BLOCK * ITERS
    tl.multiple_of(base, 256)
    r = tl.arange(0, BLOCK)

    acc = tl.zeros((), dtype=tl.float32)
    for i in tl.static_range(0, ITERS):
        offsets = base + i * BLOCK + r
        x = tl.load(x_ptr + offsets, mask=offsets < n_elements, other=0.0, cache_modifier=".cg")
        acc += tl.sum(x, axis=0)

    tl.atomic_add(out_ptr, acc)


def _pick_iters(n: int) -> int:
    if n >= 20_000_000:
        return 16
    if n >= 2_000_000:
        return 8
    return 4


def custom_kernel(data: input_t) -> output_t:
    x, out = data
    n = x.numel()

    _zero_scalar[(1,)](out, num_warps=1, num_stages=1)

    BLOCK = 1024
    iters = _pick_iters(n)
    grid = (triton.cdiv(n, BLOCK * iters),)

    _sum_atomic_chunked[grid](
        x,
        out,
        n,
        BLOCK=BLOCK,
        ITERS=iters,
        num_warps=4,
        num_stages=4,
    )
    return out[0]

scrolls · 64 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 510378.

  ⋯ 32 unchanged lines
   
   
  def _pick_iters(n: int) -> int:
-     # Slightly higher ITERS for mid-size vectors reduces atomic pressure on L4.
      if n >= 20_000_000:
          return 16
      if n >= 2_000_000:
  ⋯ 17 unchanged lines
          n,
          BLOCK=BLOCK,
          ITERS=iters,
-         num_warps=8,
+         num_warps=4,
          num_stages=4,
      )
      return out[0]

Best evidence level for this revision: reported

JSON