L1 P30 · popcorn

correct 30_broadcast_masked_from_row0.py

model

grok-4.3

turns / max

4 / 12

tool calls

tokens (total)

12391

wall clock

1146.5s

truncated

backend

cuda fp32

variant

popcorn

runtime

0.01 μs

ref runtime

0.01 μs

speedup

1.05x

started

2026-05-11T11:50:43Z

timingkernelbench.timing

kernel runtime

0.01 μs

ref runtime

0.01 μs

speedup

1.05x

kernel mean / median / std

0.01 / — / 1.70e-04 μs

ref mean / median / std

0.01 / — / 1.63e-04 μs

numerical precisioneval_kernel_against_ref

max abs error

0.00

mean abs error

0.00

max rel error

0.00

mean rel error

0.00

gpu memoryextended_metrics.measure_memory

kernel peak

33.01 MB

ref peak

33.01 MB

kernel / ref

1.00x

kernel launches / fusiontorch.profiler

kernel launches

ref launches

fusion ratio (ref / kernel)

3.67x

total cuda time

15 μs

SOL (speed-of-light)nsight_hardware_counters

SOL score (0–1)

0.001

dram util %

0.1%

compute util %

0.0%

bottleneck

unknown

dominant pipe

None

arithmetic intensity

—

ridge point

20.00

roofline / occupancynsight_hardware_counters

dram bandwidth

3.35 GB/s

dram util %

0.1%

fp32 tflops

—

fp32 util %

—

fp16 tflops

—

occupancy %

10.7%

bottleneck

unknown

dominant pipe %

—

arithmetic intensity

—

ridge point

20.00

L1 hit %

—

L2 hit %

99.1%

ld sectors / req

3.53

st sectors / req

—

regs / thread

smem / block

1024 B

block size

256

peak bw

3350.00 GB/s

peak fp32 tflops

67.00

peak fp16 tflops

1979.00

warp stalls (top 5)nsight

long_scoreboard

2815.8%

short_scoreboard

792.3%

reference PyTorch source (the prompt the agent sees)

"""
Simulated broadcast of row 0 into masked columns for all rows: where mask[j] is True,
every row matches x[0, j]; elsewhere keep x[b, j].
"""

import torch
import torch.nn as nn

from kernelbench.distributed_collectives import default_device

B = 18
D = 44

# Fixed mask (not all True, not all False)
_MASK = torch.tensor(
    [1, 0, 1, 1, 0, 0, 1, 0, 1, 1, 0, 1, 1, 0, 1, 0, 1, 1, 0, 1, 1, 0, 1, 0, 1, 1, 1, 0, 1, 0, 1, 1, 0, 1, 1, 0, 1, 0, 1, 1, 0, 1, 0, 1],
    dtype=torch.bool,
)


class Model(nn.Module):
    def __init__(self):
        super().__init__()
        self.register_buffer("row_mask", _MASK.clone(), persistent=False)

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        m = self.row_mask.to(device=x.device).view(1, D).expand(B, D)
        row0 = x[0:1].expand(B, D)
        return torch.where(m, row0, x)


def get_inputs():
    dev = default_device()
    return [torch.randn(popcorn_pri.jitter_int(B), D, device=dev)]


def get_init_inputs():
    return []

turn 0583.7s turn · 2.1s LLM · 581.6s tools · tokens in=1,606 · out=6 · total=1,721

reasoning

The task is to replace the forward computation with a custom CUDA kernel. The operation is:

call → get_gpu_specs

callget_gpu_specs {}

result ← get_gpu_specs (ok)

okget_gpu_specs GPU specs for NVIDIA H100 80GB HBM3: total memory (runtime): 79.1 GB spec entry: H100 GPU Architecture: Hopper GPU Memory: 80GB Memory Bandwidth: 3.35 TB/s FP64 TFLOPS: 34 FP64 Tensor Core TFLOPS: 67 FP32 TFLOPS: 67 TF32 Tensor Core TFLOPS: 989 with sparsity BFLOAT16 Tensore Core TFLOPS: 1979 with sparsity FP16 Tensor Core TFLOPS: 1979 with sparsity FP8 Tensor Core TFLOPS: 3958 with sparsity INT8 Tensor Core TOPS: 3958 with sparsity Register File Size: 64K 32-bit registers per SM Maximum number of registers per thread: 255 Maximum number of thread blocks per SM: 32 Shared memory capacity per SM: 228 KB Maximum shared memory per thread block: 227 KB

turn 1104.0s turn · 9.3s LLM · 94.7s tools · tokens in=1,954 · out=537 · total=2,965

reasoning

It's an H100, very fast memory bandwidth 3.35 TB/s. The data is tiny: 18*44*4 bytes ~ 3KB, so this will be launch overhead dominated probably, but we need to beat PyTorch. I need to replace the forward computation of the PyTorch model with a custom CUDA kernel. The `ModelNew` class must accept the same arguments as `Model`, implement `forward()` with the same signature, and produce the same output numerically.

call → compile_kernel

callcompile_kernel { "kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = \"\"\"\n#include <torch/extension.h>\n\n__global__ void broadcast_mask_kernel(const float* __restrict__ x, const bool* __restrict__ mask, float* __restrict__ out, int B, int D) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * D;\n if (idx < total) {\n int j = idx % D;\n if (mask[j]) {\n out[idx] = x[j];\n } else {\n out[idx] = x[idx];\n }\n }\n}\n\ntorch::Tensor broadcast_mask_cuda(torch::Tensor x, torch::Tensor mask) {\n auto out = torch::empty_like(x);\n int B = x.size(0);\n int D = x.size(1);\n int total = B * D;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n broadcast_mask_kernel<<<blocks, threads>>>(x.data_ptr<float>(), mask.data_ptr<bool>(), out.data_ptr<float>(), B, D);\n return out;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor broadcast_mask_cuda(torch::Tensor x, torch::Tensor mask);\n\"\"\"\n\nmodule = load_inline(\n name='broadcast_mask',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['broadcast_mask_cuda'],\n verbose=True,\n extra_cflags=['-O3']\n)\n\n_MASK = torch.tensor(\n [1, 0, 1, 1, 0, 0, 1, 0, 1, 1, 0, 1, 1, 0, 1, 0, 1, 1, 0, 1, 1, 0, 1, 0, 1, 1, 1, 0, 1, 0, 1, 1, 0, 1, 1, 0, 1, 0, 1, 1, 0, 1, 0, 1],\n dtype=torch.bool,\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.register_buffer(\"row_mask\", _MASK.clone(), persistent=False)\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n m = self.row_mask.to(device=x.device)\n return module.broadcast_mask_cuda(x, m)\n" }

result ← compile_kernel (ok)

okcompile_kernel compile_kernel PASSED: kernel compiled without errors.

turn 2226.1s turn · 5.2s LLM · 221.0s tools · tokens in=2,982 · out=538 · total=3,527

reasoning

Compiled successfully. Now check correctness.

call → run_correctness

callrun_correctness { "kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = \"\"\"\n#include <torch/extension.h>\n\n__global__ void broadcast_mask_kernel(const float* __restrict__ x, const bool* __restrict__ mask, float* __restrict__ out, int B, int D) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * D;\n if (idx < total) {\n int j = idx % D;\n if (mask[j]) {\n out[idx] = x[j];\n } else {\n out[idx] = x[idx];\n }\n }\n}\n\ntorch::Tensor broadcast_mask_cuda(torch::Tensor x, torch::Tensor mask) {\n auto out = torch::empty_like(x);\n int B = x.size(0);\n int D = x.size(1);\n int total = B * D;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n broadcast_mask_kernel<<<blocks, threads>>>(x.data_ptr<float>(), mask.data_ptr<bool>(), out.data_ptr<float>(), B, D);\n return out;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor broadcast_mask_cuda(torch::Tensor x, torch::Tensor mask);\n\"\"\"\n\nmodule = load_inline(\n name='broadcast_mask',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['broadcast_mask_cuda'],\n verbose=True,\n extra_cflags=['-O3']\n)\n\n_MASK = torch.tensor(\n [1, 0, 1, 1, 0, 0, 1, 0, 1, 1, 0, 1, 1, 0, 1, 0, 1, 1, 0, 1, 1, 0, 1, 0, 1, 1, 1, 0, 1, 0, 1, 1, 0, 1, 1, 0, 1, 0, 1, 1, 0, 1, 0, 1],\n dtype=torch.bool,\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.register_buffer(\"row_mask\", _MASK.clone(), persistent=False)\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n m = self.row_mask.to(device=x.device)\n return module.broadcast_mask_cuda(x, m)\n" }

result ← run_correctness (ok)

okrun_correctness run_correctness PASSED: (5 / 5) trials all matched the reference. Numerical precision: max_abs_err=0.00e+00 mean_abs_err=0.00e+00 max_rel_err=0.00e+00

turn 3232.6s turn · 7.4s LLM · 225.2s tools · tokens in=3,587 · out=537 · total=4,178 · FINAL

reasoning

Correct. Now, to check if faster, I should submit, but since tiny, maybe optimize a bit, but it should be fine.

call → submit_kernel

callsubmit_kernel { "kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = \"\"\"\n#include <torch/extension.h>\n\n__global__ void broadcast_mask_kernel(const float* __restrict__ x, const bool* __restrict__ mask, float* __restrict__ out, int B, int D) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * D;\n if (idx < total) {\n int j = idx % D;\n if (mask[j]) {\n out[idx] = x[j];\n } else {\n out[idx] = x[idx];\n }\n }\n}\n\ntorch::Tensor broadcast_mask_cuda(torch::Tensor x, torch::Tensor mask) {\n auto out = torch::empty_like(x);\n int B = x.size(0);\n int D = x.size(1);\n int total = B * D;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n broadcast_mask_kernel<<<blocks, threads>>>(x.data_ptr<float>(), mask.data_ptr<bool>(), out.data_ptr<float>(), B, D);\n return out;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor broadcast_mask_cuda(torch::Tensor x, torch::Tensor mask);\n\"\"\"\n\nmodule = load_inline(\n name='broadcast_mask',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['broadcast_mask_cuda'],\n verbose=True,\n extra_cflags=['-O3']\n)\n\n_MASK = torch.tensor(\n [1, 0, 1, 1, 0, 0, 1, 0, 1, 1, 0, 1, 1, 0, 1, 0, 1, 1, 0, 1, 1, 0, 1, 0, 1, 1, 1, 0, 1, 0, 1, 1, 0, 1, 1, 0, 1, 0, 1, 1, 0, 1, 0, 1],\n dtype=torch.bool,\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.register_buffer(\"row_mask\", _MASK.clone(), persistent=False)\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n m = self.row_mask.to(device=x.device)\n return module.broadcast_mask_cuda(x, m)\n" }

result ← submit_kernel (ok)

oksubmit_kernel submit_kernel PASSED: (5 / 5) correctness trials all passed. Kernel runtime: 0.01 μs Runtime stats: mean=0.01μs median=?μs std=0.00μs Numerical precision: max_abs_err=0.00e+00 mean_abs_err=0.00e+00 max_rel_err=0.00e+00 Memory: 33.01 MB peak (ref: 33.01 MB, ratio: 1.0x) Kernel launches: 3 (ref: 11, fusion ratio: 3.6667) SOL score: 0.001

final submitted kernel (59 lines, 1,690 chars) — full source, untruncated

import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline

cuda_source = """
#include <torch/extension.h>

__global__ void broadcast_mask_kernel(const float* __restrict__ x, const bool* __restrict__ mask, float* __restrict__ out, int B, int D) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int total = B * D;
    if (idx < total) {
        int j = idx % D;
        if (mask[j]) {
            out[idx] = x[j];
        } else {
            out[idx] = x[idx];
        }
    }
}

torch::Tensor broadcast_mask_cuda(torch::Tensor x, torch::Tensor mask) {
    auto out = torch::empty_like(x);
    int B = x.size(0);
    int D = x.size(1);
    int total = B * D;
    int threads = 256;
    int blocks = (total + threads - 1) / threads;
    broadcast_mask_kernel<<<blocks, threads>>>(x.data_ptr<float>(), mask.data_ptr<bool>(), out.data_ptr<float>(), B, D);
    return out;
}
"""

cpp_source = """
torch::Tensor broadcast_mask_cuda(torch::Tensor x, torch::Tensor mask);
"""

module = load_inline(
    name='broadcast_mask',
    cpp_sources=[cpp_source],
    cuda_sources=[cuda_source],
    functions=['broadcast_mask_cuda'],
    verbose=True,
    extra_cflags=['-O3']
)

_MASK = torch.tensor(
    [1, 0, 1, 1, 0, 0, 1, 0, 1, 1, 0, 1, 1, 0, 1, 0, 1, 1, 0, 1, 1, 0, 1, 0, 1, 1, 1, 0, 1, 0, 1, 1, 0, 1, 1, 0, 1, 0, 1, 1, 0, 1, 0, 1],
    dtype=torch.bool,
)

class ModelNew(nn.Module):
    def __init__(self):
        super().__init__()
        self.register_buffer("row_mask", _MASK.clone(), persistent=False)

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        m = self.row_mask.to(device=x.device)
        return module.broadcast_mask_cuda(x, m)