submission 67541

#!POPCORN leaderboard grayscale_v2

from task import input_t, output_t
from utils import DeterministicContext
import triton
import triton.language as tl


@triton.jit
def _grayscale_kernel(
    rgb_ptr,
    output_ptr,
    height: tl.int32,
    width: tl.int32,
    stride_h: tl.int32,
    stride_w: tl.int32,
    stride_c: tl.int32,
    out_stride_h: tl.int32,
    out_stride_w: tl.int32,
    BLOCK_SIZE: tl.constexpr,
):
    pid = tl.program_id(0)

    offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
    n_pixels = height * width
    mask = offsets < n_pixels

    rows = offsets // width
    cols = offsets % width

    rgb_index = rows * stride_h + cols * stride_w

    r = tl.load(rgb_ptr + rgb_index + 0 * stride_c, mask=mask, other=0.0)
    g = tl.load(rgb_ptr + rgb_index + 1 * stride_c, mask=mask, other=0.0)
    b = tl.load(rgb_ptr + rgb_index + 2 * stride_c, mask=mask, other=0.0)

    grayscale = 0.2989 * r + 0.5870 * g + 0.1140 * b

    out_index = rows * out_stride_h + cols * out_stride_w
    tl.store(output_ptr + out_index, grayscale, mask=mask)


def custom_kernel(data: input_t) -> output_t:
    with DeterministicContext():
        rgb, output = data
        height, width, channels = rgb.shape
        if channels != 3:
            raise ValueError(f"Expected last dimension to be 3, got {channels}")

        BLOCK_SIZE = 1024

        grid = (triton.cdiv(height * width, BLOCK_SIZE),)

        _grayscale_kernel[grid](
            rgb,
            output,
            height,
            width,
            rgb.stride(0),
            rgb.stride(1),
            rgb.stride(2),
            output.stride(0),
            output.stride(1),
            BLOCK_SIZE=BLOCK_SIZE,
            num_warps=4,
            num_stages=1,
        )
        return output