submission 677145 · KernelIndex

Use it

Vendorable · source mirrored · license unknownView source →

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

submission.py

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

12.2µs

#6 of 54

2026-03-31

Run →

Reported · How evidence levels are derived →

Source and license

sourceavailable

revision digestsha256:0a50f2f3012571dedcee01bedd2ef4ea2528ea8dc83f8cb8d8a97ae728dfbb66

license declaredunknown

license concludedunknown

authorsngolhn

imported2026-08-15

Techniques

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

shared-memoryextern __shared__ unsigned int smem_hist[];

vector-width = uint4reinterpret_cast<uint4*>(output)[tid] = make_uint4(0u, 0u, 0u, 0u);

Kernel source

submission.py86 lines

#!POPCORN leaderboard histogram_v2
#!POPCORN gpu B200

import torch
from task import input_t, output_t
from torch.utils.cpp_extension import load_inline

cuda_src = r"""
#include <torch/extension.h>
#include <cuda_runtime.h>

__global__ void __launch_bounds__(256, 4)
histogram_kernel(const uint8_t* __restrict__ data, int64_t* __restrict__ output, int N) {
    extern __shared__ unsigned int smem_hist[];
    const int tid = threadIdx.x;

    smem_hist[tid] = 0u;

    // Vectorized output zeroing with uint4 (16 bytes per store)
    if (blockIdx.x == 0) {
        if (tid < 128) {
            reinterpret_cast<uint4*>(output)[tid] = make_uint4(0u, 0u, 0u, 0u);
        }
    }
    __syncthreads();

    const int vec_n = N >> 4;
    const uint4* data_vec = reinterpret_cast<const uint4*>(data);
    int idx = blockIdx.x * blockDim.x + tid;
    const int stride = blockDim.x * gridDim.x;

    // Process pairs of uint4 (32 bytes = 32 elements per iteration)
    const int vec_n_pairs = vec_n >> 1;
    for (int i = idx; i < vec_n_pairs; i += stride) {
        uint4 val0 = __ldg(&data_vec[i * 2]);
        uint4 val1 = __ldg(&data_vec[i * 2 + 1]);
        const uint8_t* b0 = reinterpret_cast<const uint8_t*>(&val0);
        const uint8_t* b1 = reinterpret_cast<const uint8_t*>(&val1);
        #pragma unroll
        for (int j = 0; j < 16; j++) {
            atomicAdd(&smem_hist[b0[j]], 1u);
        }
        #pragma unroll
        for (int j = 0; j < 16; j++) {
            atomicAdd(&smem_hist[b1[j]], 1u);
        }
    }

    __syncthreads();

    if (smem_hist[tid] > 0u) {
        atomicAdd(reinterpret_cast<unsigned long long*>(&output[tid]),
                  static_cast<unsigned long long>(smem_hist[tid]));
    }
}

void histogram_inplace(torch::Tensor data, torch::Tensor output) {
    const int N = data.numel();
    int num_blocks = min(256, max(1, (N + 256*32 - 1) / (256*32)));
    histogram_kernel<<<num_blocks, 256, 256*sizeof(unsigned int)>>>(
        data.data_ptr<uint8_t>(), output.data_ptr<int64_t>(), N);
}
"""

cpp_src = r"""
void histogram_inplace(torch::Tensor data, torch::Tensor output);
"""

_ext = load_inline(
    name="histogram_sm100_notail",
    cpp_sources=cpp_src,
    cuda_sources=cuda_src,
    functions=["histogram_inplace"],
    with_cuda=True,
    extra_cflags=["-O3", "-std=c++17"],
    extra_cuda_cflags=["-O3", "--use_fast_math", "-std=c++17",
                       "-gencode=arch=compute_100,code=sm_100"],
    verbose=False,
)


def custom_kernel(data: input_t) -> output_t:
    data_tensor, output_tensor = data
    _ext.histogram_inplace(data_tensor, output_tensor)
    return output_tensor

scrolls · 86 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 676997.

  ⋯ 14 unchanged lines
      const int tid = threadIdx.x;
   
      smem_hist[tid] = 0u;
+  
+     // Vectorized output zeroing with uint4 (16 bytes per store)
      if (blockIdx.x == 0) {
-         output[tid] = 0;
+         if (tid < 128) {
+             reinterpret_cast<uint4*>(output)[tid] = make_uint4(0u, 0u, 0u, 0u);
+         }
      }
      __syncthreads();
   
  ⋯ 2 unchanged lines
      int idx = blockIdx.x * blockDim.x + tid;
      const int stride = blockDim.x * gridDim.x;
   
-     // 2x uint4 loads with __ldg for read-only cache path
-     const int vec_n_pairs = vec_n & ~1;
-     for (int i = idx * 2; i < vec_n_pairs; i += stride * 2) {
-         uint4 val0 = __ldg(&data_vec[i]);
-         uint4 val1 = __ldg(&data_vec[i + 1]);
+     // Process pairs of uint4 (32 bytes = 32 elements per iteration)
+     const int vec_n_pairs = vec_n >> 1;
+     for (int i = idx; i < vec_n_pairs; i += stride) {
+         uint4 val0 = __ldg(&data_vec[i * 2]);
+         uint4 val1 = __ldg(&data_vec[i * 2 + 1]);
          const uint8_t* b0 = reinterpret_cast<const uint8_t*>(&val0);
          const uint8_t* b1 = reinterpret_cast<const uint8_t*>(&val1);
          #pragma unroll
  ⋯ 6 unchanged lines
          }
      }
   
-     // Handle odd vec element
-     if ((vec_n & 1) && idx == 0) {
-         uint4 val = __ldg(&data_vec[vec_n - 1]);
-         const uint8_t* b = reinterpret_cast<const uint8_t*>(&val);
-         #pragma unroll
-         for (int j = 0; j < 16; j++) {
-             atomicAdd(&smem_hist[b[j]], 1u);
-         }
-     }
-  
-     // Byte-level tail
-     int tail_start = vec_n * 16;
-     for (int i = tail_start + idx; i < N; i += stride) {
-         atomicAdd(&smem_hist[__ldg(&data[i])], 1u);
-     }
-  
      __syncthreads();
   
      if (smem_hist[tid] > 0u) {
  ⋯ 15 unchanged lines
  """
   
  _ext = load_inline(
-     name="histogram_ldg_2xu4_256blk",
+     name="histogram_sm100_notail",
      cpp_sources=cpp_src,
      cuda_sources=cuda_src,
      functions=["histogram_inplace"],
      with_cuda=True,
      extra_cflags=["-O3", "-std=c++17"],
-     extra_cuda_cflags=["-O3", "--use_fast_math", "-std=c++17"],
+     extra_cuda_cflags=["-O3", "--use_fast_math", "-std=c++17",
+                        "-gencode=arch=compute_100,code=sm_100"],
      verbose=False,
  )

scrolls · 71 diff lines total

Best evidence level for this revision: reported

JSON