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