Kimi-K2.6 · original · l12_default_gpt
error 93_masked_cumsum.py
model
Kimi-K2.6
turns / max
2 / 10
tool calls
1
tokens (total)
0
wall clock
—
truncated
no
backend
cuda fp32
variant
original
runtime
—
ref runtime
0.00 μs
speedup
—
started
2026-05-02T20:35:32Z
reference PyTorch source (the prompt the agent sees)
import torch
import torch.nn as nn

class Model(nn.Module):
    """
    A model that performs a masked cumulative sum, only summing elements that satisfy a condition.

    Parameters:
        dim (int): The dimension along which to perform the masked cumulative sum.
    """

    def __init__(self, dim):
        super(Model, self).__init__()
        self.dim = dim

    def forward(self, x, mask):
        """
        Args:
            x (torch.Tensor): Input tensor of shape (batch_size, *input_shape).
            mask (torch.Tensor): Boolean mask of the same shape as x.

        Returns:
            torch.Tensor: Cumulative sum of elements where mask is True.
        """
        return torch.cumsum(x * mask, dim=self.dim)

batch_size = 32768
input_shape = (32768,)
dim = 1

def get_inputs():
    x = torch.rand(batch_size, *input_shape)
    mask = torch.randint(0, 2, x.shape).bool()  # Random boolean mask
    return [x, mask]

def get_init_inputs():
    return [dim]
turn 06.0s turn · 6.0s LLM · 0.0s tools
reasoning
We need to replace the PyTorch reference `Model` with a custom CUDA kernel `ModelNew`. The reference computation: `torch.cumsum(x * mask, dim=self.dim)`, where `x` is fp32 (I think by default `torch.rand` gives fp32), mask is bool. Input shape: `[batch_size, input_shape] = [32768, 32768]`, so a 2D tensor of shape [32768, 32768]. `dim=1`. So we need to compute, for each row independently, the cumulative sum along dimension 1, but only including elements where mask is True (i.e., multiply x by mask, then cumsum). The reference does `x * mask`, where mask is bool, so mask gets promoted to same dtype as x, i.e., fp32, values 0.0 or 1.0, then cumsum along dim=1. So the core computation: for each i in [0, batch_size), for each j in [0, N), compute sum_{k=0..j} x[i,k] * mask[i,k]. We need a custom CUDA kernel that is faster than PyTorch's cumsum. Since the shape is huge: 32768x32768 = about 1.07 billion elements, ~4 GB of fp32 data. Cumsum along rows of a large 2D matrix. PyTorch's cumsum likely uses some parallel scan algorithm, but for such a large tensor it may be memory bandwidth bound. However, we also need to incorporate the mask multiplication. Doing that as a separate kernel (or fused) can save bandwidth. We need to write a CUDA kernel using torch.utils.cpp_extension.load_inline. Also need to consider if we can use cub::DeviceScan (CUB) for scan per row. Since we have many rows (32768), we can process each row with CUB block scan. The number of elements per row is also 32768. A 2D grid where each block handles a row (or multiple rows) using CUB's BlockScan. We need to get high memory bandwidth and occupancy. We can load elements from x and mask. Since mask is bool (1 byte each), we should load it as a vector or at least contiguous. Given row-major layout (dim=1 is contiguous? Actually a [32768, 32768] tensor in PyTorch default memory layout is row-major? PyTorch uses strided layout, but for a 2D tensor created by `torch.rand(batch_size, *input_shape)`, the strides are [32768, 1] (row-major). So dim=1 is the contiguous dimension. That means each row is contiguous in memory. This is good. So for each row i, elements x[i, j] are contiguous. We can assign one thread block per row, or multiple blocks per row. 32768 rows; 32768 elements per row. This is huge. Let's consider hardware. Let's get GPU specs first. We need to call get_gpu_specs. Then design kernel. We must write a complete Python module that defines ModelNew. At module load time, compile inline CUDA. Then instantiate ModelNew with same args as Model. We need to be careful: The evaluation may run on a specific GPU; we need to use `get_gpu_specs` to know its capabilities. Let's get GPU specs.
call → get_gpu_specs
callget_gpu_specs {}
result ← get_gpu_specs (fail)
failget_gpu_specs get_gpu_specs FAILED: unexpected error during tool execution. _WorkItemTimeout: work item exceeded soft timeout budget Traceback (most recent call last): File "/scratch/tejas/PopcornBench/src/kernelbench/agent/agent.py", line 627, in _execute_tool return tool.execute(self.ctx, **args) File "/scratch/tejas/PopcornBench/src/kernelbench/agent/tools.py", line 656, in execute return ctx.eval_client.get_gpu_specs(ctx) File "/scratch/tejas/PopcornBench/src/kernelbench/agent/eval_client.py", line 62, in get_gpu_specs return self._call("specs", args) File "/scratch/tejas/PopcornBench/src/kernelbench/agent/eval_client.py", line 117, in _call payload = self._response_q.get(timeout=remaining) File "<string>", line 2, in get File "/home/prabhune/.local/share/uv/python/cpython-3.10-linux-x86_64-gnu/lib/python3.10/multiprocessing/managers.py", line 818, in _callmethod kind, result = conn.recv() File "/home/prabhune/.local/share/uv/python/cpython-3.10-linux-x86_64-gnu/lib/python3.10/multiprocessing/connection.py", line 250, in recv buf = self._recv_bytes() File "/home/prabhune/.local/share/uv/python/cpython-3.10-linux-x86_64-gnu/lib/python3.10/multiprocessing/connection.py", line 414, in _recv_bytes buf = self._recv(4) File "/home/prabhune/.local/share/uv/python/cpython-3.10-linux-x86_64-gnu/lib/python3.10/multiprocessing/connection.py", line 379, in _recv chunk = read(handle, remaining) File "/scratch/tejas/PopcornBench/scripts/run_sweep.py", line 161, in _alarm_handler raise _WorkItemTimeout("work item exceeded soft timeout budget") __mp_main__._WorkItemTimeout: work item exceeded soft timeout budget
turn 16.0s turn · 6.0s LLM · 0.0s tools
feedbackLLM call failed: Empty: