Skip to content
KernelIndex
Search⌘K
Sign in

submission 682179

ngolhn · python · License unknown
Kernel source · 96 lines ↓

Use it

Vendorable · source mirrored · license unknownView source →

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

submission.py
curl "https://kernelindex.com/api/v1/implementations/kernelbot-vectoradd-v2-682179?include=source"
interfacepython
Compatibility
measured onNVIDIA B200
declared hardwareNVIDIA B200
architecturessm_100
dtypesfp16

Benchmark evidence

1 measurement across 1 GPU, fastest first.

Operation / workload
Hardware
Latency
Rank
Observed
FP16 vector additionsuite of 5 cases
NVIDIA B200
233.5µs
#11 of 66
2026-03-31
Run →

Reported · How evidence levels are derived →

Source and license

sourceavailable
revision digestsha256:5c4fadb0291b0d6284bb025619387a1590ca22611f3288830cdfeefd16d78255
license declaredunknown
license concludedunknown
authorsngolhn
imported2026-08-15

Techniques

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

vector-width = float4void vectoradd_kernel(const float4* __restrict__ A, const float4* __restrict__ B,

Kernel source

submission.py96 lines
#!POPCORN leaderboard vectoradd_v2
#!POPCORN gpu B200

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

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

__global__ __launch_bounds__(512, 4)
void vectoradd_kernel(const float4* __restrict__ A, const float4* __restrict__ B,
                      float4* __restrict__ C, const int n4) {
    const int idx = blockIdx.x * 512 + threadIdx.x;
    if (idx < n4) {
        float4 a = A[idx];
        float4 b = B[idx];

        half2* ah = reinterpret_cast<half2*>(&a);
        half2* bh = reinterpret_cast<half2*>(&b);
        float4 c;
        half2* ch = reinterpret_cast<half2*>(&c);

        ch[0] = __hadd2(ah[0], bh[0]);
        ch[1] = __hadd2(ah[1], bh[1]);
        ch[2] = __hadd2(ah[2], bh[2]);
        ch[3] = __hadd2(ah[3], bh[3]);

        C[idx] = c;
    }
}

__global__ void vectoradd_tail(const __half* __restrict__ A, const __half* __restrict__ B,
                               __half* __restrict__ C, const int start, const int n) {
    const int idx = start + blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        C[idx] = __hadd(A[idx], B[idx]);
    }
}

void vectoradd_raw(int64_t a_ptr, int64_t b_ptr, int64_t c_ptr, int N) {
    const int n8 = N / 8;
    const int remainder = N - n8 * 8;

    if (n8 > 0) {
        const int blocks = (n8 + 511) / 512;
        vectoradd_kernel<<<blocks, 512>>>(
            reinterpret_cast<const float4*>(a_ptr),
            reinterpret_cast<const float4*>(b_ptr),
            reinterpret_cast<float4*>(c_ptr),
            n8);
    }

    if (remainder > 0) {
        const int start = n8 * 8;
        const int blocks = (remainder + 511) / 512;
        vectoradd_tail<<<blocks, 512>>>(
            reinterpret_cast<const __half*>(a_ptr),
            reinterpret_cast<const __half*>(b_ptr),
            reinterpret_cast<__half*>(c_ptr),
            start, N);
    }
}
"""

cpp_src = r"""
void vectoradd_raw(int64_t a_ptr, int64_t b_ptr, int64_t c_ptr, int N);
"""

_ext = load_inline(
    name="vectoradd_w512",
    cpp_sources=cpp_src,
    cuda_sources=cuda_src,
    functions=["vectoradd_raw"],
    with_cuda=True,
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3", "--use_fast_math", "-arch=sm_100a", "-maxrregcount=32"],
    verbose=False,
)

# Small warmup
_wa = torch.randn(128, 128, device="cuda", dtype=torch.float16)
_wb = torch.randn(128, 128, device="cuda", dtype=torch.float16)
_wc = torch.empty(128, 128, device="cuda", dtype=torch.float16)
_ext.vectoradd_raw(_wa.data_ptr(), _wb.data_ptr(), _wc.data_ptr(), _wa.numel())
torch.cuda.synchronize()
del _wa, _wb, _wc


def custom_kernel(data: input_t) -> output_t:
    A, B, output = data
    _ext.vectoradd_raw(A.data_ptr(), B.data_ptr(), output.data_ptr(), A.numel())
    return output
scrolls · 96 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 682119.

  ⋯ 9 unchanged lines
  #include <cuda_runtime.h>
  #include <cuda_fp16.h>
   
- __global__ __launch_bounds__(256, 16)
+ __global__ __launch_bounds__(512, 4)
  void vectoradd_kernel(const float4* __restrict__ A, const float4* __restrict__ B,
                        float4* __restrict__ C, const int n4) {
-     const int idx = blockIdx.x * 256 + threadIdx.x;
-     if (idx >= n4) return;
+     const int idx = blockIdx.x * 512 + threadIdx.x;
+     if (idx < n4) {
+         float4 a = A[idx];
+         float4 b = B[idx];
   
-     if (idx + 256 < n4) {
-         const void* pa = reinterpret_cast<const void*>(&A[idx + 256]);
-         const void* pb = reinterpret_cast<const void*>(&B[idx + 256]);
-         asm volatile("prefetch.global.L2 [%0];" :: "l"(pa));
-         asm volatile("prefetch.global.L2 [%0];" :: "l"(pb));
-     }
+         half2* ah = reinterpret_cast<half2*>(&a);
+         half2* bh = reinterpret_cast<half2*>(&b);
+         float4 c;
+         half2* ch = reinterpret_cast<half2*>(&c);
   
-     float4 a, b;
-     const float4* addr_a = &A[idx];
-     const float4* addr_b = &B[idx];
+         ch[0] = __hadd2(ah[0], bh[0]);
+         ch[1] = __hadd2(ah[1], bh[1]);
+         ch[2] = __hadd2(ah[2], bh[2]);
+         ch[3] = __hadd2(ah[3], bh[3]);
   
-     asm volatile("ld.global.v4.b32 {%0, %1, %2, %3}, [%4];"
-                  : "=r"(reinterpret_cast<unsigned int*>(&a)[0]),
-                    "=r"(reinterpret_cast<unsigned int*>(&a)[1]),
-                    "=r"(reinterpret_cast<unsigned int*>(&a)[2]),
-                    "=r"(reinterpret_cast<unsigned int*>(&a)[3])
-                  : "l"(addr_a));
-  
-     asm volatile("ld.global.v4.b32 {%0, %1, %2, %3}, [%4];"
-                  : "=r"(reinterpret_cast<unsigned int*>(&b)[0]),
-                    "=r"(reinterpret_cast<unsigned int*>(&b)[1]),
-                    "=r"(reinterpret_cast<unsigned int*>(&b)[2]),
-                    "=r"(reinterpret_cast<unsigned int*>(&b)[3])
-                  : "l"(addr_b));
-  
-     half2* ah = reinterpret_cast<half2*>(&a);
-     half2* bh = reinterpret_cast<half2*>(&b);
-     float4 c;
-     half2* ch = reinterpret_cast<half2*>(&c);
-     ch[0] = __hadd2(ah[0], bh[0]);
-     ch[1] = __hadd2(ah[1], bh[1]);
-     ch[2] = __hadd2(ah[2], bh[2]);
-     ch[3] = __hadd2(ah[3], bh[3]);
-  
-     float4* addr_c = &C[idx];
-     asm volatile("st.global.v4.b32 [%0], {%1, %2, %3, %4};"
-                  :: "l"(addr_c),
-                     "r"(reinterpret_cast<unsigned int*>(&c)[0]),
-                     "r"(reinterpret_cast<unsigned int*>(&c)[1]),
-                     "r"(reinterpret_cast<unsigned int*>(&c)[2]),
-                     "r"(reinterpret_cast<unsigned int*>(&c)[3]));
+         C[idx] = c;
+     }
  }
   
  __global__ void vectoradd_tail(const __half* __restrict__ A, const __half* __restrict__ B,
  ⋯ 9 unchanged lines
      const int remainder = N - n8 * 8;
   
      if (n8 > 0) {
-         const int blocks = (n8 + 255) / 256;
-         vectoradd_kernel<<<blocks, 256>>>(
+         const int blocks = (n8 + 511) / 512;
+         vectoradd_kernel<<<blocks, 512>>>(
              reinterpret_cast<const float4*>(a_ptr),
              reinterpret_cast<const float4*>(b_ptr),
              reinterpret_cast<float4*>(c_ptr),
  ⋯ 2 unchanged lines
   
      if (remainder > 0) {
          const int start = n8 * 8;
-         const int blocks = (remainder + 255) / 256;
-         vectoradd_tail<<<blocks, 256>>>(
+         const int blocks = (remainder + 511) / 512;
+         vectoradd_tail<<<blocks, 512>>>(
              reinterpret_cast<const __half*>(a_ptr),
              reinterpret_cast<const __half*>(b_ptr),
              reinterpret_cast<__half*>(c_ptr),
  ⋯ 7 unchanged lines
  """
   
  _ext = load_inline(
-     name="vectoradd_warmup",
+     name="vectoradd_w512",
      cpp_sources=cpp_src,
      cuda_sources=cuda_src,
      functions=["vectoradd_raw"],
  ⋯ 3 unchanged lines
      verbose=False,
  )
   
- # Warmup: run kernel once at import time to ensure CUDA context is fully initialized,
- # kernel is loaded, and any lazy initialization is done before timing starts
- _warmup_a = torch.randn(128, 128, device="cuda", dtype=torch.float16)
- _warmup_b = torch.randn(128, 128, device="cuda", dtype=torch.float16)
- _warmup_c = torch.empty(128, 128, device="cuda", dtype=torch.float16)
- _ext.vectoradd_raw(_warmup_a.data_ptr(), _warmup_b.data_ptr(), _warmup_c.data_ptr(), _warmup_a.numel())
+ # Small warmup
+ _wa = torch.randn(128, 128, device="cuda", dtype=torch.float16)
+ _wb = torch.randn(128, 128, device="cuda", dtype=torch.float16)
+ _wc = torch.empty(128, 128, device="cuda", dtype=torch.float16)
+ _ext.vectoradd_raw(_wa.data_ptr(), _wb.data_ptr(), _wc.data_ptr(), _wa.numel())
  torch.cuda.synchronize()
- del _warmup_a, _warmup_b, _warmup_c
+ del _wa, _wb, _wc
   
   
  def custom_kernel(data: input_t) -> output_t:
scrolls · 121 diff lines total

Best evidence level for this revision: reported

JSON