submission 545123 · KernelIndex

Use it

Vendorable · source mirrored · license unknownView source →

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

grayscale_v2.py

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

interfacepython

Compatibility

measured onNVIDIA H100

declared hardwareNVIDIA H100

architecturessm_90

dtypesfp32

Benchmark evidence

1 measurement across 1 GPU, fastest first.

Operation / workload

Hardware

Latency

Rank

Observed

RGB to grayscalesuite of 6 cases

NVIDIA H100

1.37ms

#5 of 36

2026-03-13

Run →

Reported · How evidence levels are derived →

Source and license

sourceavailable

revision digestsha256:72e76f843574bc183c587f4a84350fd36a577aed7d3feae7b22f5a99e64089b7

license declaredunknown

license concludedunknown

authorsrajesh0042

imported2026-08-15

Kernel source

grayscale_v2.py35 lines

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

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

@triton.jit
def grayscale_kernel(
    data_ptr, out_ptr, n_pixels,
    BLOCK_SIZE: tl.constexpr,
):
    pid = tl.program_id(0)
    offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
    mask = offsets < n_pixels
    
    # data is (H, W, 3) = contiguous [R0, G0, B0, R1, G1, B1, ...]
    base = offsets * 3
    r = tl.load(data_ptr + base, mask=mask, other=0.0)
    g = tl.load(data_ptr + base + 1, mask=mask, other=0.0)
    b = tl.load(data_ptr + base + 2, mask=mask, other=0.0)
    
    gray = r * 0.2989 + g * 0.5870 + b * 0.1140
    tl.store(out_ptr + offsets, gray, mask=mask)

def custom_kernel(data: input_t) -> output_t:
    data, output = data
    h, w, _ = data.shape
    n_pixels = h * w
    BLOCK_SIZE = 2048
    grid = ((n_pixels + BLOCK_SIZE - 1) // BLOCK_SIZE,)
    grayscale_kernel[grid](data, output, n_pixels, BLOCK_SIZE=BLOCK_SIZE)
    return output

scrolls · 35 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 545074.

  ⋯ 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
   
- # Precompute weights
- _w = None
+ @triton.jit
+ def grayscale_kernel(
+     data_ptr, out_ptr, n_pixels,
+     BLOCK_SIZE: tl.constexpr,
+ ):
+     pid = tl.program_id(0)
+     offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
+     mask = offsets < n_pixels
+     
+     # data is (H, W, 3) = contiguous [R0, G0, B0, R1, G1, B1, ...]
+     base = offsets * 3
+     r = tl.load(data_ptr + base, mask=mask, other=0.0)
+     g = tl.load(data_ptr + base + 1, mask=mask, other=0.0)
+     b = tl.load(data_ptr + base + 2, mask=mask, other=0.0)
+     
+     gray = r * 0.2989 + g * 0.5870 + b * 0.1140
+     tl.store(out_ptr + offsets, gray, mask=mask)
   
  def custom_kernel(data: input_t) -> output_t:
-     global _w
      data, output = data
-     if _w is None or _w.device != data.device:
-         _w = torch.tensor([0.2989, 0.5870, 0.1140], device=data.device, dtype=data.dtype)
-     # Matrix multiply: reshape (H*W, 3) @ (3, 1) -> (H*W, 1) -> (H, W)
      h, w, _ = data.shape
-     output[...] = (data.reshape(-1, 3) @ _w.unsqueeze(1)).reshape(h, w)
+     n_pixels = h * w
+     BLOCK_SIZE = 2048
+     grid = ((n_pixels + BLOCK_SIZE - 1) // BLOCK_SIZE,)
+     grayscale_kernel[grid](data, output, n_pixels, BLOCK_SIZE=BLOCK_SIZE)
      return output

scrolls · 41 diff lines total

Best evidence level for this revision: reported

JSON