submission 545314 · KernelIndex

Use it

Vendorable · source mirrored · license unknownView source →

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

vectoradd_v5.py

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

236.7µs

#29 of 66

2026-03-13

Run →

Reported · How evidence levels are derived →

Source and license

sourceavailable

revision digestsha256:28ae65704275b609a5f0f35610a99d461f31c292ff5fe930162e69034dc9e62d

license declaredunknown

license concludedunknown

authorsrajesh0042

imported2026-08-15

Kernel source

vectoradd_v5.py26 lines

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

import torch
from task import input_t, output_t

# VectorAdd with warmup to amortize kernel launch overhead
# The gap is only 3µs -- every microsecond counts

# Warm up CUDA context and kernel caches
def _warmup():
    for size in [1024, 4096, 16384, 65536, 262144, 1048576]:
        a = torch.randn(size, device='cuda', dtype=torch.float16)
        b = torch.randn(size, device='cuda', dtype=torch.float16)
        c = torch.empty(size, device='cuda', dtype=torch.float16)
        torch.add(a, b, out=c)
        torch.add(a, b, out=c)
    torch.cuda.synchronize()

_warmup()

def custom_kernel(data: input_t) -> output_t:
    A, B, output = data
    torch.add(A, B, out=output)
    return output

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 545222.

  ⋯ 1 unchanged lines
  os.environ["CUBLAS_WORKSPACE_CONFIG"] = ":4096:8"
   
  import torch
- import triton
- import triton.language as tl
  from task import input_t, output_t
   
- # Vectoradd with Triton for potential better scheduling
- @triton.jit
- def vectoradd_kernel(
-     a_ptr, b_ptr, out_ptr, n_elements,
-     BLOCK_SIZE: tl.constexpr,
- ):
-     pid = tl.program_id(0)
-     offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
-     mask = offsets < n_elements
-     a = tl.load(a_ptr + offsets, mask=mask)
-     b = tl.load(b_ptr + offsets, mask=mask)
-     tl.store(out_ptr + offsets, a + b, mask=mask)
+ # VectorAdd with warmup to amortize kernel launch overhead
+ # The gap is only 3µs -- every microsecond counts
   
+ # Warm up CUDA context and kernel caches
+ def _warmup():
+     for size in [1024, 4096, 16384, 65536, 262144, 1048576]:
+         a = torch.randn(size, device='cuda', dtype=torch.float16)
+         b = torch.randn(size, device='cuda', dtype=torch.float16)
+         c = torch.empty(size, device='cuda', dtype=torch.float16)
+         torch.add(a, b, out=c)
+         torch.add(a, b, out=c)
+     torch.cuda.synchronize()
+  
+ _warmup()
+  
  def custom_kernel(data: input_t) -> output_t:
      A, B, output = data
-     n = A.numel()
-     BLOCK_SIZE = 1024
-     grid = ((n + BLOCK_SIZE - 1) // BLOCK_SIZE,)
-     vectoradd_kernel[grid](A, B, output, n, BLOCK_SIZE=BLOCK_SIZE)
+     torch.add(A, B, out=output)
      return output

scrolls · 43 diff lines total

Best evidence level for this revision: reported

JSON