Skip to content
KernelIndex
Search⌘K

submission 68301

P · python · License unknown

Use it

Vendorable · source mirrored · license unknownView source →

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

submission.py
curl "https://kernelindex.com/api/v1/implementations/kernelbot-histogram-v2-68301?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
281.0µs
#50 of 54
2025-11-08

Reported · How evidence levels are derived →

Source and license

sourceavailable
revision digestsha256:72605cdf970ca07423a0e5fbfd2b73542c15087e68e5e8577ecc2db89c0e7519
license declaredunknown
license concludedunknown
authorsP
imported2026-08-15

Techniques

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

shared-memory__shared__ unsigned long long T[MAX_VAL];

Kernel source

submission.py73 lines
import torch
from utils import DeterministicContext
from torch.utils.cpp_extension import load_inline
from typing import List
from task import input_t, output_t

cuda_histogram = """
#define BLOCK_SIZE 1024
#define MAX_VAL 256
#include <ATen/ATen.h>
#include <torch/extension.h>
#include <cstdint>

__global__ void histogram(const unsigned char* __restrict__ A, unsigned long long* __restrict__ output, int len) {
    __shared__ unsigned long long T[MAX_VAL];

    int tid = threadIdx.x;

    // Initialize shared memory histogram
    for (int i = tid; i < MAX_VAL; i += blockDim.x) T[i] = 0;
    __syncthreads();

    // Build histogram in shared memory
    for (int i = blockIdx.x * blockDim.x + tid; i < len; i += blockDim.x * gridDim.x) {
        atomicAdd(&T[A[i]], 1ULL);
    }
    __syncthreads();

    // Accumulate shared memory histogram into global memory
    for (int i = tid; i < MAX_VAL; i += blockDim.x) {
        atomicAdd(&output[i], T[i]);
    }
}

torch::Tensor& histogram(torch::Tensor& A, torch::Tensor& B) {
    int N = A.numel();

    int blocks = (N + BLOCK_SIZE - 1) / BLOCK_SIZE;
    histogram<<<blocks, BLOCK_SIZE>>>(
        A.data_ptr<unsigned char>(),
        reinterpret_cast<unsigned long long*>(B.data_ptr<int64_t>()),
        N
    );
    return B;
}
"""

abi = torch._C._GLIBCXX_USE_CXX11_ABI
extra_compile_args={'cxx': [f"-O3", f"-D_GLIBCXX_USE_CXX11_ABI={abi}"],
                    'nvcc': ["-O3"]}

histogram_md = load_inline(
    name='sum_cuda_ext',
    cpp_sources="torch::Tensor& histogram(torch::Tensor& A, torch::Tensor& B);",
    cuda_sources=cuda_histogram,
    functions=['histogram'],
    extra_cflags=extra_compile_args['cxx'],
    extra_cuda_cflags=extra_compile_args['nvcc'],
    verbose=True,
)

def custom_kernel(data: input_t) -> output_t:
    """
    Custom implementation of vector addition using CUDA.
    Args:
        inputs: List of pairs of tensors [A, B] to be added.
    Returns:
        Tensor containing element-wise sum.
    """
    A, B = data
    B.zero_()  # Initialize output tensor to zero
    return histogram_md.histogram(A, B)
scrolls · 73 lines total

Source code from GPU Mode and the KernelBot dataset · June 9 Researcher Reciprocity License v1.0

Best evidence level for this revision: reported

JSON