Skip to content
KernelIndex
Search⌘K

submission 512893

JordanNanos · python · License unknown

Use it

Vendorable · source mirrored · license unknownView source →

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

vectoradd_py_submission.py
curl "https://kernelindex.com/api/v1/implementations/kernelbot-vectoradd-v2-512893?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
240.6µs
#46 of 66
2026-03-05

Reported · How evidence levels are derived →

Source and license

sourceavailable
revision digestsha256:429b62eae8e1e4a65a54059668d3ed8d7491311c2103a2f499591afe9584ae20
license declaredunknown
license concludedunknown
authorsJordanNanos
imported2026-08-15

Techniques

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

autotune@triton.autotune(
num-warps = 4triton.Config({'BLOCK_SIZE': 1024}, num_warps=4, num_stages=2),
stages = 2triton.Config({'BLOCK_SIZE': 1024}, num_warps=4, num_stages=2),

Kernel source

vectoradd_py_submission.py58 lines
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,
    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, 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")


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
scrolls · 58 lines total

Source code from GPU Mode and the KernelBot dataset · June 9 Researcher Reciprocity License v1.0

Best evidence level for this revision: reported

JSON