Skip to content
KernelIndex
Search⌘K
Sign in

submission 513370

JordanNanos · python · License unknown
Kernel source · 28 lines ↓

Use it

Vendorable · source mirrored · license unknownView source →

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

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

Benchmark evidence

1 measurement across 1 GPU, fastest first.

Operation / workload
Hardware
Latency
Rank
Observed
FP16 vector additionsuite of 5 cases
NVIDIA H100
524.7µs
#14 of 44
2026-03-06
Run →

Reported · How evidence levels are derived →

Source and license

sourceavailable
revision digestsha256:10782a754c396750a262bb458ec3f19bc456ad65139c056ef21678b5a96687eb
license declaredunknown
license concludedunknown
authorsJordanNanos
imported2026-08-15

Kernel source

submission.py28 lines
import torch
import triton
import triton.language as tl
from task import input_t, output_t


@triton.jit
def vector_add_kernel(
    a_ptr,
    b_ptr,
    output_ptr,
    n_elements,
    BLOCK_SIZE: tl.constexpr,
):
    pid = tl.program_id(axis=0)
    block_start = pid * BLOCK_SIZE
    offsets = block_start + 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)
    output = a + b
    tl.store(output_ptr + offsets, output, mask=mask)


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

+ import torch
  import triton
  import triton.language as tl
- import torch
  from task import input_t, output_t
   
   
- @triton.autotune(
-     configs=[
-         triton.Config({'BLOCK_SIZE': 1024}, num_warps=4, num_stages=2),
-         triton.Config({'BLOCK_SIZE': 1024}, num_warps=8, num_stages=2),
-         triton.Config({'BLOCK_SIZE': 2048}, num_warps=8, num_stages=2),
-         triton.Config({'BLOCK_SIZE': 2048}, num_warps=16, num_stages=2),
-         triton.Config({'BLOCK_SIZE': 4096}, num_warps=16, num_stages=2),
-         triton.Config({'BLOCK_SIZE': 4096}, num_warps=32, num_stages=2),
-         triton.Config({'BLOCK_SIZE': 8192}, num_warps=32, num_stages=2),
-         triton.Config({'BLOCK_SIZE': 8192}, num_warps=16, num_stages=2),
-         triton.Config({'BLOCK_SIZE': 16384}, num_warps=32, num_stages=2),
-         triton.Config({'BLOCK_SIZE': 32768}, num_warps=32, num_stages=2),
-         triton.Config({'BLOCK_SIZE': 1024}, num_warps=4, num_stages=4),
-         triton.Config({'BLOCK_SIZE': 2048}, num_warps=8, num_stages=4),
-         triton.Config({'BLOCK_SIZE': 4096}, num_warps=16, num_stages=4),
-         triton.Config({'BLOCK_SIZE': 4096}, num_warps=32, num_stages=4),
-         triton.Config({'BLOCK_SIZE': 8192}, num_warps=32, num_stages=4),
-         triton.Config({'BLOCK_SIZE': 16384}, num_warps=32, num_stages=4),
-         triton.Config({'BLOCK_SIZE': 32768}, num_warps=32, num_stages=4),
-         triton.Config({'BLOCK_SIZE': 65536}, num_warps=32, num_stages=2),
-         triton.Config({'BLOCK_SIZE': 65536}, num_warps=32, num_stages=4),
-     ],
-     key=['n_elements'],
- )
  @triton.jit
  def vector_add_kernel(
      a_ptr,
      b_ptr,
-     c_ptr,
+     output_ptr,
      n_elements,
      BLOCK_SIZE: tl.constexpr,
  ):
  ⋯ 1 unchanged lines
      block_start = pid * BLOCK_SIZE
      offsets = block_start + tl.arange(0, BLOCK_SIZE)
      mask = offsets < n_elements
-     a = tl.load(a_ptr + offsets, mask=mask, cache_modifier=".cg")
-     b = tl.load(b_ptr + offsets, mask=mask, cache_modifier=".cg")
-     c = a + b
-     tl.store(c_ptr + offsets, c, mask=mask, cache_modifier=".cg")
+     a = tl.load(a_ptr + offsets, mask=mask)
+     b = tl.load(b_ptr + offsets, mask=mask)
+     output = a + b
+     tl.store(output_ptr + offsets, output, mask=mask)
   
   
  def custom_kernel(data: input_t) -> output_t:
      A, B, output = data
-     n_elements = A.numel()
-     grid = lambda meta: (triton.cdiv(n_elements, meta['BLOCK_SIZE']),)
-     vector_add_kernel[grid](
-         A, B, output,
-         n_elements,
-     )
-     return output
+     torch.add(A, B, out=output)
+     return output
   No newline at end of file
scrolls · 66 diff lines total

Best evidence level for this revision: reported

JSON