submission 614301 · KernelIndex

Use it

Vendorable · source mirrored · license unknownView source →

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

submission.py

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

11.5µs

#3 of 54

2026-03-23

Run →

Reported · How evidence levels are derived →

Source and license

sourceavailable

revision digestsha256:cb3bc116a25dd1a19b82fe54b495af07b47db8edfddd6558b8df366934a44f19

license declaredunknown

license concludedunknown

authorsdannywillowliu-uchi

imported2026-08-15

Techniques

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

shared-memory__shared__ uint32_t smem[256];

vector-width = int4((int4*)output)[threadIdx.x] = make_int4(0, 0, 0, 0);

Kernel source

submission.py124 lines

import os
os.environ["CUBLAS_WORKSPACE_CONFIG"] = ":4096:8"

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

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

__global__ __launch_bounds__(512)
void histogram_custom(
    const uint8_t* __restrict__ data,
    int64_t* __restrict__ output,
    const int64_t n
) {
    __shared__ uint32_t smem[256];

    // Fused zero: block 0 zeros output using vectorized int4 writes
    if (blockIdx.x == 0 && threadIdx.x < 128) {
        ((int4*)output)[threadIdx.x] = make_int4(0, 0, 0, 0);
    }

    // Zero shared memory (256 uint32s, only need first 256 threads)
    if (threadIdx.x < 256) {
        smem[threadIdx.x] = 0;
    }
    __syncthreads();

    const int64_t tid = blockIdx.x * 512 + threadIdx.x;
    const int64_t grid_stride = 512LL * gridDim.x;

    const int64_t n16 = n >> 4;
    const uint4* data_vec = (const uint4*)data;

    for (int64_t i = tid; i < n16; i += grid_stride) {
        uint4 vals = __ldg(&data_vec[i]);

        atomicAdd(&smem[(vals.x      ) & 0xFF], 1u);
        atomicAdd(&smem[(vals.x >>  8) & 0xFF], 1u);
        atomicAdd(&smem[(vals.x >> 16) & 0xFF], 1u);
        atomicAdd(&smem[(vals.x >> 24)       ], 1u);
        atomicAdd(&smem[(vals.y      ) & 0xFF], 1u);
        atomicAdd(&smem[(vals.y >>  8) & 0xFF], 1u);
        atomicAdd(&smem[(vals.y >> 16) & 0xFF], 1u);
        atomicAdd(&smem[(vals.y >> 24)       ], 1u);
        atomicAdd(&smem[(vals.z      ) & 0xFF], 1u);
        atomicAdd(&smem[(vals.z >>  8) & 0xFF], 1u);
        atomicAdd(&smem[(vals.z >> 16) & 0xFF], 1u);
        atomicAdd(&smem[(vals.z >> 24)       ], 1u);
        atomicAdd(&smem[(vals.w      ) & 0xFF], 1u);
        atomicAdd(&smem[(vals.w >>  8) & 0xFF], 1u);
        atomicAdd(&smem[(vals.w >> 16) & 0xFF], 1u);
        atomicAdd(&smem[(vals.w >> 24)       ], 1u);
    }

    {
        int64_t tail_start = n16 * 16;
        for (int64_t i = tail_start + tid; i < n; i += grid_stride) {
            atomicAdd(&smem[__ldg(&data[i])], 1u);
        }
    }

    __syncthreads();

    if (threadIdx.x < 256) {
        uint32_t v = smem[threadIdx.x];
        if (v > 0) {
            atomicAdd((unsigned long long*)&output[threadIdx.x], (unsigned long long)v);
        }
    }
}

static int _sm_count = -1;

int get_sm_count() {
    if (_sm_count < 0) {
        cudaDeviceProp prop;
        cudaGetDeviceProperties(&prop, 0);
        _sm_count = prop.multiProcessorCount;
    }
    return _sm_count;
}

torch::Tensor histogram_cuda(torch::Tensor data, torch::Tensor output) {
    const int64_t n = data.numel();

    histogram_custom<<<get_sm_count(), 512>>>(
        data.data_ptr<uint8_t>(),
        output.data_ptr<int64_t>(),
        n
    );

    return output;
}
""";

cpp_source = r"""
torch::Tensor histogram_cuda(torch::Tensor data, torch::Tensor output);
""";

module = load_inline(
    name="histogram_submit_v2",
    cpp_sources=[cpp_source],
    cuda_sources=[cuda_source],
    functions=["histogram_cuda"],
    verbose=False,
    extra_cuda_cflags=["-O3", "--use_fast_math"],
)

_wd = torch.randint(0, 256, (1024,), device="cuda", dtype=torch.uint8)
_wo = torch.zeros(256, device="cuda", dtype=torch.int64)
module.histogram_cuda(_wd, _wo)
torch.cuda.synchronize()
del _wd, _wo


def custom_kernel(data: input_t) -> output_t:
    data_tensor, output = data
    module.histogram_cuda(data_tensor, output)
    return output

scrolls · 124 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 614241.

  ⋯ 15 unchanged lines
      int64_t* __restrict__ output,
      const int64_t n
  ) {
-     constexpr int WSUB = 8;
-     __shared__ uint32_t smem[WSUB * 256];
+     __shared__ uint32_t smem[256];
   
-     const int warp_id = threadIdx.x >> 5;
-  
-     // Fused zero: block 0 zeros output (eliminates separate memset kernel)
-     if (blockIdx.x == 0 && threadIdx.x < 256) {
-         output[threadIdx.x] = 0;
+     // Fused zero: block 0 zeros output using vectorized int4 writes
+     if (blockIdx.x == 0 && threadIdx.x < 128) {
+         ((int4*)output)[threadIdx.x] = make_int4(0, 0, 0, 0);
      }
   
-     #pragma unroll
-     for (int i = threadIdx.x; i < WSUB * 256; i += 512) {
-         smem[i] = 0;
+     // Zero shared memory (256 uint32s, only need first 256 threads)
+     if (threadIdx.x < 256) {
+         smem[threadIdx.x] = 0;
      }
      __syncthreads();
   
-     if (blockIdx.x == 0 && threadIdx.x == 0) {
-         __threadfence();
-     }
-  
-     uint32_t* my_hist = smem + (warp_id & 7) * 256;
-  
      const int64_t tid = blockIdx.x * 512 + threadIdx.x;
      const int64_t grid_stride = 512LL * gridDim.x;
   
  ⋯ 3 unchanged lines
      for (int64_t i = tid; i < n16; i += grid_stride) {
          uint4 vals = __ldg(&data_vec[i]);
   
-         atomicAdd(&my_hist[(vals.x      ) & 0xFF], 1u);
-         atomicAdd(&my_hist[(vals.x >>  8) & 0xFF], 1u);
-         atomicAdd(&my_hist[(vals.x >> 16) & 0xFF], 1u);
-         atomicAdd(&my_hist[(vals.x >> 24)       ], 1u);
-         atomicAdd(&my_hist[(vals.y      ) & 0xFF], 1u);
-         atomicAdd(&my_hist[(vals.y >>  8) & 0xFF], 1u);
-         atomicAdd(&my_hist[(vals.y >> 16) & 0xFF], 1u);
-         atomicAdd(&my_hist[(vals.y >> 24)       ], 1u);
-         atomicAdd(&my_hist[(vals.z      ) & 0xFF], 1u);
-         atomicAdd(&my_hist[(vals.z >>  8) & 0xFF], 1u);
-         atomicAdd(&my_hist[(vals.z >> 16) & 0xFF], 1u);
-         atomicAdd(&my_hist[(vals.z >> 24)       ], 1u);
-         atomicAdd(&my_hist[(vals.w      ) & 0xFF], 1u);
-         atomicAdd(&my_hist[(vals.w >>  8) & 0xFF], 1u);
-         atomicAdd(&my_hist[(vals.w >> 16) & 0xFF], 1u);
-         atomicAdd(&my_hist[(vals.w >> 24)       ], 1u);
+         atomicAdd(&smem[(vals.x      ) & 0xFF], 1u);
+         atomicAdd(&smem[(vals.x >>  8) & 0xFF], 1u);
+         atomicAdd(&smem[(vals.x >> 16) & 0xFF], 1u);
+         atomicAdd(&smem[(vals.x >> 24)       ], 1u);
+         atomicAdd(&smem[(vals.y      ) & 0xFF], 1u);
+         atomicAdd(&smem[(vals.y >>  8) & 0xFF], 1u);
+         atomicAdd(&smem[(vals.y >> 16) & 0xFF], 1u);
+         atomicAdd(&smem[(vals.y >> 24)       ], 1u);
+         atomicAdd(&smem[(vals.z      ) & 0xFF], 1u);
+         atomicAdd(&smem[(vals.z >>  8) & 0xFF], 1u);
+         atomicAdd(&smem[(vals.z >> 16) & 0xFF], 1u);
+         atomicAdd(&smem[(vals.z >> 24)       ], 1u);
+         atomicAdd(&smem[(vals.w      ) & 0xFF], 1u);
+         atomicAdd(&smem[(vals.w >>  8) & 0xFF], 1u);
+         atomicAdd(&smem[(vals.w >> 16) & 0xFF], 1u);
+         atomicAdd(&smem[(vals.w >> 24)       ], 1u);
      }
   
      {
          int64_t tail_start = n16 * 16;
          for (int64_t i = tail_start + tid; i < n; i += grid_stride) {
-             atomicAdd(&my_hist[__ldg(&data[i])], 1u);
+             atomicAdd(&smem[__ldg(&data[i])], 1u);
          }
      }
   
      __syncthreads();
   
      if (threadIdx.x < 256) {
-         uint32_t total = 0;
-         #pragma unroll
-         for (int w = 0; w < WSUB; w++) {
-             total += smem[w * 256 + threadIdx.x];
+         uint32_t v = smem[threadIdx.x];
+         if (v > 0) {
+             atomicAdd((unsigned long long*)&output[threadIdx.x], (unsigned long long)v);
          }
-         if (total > 0) {
-             atomicAdd((unsigned long long*)&output[threadIdx.x], (unsigned long long)total);
-         }
      }
  }
   
  ⋯ 11 unchanged lines
  torch::Tensor histogram_cuda(torch::Tensor data, torch::Tensor output) {
      const int64_t n = data.numel();
   
-     int blocks = get_sm_count();
-  
-     histogram_custom<<<blocks, 512>>>(
+     histogram_custom<<<get_sm_count(), 512>>>(
          data.data_ptr<uint8_t>(),
          output.data_ptr<int64_t>(),
          n
  ⋯ 8 unchanged lines
  """;
   
  module = load_inline(
-     name="histogram_submit_fused",
+     name="histogram_submit_v2",
      cpp_sources=[cpp_source],
      cuda_sources=[cuda_source],
      functions=["histogram_cuda"],

scrolls · 119 diff lines total

Best evidence level for this revision: reported

JSON