submission 763185 · KernelIndex

Use it

Vendorable · source mirrored · license unknownView source →

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

submission_cuda_inline_A100.py

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

interfacepython

Compatibility

measured onNVIDIA A100

declared hardwareNVIDIA A100

architecturessm_80

dtypesfp32

Benchmark evidence

1 measurement across 1 GPU, fastest first.

Operation / workload

Hardware

Latency

Rank

Observed

Vector sum reductionsuite of 6 cases

NVIDIA A100

142.1µs

#19 of 96

2026-04-11

Run →

Reported · How evidence levels are derived →

Source and license

sourceavailable

revision digestsha256:fd394dacb7fab8b2ca007070397eba145a501dc31c8e29115bdb91d6018e297b

license declaredunknown

license concludedunknown

authorsCaptnJackSparrow

imported2026-08-15

Techniques

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

shared-memory__shared__ double warp_sums[16];

vector-width = float4const float4* input4 = reinterpret_cast<const float4*>(input);

Kernel source

submission_cuda_inline_A100.py127 lines

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

sum_cuda_source = """
#include <cuda_fp16.h>
#include <cuda_runtime.h>

#define MAX_BLOCKS 2048

static double* d_partial = nullptr;
__device__ unsigned int g_retirement_count = 0;

__global__ void __launch_bounds__(512)
sum_kernel(const float* __restrict__ input,
           double* __restrict__ partial_sums,
           float* __restrict__ output,
           int N) {
    double s0 = 0.0, s1 = 0.0, s2 = 0.0, s3 = 0.0;

    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = blockDim.x * gridDim.x;

    int N4 = N / 4;
    const float4* input4 = reinterpret_cast<const float4*>(input);

    int vec_idx = idx;
    for (; vec_idx + 3 * stride < N4; vec_idx += 4 * stride) {
        float4 v0 = __ldg(&input4[vec_idx]);
        float4 v1 = __ldg(&input4[vec_idx + stride]);
        float4 v2 = __ldg(&input4[vec_idx + 2 * stride]);
        float4 v3 = __ldg(&input4[vec_idx + 3 * stride]);
        s0 += (double)v0.x + (double)v0.y + (double)v0.z + (double)v0.w;
        s1 += (double)v1.x + (double)v1.y + (double)v1.z + (double)v1.w;
        s2 += (double)v2.x + (double)v2.y + (double)v2.z + (double)v2.w;
        s3 += (double)v3.x + (double)v3.y + (double)v3.z + (double)v3.w;
    }
    for (; vec_idx < N4; vec_idx += stride) {
        float4 v = __ldg(&input4[vec_idx]);
        s0 += (double)v.x + (double)v.y + (double)v.z + (double)v.w;
    }

    double thread_sum = s0 + s1 + s2 + s3;

    int scalar_start = N4 * 4;
    for (int i = scalar_start + threadIdx.x; i < N; i += blockDim.x) {
        thread_sum += (double)__ldg(&input[i]);
    }

    unsigned mask = 0xffffffff;
    for (int offset = 16; offset > 0; offset >>= 1)
        thread_sum += __shfl_down_sync(mask, thread_sum, offset);

    __shared__ double warp_sums[16];
    int lane = threadIdx.x & 31;
    int warp_id = threadIdx.x >> 5;

    if (lane == 0) warp_sums[warp_id] = thread_sum;
    __syncthreads();

    double block_sum = 0.0;
    if (warp_id == 0) {
        thread_sum = (lane < (blockDim.x >> 5)) ? warp_sums[lane] : 0.0;
        for (int offset = 16; offset > 0; offset >>= 1)
            thread_sum += __shfl_down_sync(mask, thread_sum, offset);
        block_sum = thread_sum;
    }

    __shared__ bool s_is_last;
    if (threadIdx.x == 0) {
        partial_sums[blockIdx.x] = block_sum;
        __threadfence();
        unsigned int ticket = atomicAdd(&g_retirement_count, 1);
        s_is_last = (ticket == gridDim.x - 1);
    }
    __syncthreads();

    if (s_is_last) {
        double final_sum = 0.0;
        for (int i = threadIdx.x; i < gridDim.x; i += blockDim.x)
            final_sum += partial_sums[i];
        for (int offset = 16; offset > 0; offset >>= 1)
            final_sum += __shfl_down_sync(mask, final_sum, offset);
        if (lane == 0) warp_sums[warp_id] = final_sum;
        __syncthreads();
        if (warp_id == 0) {
            final_sum = (lane < (blockDim.x >> 5)) ? warp_sums[lane] : 0.0;
            for (int offset = 16; offset > 0; offset >>= 1)
                final_sum += __shfl_down_sync(mask, final_sum, offset);
            if (lane == 0) {
                output[0] = (float)final_sum;
                g_retirement_count = 0;
            }
        }
    }
}

void sum_cuda(torch::Tensor input, torch::Tensor output) {
    int N = input.numel();
    const int threads = 512;
    int blocks = min((N / 4 + threads - 1) / threads, MAX_BLOCKS);
    if (blocks < 1) blocks = 1;
    if (!d_partial) cudaMalloc(&d_partial, MAX_BLOCKS * sizeof(double));
    sum_kernel<<<blocks, threads>>>(
        input.data_ptr<float>(), d_partial, output.data_ptr<float>(), N);
}
"""

sum_cpp_source = """
#include <torch/extension.h>
void sum_cuda(torch::Tensor input, torch::Tensor output);
"""

sum_module = load_inline(
    name='sum_cuda',
    cpp_sources=sum_cpp_source,
    cuda_sources=sum_cuda_source,
    functions=['sum_cuda'],
    verbose=True,
    extra_cuda_cflags=['-O3', '--use_fast_math', '-gencode', 'arch=compute_80,code=sm_80'],
)

def custom_kernel(data: input_t) -> output_t:
    data, output = data
    sum_module.sum_cuda(data, output)
    return output[0]

scrolls · 127 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 763184.

  ⋯ 15 unchanged lines
             double* __restrict__ partial_sums,
             float* __restrict__ output,
             int N) {
-     double thread_sum = 0.0;
+     double s0 = 0.0, s1 = 0.0, s2 = 0.0, s3 = 0.0;
   
      int idx = blockIdx.x * blockDim.x + threadIdx.x;
      int stride = blockDim.x * gridDim.x;
  ⋯ 2 unchanged lines
      const float4* input4 = reinterpret_cast<const float4*>(input);
   
      int vec_idx = idx;
-     for (; vec_idx + stride < N4; vec_idx += 2 * stride) {
+     for (; vec_idx + 3 * stride < N4; vec_idx += 4 * stride) {
          float4 v0 = __ldg(&input4[vec_idx]);
          float4 v1 = __ldg(&input4[vec_idx + stride]);
-         thread_sum += (double)v0.x + (double)v0.y + (double)v0.z + (double)v0.w;
-         thread_sum += (double)v1.x + (double)v1.y + (double)v1.z + (double)v1.w;
+         float4 v2 = __ldg(&input4[vec_idx + 2 * stride]);
+         float4 v3 = __ldg(&input4[vec_idx + 3 * stride]);
+         s0 += (double)v0.x + (double)v0.y + (double)v0.z + (double)v0.w;
+         s1 += (double)v1.x + (double)v1.y + (double)v1.z + (double)v1.w;
+         s2 += (double)v2.x + (double)v2.y + (double)v2.z + (double)v2.w;
+         s3 += (double)v3.x + (double)v3.y + (double)v3.z + (double)v3.w;
      }
      for (; vec_idx < N4; vec_idx += stride) {
          float4 v = __ldg(&input4[vec_idx]);
-         thread_sum += (double)v.x + (double)v.y + (double)v.z + (double)v.w;
+         s0 += (double)v.x + (double)v.y + (double)v.z + (double)v.w;
      }
   
+     double thread_sum = s0 + s1 + s2 + s3;
+  
      int scalar_start = N4 * 4;
      for (int i = scalar_start + threadIdx.x; i < N; i += blockDim.x) {
          thread_sum += (double)__ldg(&input[i]);
      }
   
      unsigned mask = 0xffffffff;
-     for (int offset = 16; offset > 0; offset >>= 1) {
+     for (int offset = 16; offset > 0; offset >>= 1)
          thread_sum += __shfl_down_sync(mask, thread_sum, offset);
-     }
   
      __shared__ double warp_sums[16];
      int lane = threadIdx.x & 31;
      int warp_id = threadIdx.x >> 5;
   
-     if (lane == 0) {
-         warp_sums[warp_id] = thread_sum;
-     }
+     if (lane == 0) warp_sums[warp_id] = thread_sum;
      __syncthreads();
   
      double block_sum = 0.0;
      if (warp_id == 0) {
-         int num_warps = blockDim.x >> 5;
-         thread_sum = (lane < num_warps) ? warp_sums[lane] : 0.0;
-         for (int offset = 16; offset > 0; offset >>= 1) {
+         thread_sum = (lane < (blockDim.x >> 5)) ? warp_sums[lane] : 0.0;
+         for (int offset = 16; offset > 0; offset >>= 1)
              thread_sum += __shfl_down_sync(mask, thread_sum, offset);
-         }
          block_sum = thread_sum;
      }
   
  ⋯ 8 unchanged lines
   
      if (s_is_last) {
          double final_sum = 0.0;
-         for (int i = threadIdx.x; i < gridDim.x; i += blockDim.x) {
+         for (int i = threadIdx.x; i < gridDim.x; i += blockDim.x)
              final_sum += partial_sums[i];
-         }
-  
-         for (int offset = 16; offset > 0; offset >>= 1) {
+         for (int offset = 16; offset > 0; offset >>= 1)
              final_sum += __shfl_down_sync(mask, final_sum, offset);
-         }
-  
-         if (lane == 0) {
-             warp_sums[warp_id] = final_sum;
-         }
+         if (lane == 0) warp_sums[warp_id] = final_sum;
          __syncthreads();
-  
          if (warp_id == 0) {
-             int num_warps = blockDim.x >> 5;
-             final_sum = (lane < num_warps) ? warp_sums[lane] : 0.0;
-             for (int offset = 16; offset > 0; offset >>= 1) {
+             final_sum = (lane < (blockDim.x >> 5)) ? warp_sums[lane] : 0.0;
+             for (int offset = 16; offset > 0; offset >>= 1)
                  final_sum += __shfl_down_sync(mask, final_sum, offset);
-             }
              if (lane == 0) {
                  output[0] = (float)final_sum;
                  g_retirement_count = 0;
  ⋯ 7 unchanged lines
      const int threads = 512;
      int blocks = min((N / 4 + threads - 1) / threads, MAX_BLOCKS);
      if (blocks < 1) blocks = 1;
-  
-     if (!d_partial) {
-         cudaMalloc(&d_partial, MAX_BLOCKS * sizeof(double));
-     }
-  
+     if (!d_partial) cudaMalloc(&d_partial, MAX_BLOCKS * sizeof(double));
      sum_kernel<<<blocks, threads>>>(
-         input.data_ptr<float>(),
-         d_partial,
-         output.data_ptr<float>(),
-         N
-     );
+         input.data_ptr<float>(), d_partial, output.data_ptr<float>(), N);
  }
  """

scrolls · 118 diff lines total

Best evidence level for this revision: reported

JSON