Skip to content
KernelIndex
Search⌘K
Sign in

submission 545221

rajesh0042 · python · License unknown
Kernel source · 37 lines ↓

Use it

Vendorable · source mirrored · license unknownView source →

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

grayscale_v3.py
curl "https://kernelindex.com/api/v1/implementations/kernelbot-grayscale-v2-545221?include=source"
interfacepython
Compatibility
measured onNVIDIA A100
declared hardwareNVIDIA A100
architecturessm_80
dtypesfp32

Benchmark evidence

1 measurement across 1 GPU, fastest first.

Operation / workload
Hardware
Latency
Rank
Observed
RGB to grayscalesuite of 6 cases
NVIDIA A100
2.51ms
#21 of 137
2026-03-13
Run →

Reported · How evidence levels are derived →

Source and license

sourceavailable
revision digestsha256:4b03f1d55e28e7984e8e8d52453d961aa2eb740c828459879f0ce6ca1586ca62
license declaredunknown
license concludedunknown
authorsrajesh0042
imported2026-08-15

Kernel source

grayscale_v3.py37 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

# Tighter grayscale with larger block and vectorized access pattern
@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)

    # Use exact FP32 constants matching reference
    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 = 4096  # Larger block for better occupancy
    grid = ((n_pixels + BLOCK_SIZE - 1) // BLOCK_SIZE,)
    grayscale_kernel[grid](data, output, n_pixels, BLOCK_SIZE=BLOCK_SIZE)
    return output
scrolls · 37 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 545196.

  ⋯ 5 unchanged lines
  import triton.language as tl
  from task import input_t, output_t
   
+ # Tighter grayscale with larger block and vectorized access pattern
  @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 = 2048
+     BLOCK_SIZE = 4096  # Larger block for better occupancy
      grid = ((n_pixels + BLOCK_SIZE - 1) // BLOCK_SIZE,)
      grayscale_kernel[grid](data, output, n_pixels, BLOCK_SIZE=BLOCK_SIZE)
      return output
scrolls · 34 diff lines total

Best evidence level for this revision: reported

JSON