submission 68488 · KernelIndex

Use it

Vendorable · source mirrored · license unknownView source →

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

submission_triton_v0.py

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

63.5µs

#35 of 54

2025-11-09

Run →

Reported · How evidence levels are derived →

Source and license

sourceavailable

revision digestsha256:809d63b3386b026c56e93c230b36b8150861c4374c90c40e4443588205401537

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.py50 lines

#!POPCORN leaderboard histogram_v2

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


@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)  # old triton doesn't have mask for histogram

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

    # compensation since we use 0 for masked elements
    if pid == 0:
        compensate = size - num_iters * BLOCK_SIZE * num_pids
        tl.atomic_add(output_ptr, compensate)


# NUM_SMS = torch.cuda.get_device_properties().multi_processor_count
NUM_SMS = 264


def custom_kernel(data: input_t) -> output_t:
    data, output = data

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

    return output

scrolls · 50 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 68477.

  #!POPCORN leaderboard histogram_v2
   
  import torch
- from task import input_t, output_t
  import triton
  import triton.language as tl
+ from task import input_t, output_t
   
   
  @triton.jit
  ⋯ 14 unchanged lines
          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?
+         acc += tl.histogram(data, NUM_BINS)  # old triton doesn't have mask for histogram
   
      # NOTE: output_ptr is i64 type
      tl.atomic_add(output_ptr + tl.arange(0, NUM_BINS), acc)
   
+     # compensation since we use 0 for masked elements
      if pid == 0:
          compensate = size - num_iters * BLOCK_SIZE * num_pids
          tl.atomic_add(output_ptr, compensate)
   
   
+ # NUM_SMS = torch.cuda.get_device_properties().multi_processor_count
+ NUM_SMS = 264
+  
+  
  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)
+     kernel[(NUM_SMS,)](data, output, data.shape[0], BLOCK_SIZE)
   
      return output

scrolls · 41 diff lines total

Best evidence level for this revision: reported

JSON