submission 510378 · KernelIndex

Use it

Vendorable · source mirrored · license unknownView source →

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

submission_atomic_chunked_v2.py

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

945.5µs

#10 of 26

2026-02-27

Run →

Reported · How evidence levels are derived →

Source and license

sourceavailable

revision digestsha256:60f4d786b28b92fdf4538d7452540803aed01d35d06b2e88211f708a271cd859

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_v2.py65 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:
    # Slightly higher ITERS for mid-size vectors reduces atomic pressure on L4.
    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=8,
        num_stages=4,
    )
    return out[0]

scrolls · 65 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 510371.

   
- import torch
  import triton
  import triton.language as tl
   
  ⋯ 1 unchanged lines
   
   
  @triton.jit
+ def _zero_scalar(out_ptr):
+     tl.store(out_ptr, 0.0)
+  
+  
+ @triton.jit
  def _sum_atomic_chunked(
      x_ptr,
      out_ptr,
  ⋯ 3 unchanged lines
  ):
      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
-         mask = offsets < n_elements
-         x = tl.load(x_ptr + offsets, mask=mask, other=0.0)
+         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)
   
   
- @triton.jit
- def _zero_scalar(out_ptr):
-     tl.store(out_ptr, 0.0)
-  
-  
  def _pick_iters(n: int) -> int:
-     # Reduce atomic traffic as N grows (baseline does 1 atomic per 1024 elems).
+     # Slightly higher ITERS for mid-size vectors reduces atomic pressure on L4.
      if n >= 20_000_000:
          return 16
-     if n >= 6_000_000:
+     if n >= 2_000_000:
          return 8
      return 4
   
  ⋯ 1 unchanged lines
  def custom_kernel(data: input_t) -> output_t:
      x, out = data
      n = x.numel()
-     iters = _pick_iters(n)
   
+     _zero_scalar[(1,)](out, num_warps=1, num_stages=1)
+  
      BLOCK = 1024
+     iters = _pick_iters(n)
      grid = (triton.cdiv(n, BLOCK * iters),)
   
-     # Ensure correct accumulation for atomic reduction.
-     _zero_scalar[(1,)](out, num_warps=1, num_stages=1)
-  
      _sum_atomic_chunked[grid](
          x,
          out,
  ⋯ 4 unchanged lines
          num_stages=4,
      )
      return out[0]
+

scrolls · 73 diff lines total

Best evidence level for this revision: reported

JSON