submission 515759 · KernelIndex

Use it

Vendorable · source mirrored · license unknownView source →

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

submission.py

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

interfacepython

Compatibility

measured onNVIDIA A100

declared hardwareNVIDIA A100

architecturessm_80

dtypesfp16

Benchmark evidence

1 measurement across 1 GPU, fastest first.

Operation / workload

Hardware

Latency

Rank

Observed

FP16 vector additionsuite of 5 cases

NVIDIA A100

894.6µs

#6= of 87

2026-03-07

Run →

Reported · How evidence levels are derived →

Source and license

sourceavailable

revision digestsha256:54030dafb7c3662ab64224452f0a9f0afc0e4c53827b4f7f5b25725d0b586298

license declaredunknown

license concludedunknown

authorsweyj4

imported2026-08-15

Kernel source

submission.py27 lines

#!POPCORN leaderboard vectoradd_v2
#!POPCORN gpu A100

import torch
import triton
import triton.language as tl

@triton.jit
def vec_add_kernel(A_ptr, B_ptr, C_ptr, SIZE, BLOCK: tl.constexpr):
    pid = tl.program_id(0)
    offs = pid * BLOCK + tl.arange(0, BLOCK)
    mask = offs < SIZE
    a = tl.load(A_ptr + offs, mask=mask)
    b = tl.load(B_ptr + offs, mask=mask)
    tl.store(C_ptr + offs, a + b, mask=mask)

def custom_kernel(data):
    A, B, C = data
    SIZE = A.numel()
    BLOCK = 1024
    grid = (triton.cdiv(SIZE, BLOCK),)
    vec_add_kernel[grid](
        A.view(-1), B.view(-1), C.view(-1),
        SIZE, BLOCK
    )
    return C

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

  #!POPCORN leaderboard vectoradd_v2
  #!POPCORN gpu A100
- import subprocess
- import sys
- subprocess.check_call([sys.executable, "-m", "pip", "install", "numba"])
   
- import numba
- from numba import cuda
- import math
  import torch
+ import triton
+ import triton.language as tl
   
- @cuda.jit
- def vec_add_fp16_kernel(A, B, C, SIZE):
-     idx = cuda.threadIdx.x + cuda.blockIdx.x * cuda.blockDim.x
-     if idx < SIZE:
-         C[idx] = A[idx] + B[idx]
+ @triton.jit
+ def vec_add_kernel(A_ptr, B_ptr, C_ptr, SIZE, BLOCK: tl.constexpr):
+     pid = tl.program_id(0)
+     offs = pid * BLOCK + tl.arange(0, BLOCK)
+     mask = offs < SIZE
+     a = tl.load(A_ptr + offs, mask=mask)
+     b = tl.load(B_ptr + offs, mask=mask)
+     tl.store(C_ptr + offs, a + b, mask=mask)
   
  def custom_kernel(data):
      A, B, C = data
      SIZE = A.numel()
-     A_flat = cuda.as_cuda_array(A.view(-1))
-     B_flat = cuda.as_cuda_array(B.view(-1))
-     C_flat = cuda.as_cuda_array(C.view(-1))
-     threads_per_block = 1024
-     blocks = math.ceil(SIZE / threads_per_block)
-     vec_add_fp16_kernel[blocks, threads_per_block](A_flat, B_flat, C_flat, SIZE)
+     BLOCK = 1024
+     grid = (triton.cdiv(SIZE, BLOCK),)
+     vec_add_kernel[grid](
+         A.view(-1), B.view(-1), C.view(-1),
+         SIZE, BLOCK
+     )
      return C

scrolls · 43 diff lines total

Best evidence level for this revision: reported

JSON