submission 545356 · KernelIndex

Use it

Vendorable · source mirrored · license unknownView source →

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

grayscale_v8.py

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

interfacepython

Compatibility

measured onNVIDIA L4

declared hardwareNVIDIA L4

architecturessm_89

dtypesfp32

Benchmark evidence

1 measurement across 1 GPU, fastest first.

Operation / workload

Hardware

Latency

Rank

Observed

RGB to grayscalesuite of 6 cases

NVIDIA L4

17.2ms

#5 of 12

2026-03-13

Run →

Reported · How evidence levels are derived →

Source and license

sourceavailable

revision digestsha256:207edf2f6c5df5c1659519eb0b8bfb789c9a124bbfaae3ece3ddafce569d408b

license declaredunknown

license concludedunknown

authorsrajesh0042

imported2026-08-15

Techniques

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

num-warps = 8grayscale_kernel[grid](data, output, n_pixels, BLOCK_SIZE=BLOCK_SIZE, num_warps=8)

Kernel source

grayscale_v8.py40 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

# Grayscale A100/H100: Triton with tuned block sizes and num_warps
# The coalescing issue: data is (H, W, 3) = stride-3 pattern
# Try restructuring: load 3*BLOCK elements, deinterleave

@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
    
    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
    # 1024 block was original, 2048 was better on some GPUs
    # Let's try 1024 with explicit num_warps
    BLOCK_SIZE = 1024
    grid = ((n_pixels + BLOCK_SIZE - 1) // BLOCK_SIZE,)
    grayscale_kernel[grid](data, output, n_pixels, BLOCK_SIZE=BLOCK_SIZE, num_warps=8)
    return output

scrolls · 40 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 545230.

  ⋯ 5 unchanged lines
  import triton.language as tl
  from task import input_t, output_t
   
- # Tighter grayscale with larger block and vectorized access pattern
+ # Grayscale A100/H100: Triton with tuned block sizes and num_warps
+ # The coalescing issue: data is (H, W, 3) = stride-3 pattern
+ # Try restructuring: load 3*BLOCK elements, deinterleave
+  
  @triton.jit
  def grayscale_kernel(
      data_ptr, out_ptr, n_pixels,
  ⋯ 2 unchanged lines
      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)
-  
-     # Use exact FP32 constants matching reference
+     
      gray = r * 0.2989 + g * 0.5870 + b * 0.1140
      tl.store(out_ptr + offsets, gray, mask=mask)
   
  ⋯ 1 unchanged lines
      data, output = data
      h, w, _ = data.shape
      n_pixels = h * w
-     BLOCK_SIZE = 4096  # Larger block for better occupancy
+     # 1024 block was original, 2048 was better on some GPUs
+     # Let's try 1024 with explicit num_warps
+     BLOCK_SIZE = 1024
      grid = ((n_pixels + BLOCK_SIZE - 1) // BLOCK_SIZE,)
-     grayscale_kernel[grid](data, output, n_pixels, BLOCK_SIZE=BLOCK_SIZE)
+     grayscale_kernel[grid](data, output, n_pixels, BLOCK_SIZE=BLOCK_SIZE, num_warps=8)
      return output

scrolls · 41 diff lines total

Best evidence level for this revision: reported

JSON