submission 68477 · KernelIndex

Use it

Vendorable · source mirrored · license unknownView source →

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

submission_triton_v0.py

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

interfacepython

Compatibility

measured onNVIDIA B200

declared hardwareNVIDIA B200

architecturessm_100

dtypesuint8

Benchmark evidence

1 measurement across 1 GPU, fastest first.

Operation / workload

Hardware

Latency

Rank

Observed

Histogramsuite of 6 cases

NVIDIA B200

64.0µs

#36 of 54

2025-11-09

Run →

Reported · How evidence levels are derived →

Source and license

sourceavailable

revision digestsha256:c6af27950a7ecc157cc69cc2b8d484b16b4e41c088c92846917aa5e531f4f3fe

license declaredunknown

license concludedunknown

authorsgau.nernst

imported2026-08-15

Techniques

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

persistent-kernelnum_pids = tl.num_programs(0)

Kernel source

submission_triton_v0.py47 lines

#!POPCORN leaderboard histogram_v2

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


@triton.jit
def kernel(
    data_ptr,  # (size,)
    output_ptr,  # (256,)
    size,
    BLOCK_SIZE: tl.constexpr,
    NUM_BINS: tl.constexpr = 256,
):
    pid = tl.program_id(0)
    num_pids = tl.num_programs(0)

    acc = tl.zeros((NUM_BINS,), dtype=tl.int32)

    num_iters = tl.cdiv(size, BLOCK_SIZE * num_pids)
    for iter_id in range(num_iters):
        offs = iter_id * (num_pids * BLOCK_SIZE) + (pid * BLOCK_SIZE) + tl.arange(0, BLOCK_SIZE)
        mask = offs < size
        data = tl.load(data_ptr + offs, mask, other=0).to(tl.int32)  # tl.histogram() doesn't work with uint8
        acc += tl.histogram(data, NUM_BINS)  # mask doesn't work?

    # NOTE: output_ptr is i64 type
    tl.atomic_add(output_ptr + tl.arange(0, NUM_BINS), acc)

    if pid == 0:
        compensate = size - num_iters * BLOCK_SIZE * num_pids
        tl.atomic_add(output_ptr, compensate)


def custom_kernel(data: input_t) -> output_t:
    data, output = data
    # output[...] = torch.bincount(data, minlength=256)

    BLOCK_SIZE = 2048
    num_blocks = 264
    output.zero_()
    kernel[(num_blocks,)](data, output, data.shape[0], BLOCK_SIZE)

    return output

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 68443.

  ⋯ 1 unchanged lines
   
  import torch
  from task import input_t, output_t
+ import triton
+ import triton.language as tl
   
   
+ @triton.jit
+ def kernel(
+     data_ptr,  # (size,)
+     output_ptr,  # (256,)
+     size,
+     BLOCK_SIZE: tl.constexpr,
+     NUM_BINS: tl.constexpr = 256,
+ ):
+     pid = tl.program_id(0)
+     num_pids = tl.num_programs(0)
+  
+     acc = tl.zeros((NUM_BINS,), dtype=tl.int32)
+  
+     num_iters = tl.cdiv(size, BLOCK_SIZE * num_pids)
+     for iter_id in range(num_iters):
+         offs = iter_id * (num_pids * BLOCK_SIZE) + (pid * BLOCK_SIZE) + tl.arange(0, BLOCK_SIZE)
+         mask = offs < size
+         data = tl.load(data_ptr + offs, mask, other=0).to(tl.int32)  # tl.histogram() doesn't work with uint8
+         acc += tl.histogram(data, NUM_BINS)  # mask doesn't work?
+  
+     # NOTE: output_ptr is i64 type
+     tl.atomic_add(output_ptr + tl.arange(0, NUM_BINS), acc)
+  
+     if pid == 0:
+         compensate = size - num_iters * BLOCK_SIZE * num_pids
+         tl.atomic_add(output_ptr, compensate)
+  
+  
  def custom_kernel(data: input_t) -> output_t:
      data, output = data
-     output[...] = torch.bincount(data, minlength=256)
+     # output[...] = torch.bincount(data, minlength=256)
+  
+     BLOCK_SIZE = 2048
+     num_blocks = 264
+     output.zero_()
+     kernel[(num_blocks,)](data, output, data.shape[0], BLOCK_SIZE)
+  
      return output

scrolls · 47 diff lines total

Best evidence level for this revision: reported

JSON