compile_fail
3_SSM2D_FusedBidirectionalCrossScan.py
started
2026-05-11T09:46:40Z
compile error: Syntax error in custom generated code or ModelNew not found
timingkernelbench.timing
kernel mean / median / std
—
reference PyTorch source (the prompt the agent sees)
import torch
import torch.nn as nn
# Default ModelNew is PyTorch-only. run_and_check: use check_kernel=False with backend=cuda
# until the candidate file contains a real CUDA kernel (see kernel_static_checker).
class Model(nn.Module):
"""
Fused **2D bidirectional cross-scan** block for robotics **SSM-hybrid** backbones
(e.g. Vision Mamba / VMamba-style encoders on overhead maps, costmaps, or
multi-sensor tiles).
One forward fuses four fixed scan permutations of ``(H, W)`` into 1D sequences,
applies a **shared depthwise Conv1d** (lightweight stand-in for a selective SSM
along the scan), maps each result back to the grid, **sums** the four
directional contributions, and adds a **residual**.
Scan order (same spirit as 2D cross-scan / S2D):
LR — row-major left→right, rows top→bottom
RL — row-major right→left per row
TB — column-wise top→bottom, columns left→right
BT — flip vertical then same as TB
"""
def __init__(self, channels: int, kernel_size: int = 3):
super().__init__()
assert channels >= 1
assert kernel_size % 2 == 1, "use odd kernel_size for same-length conv1d"
self.channels = channels
pad = kernel_size // 2
self.dw_conv = nn.Conv1d(
channels,
channels,
kernel_size,
padding=pad,
groups=channels,
bias=True,
)
def _apply_scan_conv_unscan(self, x: torch.Tensor) -> torch.Tensor:
"""Shared depthwise conv along flattened 2D, then map back to (B,C,H,W)."""
B, C, H, W = x.shape
seq = x.flatten(2)
y = self.dw_conv(seq)
return y.view(B, C, H, W)
def forward(self, x: torch.Tensor) -> torch.Tensor:
"""
Args:
x: (B, C, H, W) spatial features (map / image / fused tile).
Returns:
(B, C, H, W) after fused four-way cross-scan + residual.
"""
B, C, H, W = x.shape
assert C == self.channels
acc = torch.zeros_like(x)
# LR: row-major
acc = acc + self._apply_scan_conv_unscan(x)
# RL: flip W, scan, unflip
xr = x.flip(dims=(-1,))
acc = acc + self._apply_scan_conv_unscan(xr).flip(dims=(-1,))
# TB: columns as sequences (permute to W,H, flatten)
xt = x.permute(0, 1, 3, 2).contiguous()
Bt, Ct, Wt, Ht = xt.shape
seq_t = xt.flatten(2)
yt = self.dw_conv(seq_t).view(Bt, Ct, Wt, Ht).permute(0, 1, 3, 2).contiguous()
acc = acc + yt
# BT: flip H then same TB path
xb = x.flip(dims=(-2,))
xbt = xb.permute(0, 1, 3, 2).contiguous()
Bb, Cb, Wb, Hb = xbt.shape
seq_b = xbt.flatten(2)
yb = self.dw_conv(seq_b).view(Bb, Cb, Wb, Hb).permute(0, 1, 3, 2).contiguous().flip(
dims=(-2,)
)
acc = acc + yb
return x + acc
class ModelNew(nn.Module):
"""KernelBench candidate entry point; default delegates to ``Model``."""
def __init__(self, channels: int, kernel_size: int = 3):
super().__init__()
self._impl = Model(channels, kernel_size=kernel_size)
def forward(self, x: torch.Tensor) -> torch.Tensor:
return self._impl(x)
batch_size = 4
channels = 64
height = 32
width = 32
kernel_size = 3
def get_inputs():
p = popcorn_pri
x = torch.randn(p.jitter_int(batch_size), channels, p.jitter_int(height), p.jitter_int(width))
return [x]
def get_init_inputs():
return [channels, kernel_size]
turn 0 · FINAL · 1 tool call(s) · 87.2s turn · 37.6s LLM · tokens in=1,827 · out=2,320 · total=4,741 · 9,992 chars (click to expand)
turn 087.2s turn · 37.6s LLM · 49.5s tools · tokens in=1,827 · out=2,320 · total=4,741 · FINAL
reasoning
The task is to replace the PyTorch model's forward computation with a custom CUDA kernel that's faster but numerically equivalent.
The task is to replace the PyTorch model's forward computation with a faster but numerically equivalent custom CUDA kernel. The model involves a fused 2D bidirectional cross-scan with a shared depthwise Conv1d and a residual addition. I need to implement a ModelNew class that uses torch.utils.cpp_extension.load_inline to compile and bind a CUDA kernel.
assistant```python
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline
cuda_source = r'''
#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>
template <typename scalar_t>
__global__ void cross_scan_conv_kernel(
const scalar_t* __restrict__ x,
const scalar_t* __restrict__ weight,
const scalar_t* __restrict__ bias,
scalar_t* __restrict__ out,
int B, int C, int H, int W, int K, int pad) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
int total = B * C * H * W;
if (idx >= total) return;
int w = idx % W;
int h = (idx / W) % H;
int c = (idx / (W * H)) % C;
int b = idx / (W * H * C);
scalar_t sum = 0;
// LR: horizontal, left to right
for (int k = 0; k < K; ++k) {
int ww = w - pad + k;
scalar_t val = 0;
if (ww >= 0 && ww < W) {
val = x[((b * C + c) * H + h) * W + ww];
}
sum += val * weight[c * K + k];
}
sum += bias[c];
// RL: flip W then LR then flip
for (int k = 0; k < K; ++k) {
int ww = (W - 1 - w) - pad + k;
scalar_t val = 0;
if (ww >= 0 && ww < W) {
val = x[((b * C + c) * H + h) * W + (W - 1 - ww)];
}
sum += val * weight[c * K + k];
}
sum += bias[c];
// TB: vertical top to bottom (treat columns as seq)
for (int k = 0; k < K; ++k) {
int hh = h - pad + k;
scalar_t val = 0;
if (hh >= 0 && hh < H) {
val = x[((b * C + c) * H + hh) * W + w];
}
sum += val * weight[c * K + k];
}
sum += bias[c];
// BT: flip H then TB then flip
for (int k = 0; k < K; ++k) {
int hh = (H - 1 - h) - pad + k;
scalar_t val = 0;
if (hh >= 0 && hh < H) {
val = x[((b * C + c) * H + (H - 1 - hh)) * W + w];
}
sum += val * weight[c * K + k];
}
sum += bias[c];
// residual
scalar_t res = x[idx];
out[idx] = res + sum;
}
torch::Tensor cross_scan_conv_cuda(
torch::Tensor x,
torch::Tensor weight,
torch::Tensor bias,
int kernel_size) {
auto B = x.size(0);
auto C = x.size(1);
auto H = x.size(2);
auto W = x.size(3);
int K = kernel_size;
int pad = K / 2;
auto out = torch::empty_like(x);
int total = B * C * H * W;
int threads = 256;
int blocks = (total + threads - 1) / threads;
AT_DISPATCH_FLOATING_TYPES(x.scalar_type(), "cross_scan_conv", ([&] {
cross_scan_conv_kernel<scalar_t><<<blocks, threads>>>(
x.data_ptr<scalar_t>(),
weight.data_ptr<scalar_t>(),
bias.data_ptr<scalar_t>(),
out.data_ptr<scalar_t>(),
B, C, H, W, K, pad);
}));
return out;
}
'''
cpp_source = r'''
torch::Tensor cross_scan_conv_cuda(
torch::Tensor x,
torch::Tensor weight,
torch::Tensor bias,
int kernel_size);
'''
module = load_inline(
name='cross_scan_conv',
cpp_sources=[cpp_source],
cuda_sources=[cuda_source],
functions=['cross_scan_conv_cuda'],
verbose=False
)
class ModelNew(nn.Module):
def __init__(self, channels: int, kernel_size: int = 3):
super().__init__()
assert channels >= 1
assert kernel_size % 2 == 1
self.channels = channels
self.kernel_size = kernel_size
pad = kernel_size // 2
self.dw_conv = nn.Conv1d(
channels,
channels,
kernel_size,
padding=pad,
groups=channels,
bias=True,
)
def forward(self, x: torch.Tensor) -> torch.Tensor:
weight = self.dw_conv.weight.view(self.channels, self.kernel_size).contiguous()
bias = self.dw_conv.bias
return module.cross_scan_conv_cuda(x, weight, bias, self.kernel_size)
```
The kernel implements the four directional depthwise convolutions fused with the residual addition in a single pass over the data. This avoids multiple tensor permutations, flips, and separate conv calls, improving efficiency on GPU. The implementation is numerically equivalent by direct replication of the scan-conv-unscan logic for each direction.
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 = r'''\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n\ntemplate <typename scalar_t>\n__global__ void cross_scan_conv_kernel(\n const scalar_t* __restrict__ x,\n const scalar_t* __restrict__ weight,\n const scalar_t* __restrict__ bias,\n scalar_t* __restrict__ out,\n int B, int C, int H, int W, int K, int pad) {\n\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * C * H * W;\n if (idx >= total) return;\n\n int w = idx % W;\n int h = (idx / W) % H;\n int c = (idx / (W * H)) % C;\n int b = idx / (W * H * C);\n\n scalar_t sum = 0;\n\n // LR: horizontal, left to right\n for (int k = 0; k < K; ++k) {\n int ww = w - pad + k;\n scalar_t val = 0;\n if (ww >= 0 && ww < W) {\n val = x[((b * C + c) * H + h) * W + ww];\n }\n sum += val * weight[c * K + k];\n }\n sum += bias[c];\n\n // RL: flip W then LR then flip\n for (int k = 0; k < K; ++k) {\n int ww = (W - 1 - w) - pad + k;\n scalar_t val = 0;\n if (ww >= 0 && ww < W) {\n val = x[((b * C + c) * H + h) * W + (W - 1 - ww)];\n }\n sum += val * weight[c * K + k];\n }\n sum += bias[c];\n\n // TB: vertical top to bottom (treat columns as seq)\n for (int k = 0; k < K; ++k) {\n int hh = h - pad + k;\n scalar_t val = 0;\n if (hh >= 0 && hh < H) {\n val = x[((b * C + c) * H + hh) * W + w];\n }\n sum += val * weight[c * K + k];\n }\n sum += bias[c];\n\n // BT: flip H then TB then flip\n for (int k = 0; k < K; ++k) {\n int hh = (H - 1 - h) - pad + k;\n scalar_t val = 0;\n if (hh >= 0 && hh < H) {\n val = x[((b * C + c) * H + (H - 1 - hh)) * W + w];\n }\n sum += val * weight[c * K + k];\n }\n sum += bias[c];\n\n // residual\n scalar_t res = x[idx];\n out[idx] = res + sum;\n}\n\ntorch::Tensor cross_scan_conv_cuda(\n torch::Tensor x,\n torch::Tensor weight,\n torch::Tensor bias,\n int kernel_size) {\n\n auto B = x.size(0);\n auto C = x.size(1);\n auto H = x.size(2);\n auto W = x.size(3);\n int K = kernel_size;\n int pad = K / 2;\n\n auto out = torch::empty_like(x);\n\n int total = B * C * H * W;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n\n AT_DISPATCH_FLOATING_TYPES(x.scalar_type(), \"cross_scan_conv\", ([&] {\n cross_scan_conv_kernel<scalar_t><<<blocks, threads>>>(\n x.data_ptr<scalar_t>(),\n weight.data_ptr<scalar_t>(),\n bias.data_ptr<scalar_t>(),\n out.data_ptr<scalar_t>(),\n B, C, H, W, K, pad);\n }));\n\n return out;\n}\n'''\n\ncpp_source = r'''\ntorch::Tensor cross_scan_conv_cuda(\n torch::Tensor x,\n torch::Tensor weight,\n torch::Tensor bias,\n int kernel_size);\n\n'''\n\nmodule = load_inline(\n name='cross_scan_conv',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['cross_scan_conv_cuda'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, channels: int, kernel_size: int = 3):\n super().__init__()\n assert channels >= 1\n assert kernel_size % 2 == 1\n self.channels = channels\n self.kernel_size = kernel_size\n pad = kernel_size // 2\n self.dw_conv = nn.Conv1d(\n channels,\n channels,\n kernel_size,\n padding=pad,\n groups=channels,\n bias=True,\n )\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n weight = self.dw_conv.weight.view(self.channels, self.kernel_size).contiguous()\n bias = self.dw_conv.bias\n return module.cross_scan_conv_cuda(x, weight, bias, self.kernel_size)\n```"
}
result ← submit_kernel (fail)
failsubmit_kernel
submit_kernel FAILED: kernel did not compile.
Syntax error in custom generated code or ModelNew not found
final submitted kernel (148 lines, 3,820 chars) — full source, untruncated
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline
cuda_source = r'''
#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>
template <typename scalar_t>
__global__ void cross_scan_conv_kernel(
const scalar_t* __restrict__ x,
const scalar_t* __restrict__ weight,
const scalar_t* __restrict__ bias,
scalar_t* __restrict__ out,
int B, int C, int H, int W, int K, int pad) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
int total = B * C * H * W;
if (idx >= total) return;
int w = idx % W;
int h = (idx / W) % H;
int c = (idx / (W * H)) % C;
int b = idx / (W * H * C);
scalar_t sum = 0;
// LR: horizontal, left to right
for (int k = 0; k < K; ++k) {
int ww = w - pad + k;
scalar_t val = 0;
if (ww >= 0 && ww < W) {
val = x[((b * C + c) * H + h) * W + ww];
}
sum += val * weight[c * K + k];
}
sum += bias[c];
// RL: flip W then LR then flip
for (int k = 0; k < K; ++k) {
int ww = (W - 1 - w) - pad + k;
scalar_t val = 0;
if (ww >= 0 && ww < W) {
val = x[((b * C + c) * H + h) * W + (W - 1 - ww)];
}
sum += val * weight[c * K + k];
}
sum += bias[c];
// TB: vertical top to bottom (treat columns as seq)
for (int k = 0; k < K; ++k) {
int hh = h - pad + k;
scalar_t val = 0;
if (hh >= 0 && hh < H) {
val = x[((b * C + c) * H + hh) * W + w];
}
sum += val * weight[c * K + k];
}
sum += bias[c];
// BT: flip H then TB then flip
for (int k = 0; k < K; ++k) {
int hh = (H - 1 - h) - pad + k;
scalar_t val = 0;
if (hh >= 0 && hh < H) {
val = x[((b * C + c) * H + (H - 1 - hh)) * W + w];
}
sum += val * weight[c * K + k];
}
sum += bias[c];
// residual
scalar_t res = x[idx];
out[idx] = res + sum;
}
torch::Tensor cross_scan_conv_cuda(
torch::Tensor x,
torch::Tensor weight,
torch::Tensor bias,
int kernel_size) {
auto B = x.size(0);
auto C = x.size(1);
auto H = x.size(2);
auto W = x.size(3);
int K = kernel_size;
int pad = K / 2;
auto out = torch::empty_like(x);
int total = B * C * H * W;
int threads = 256;
int blocks = (total + threads - 1) / threads;
AT_DISPATCH_FLOATING_TYPES(x.scalar_type(), "cross_scan_conv", ([&] {
cross_scan_conv_kernel<scalar_t><<<blocks, threads>>>(
x.data_ptr<scalar_t>(),
weight.data_ptr<scalar_t>(),
bias.data_ptr<scalar_t>(),
out.data_ptr<scalar_t>(),
B, C, H, W, K, pad);
}));
return out;
}
'''
cpp_source = r'''
torch::Tensor cross_scan_conv_cuda(
torch::Tensor x,
torch::Tensor weight,
torch::Tensor bias,
int kernel_size);
'''
module = load_inline(
name='cross_scan_conv',
cpp_sources=[cpp_source],
cuda_sources=[cuda_source],
functions=['cross_scan_conv_cuda'],
verbose=False
)
class ModelNew(nn.Module):
def __init__(self, channels: int, kernel_size: int = 3):
super().__init__()
assert channels >= 1
assert kernel_size % 2 == 1
self.channels = channels
self.kernel_size = kernel_size
pad = kernel_size // 2
self.dw_conv = nn.Conv1d(
channels,
channels,
kernel_size,
padding=pad,
groups=channels,
bias=True,
)
def forward(self, x: torch.Tensor) -> torch.Tensor:
weight = self.dw_conv.weight.view(self.channels, self.kernel_size).contiguous()
bias = self.dw_conv.bias
return module.cross_scan_conv_cuda(x, weight, bias, self.kernel_size)
```