correct
3_SSM2D_FusedBidirectionalCrossScan.py
started
2026-05-11T21:03:35Z
timingkernelbench.timing
kernel mean / median / std
0.03 / — / 0.03 μs
ref mean / median / std
0.17 / — / 0.09 μs
numerical precisioneval_kernel_against_ref
gpu memoryextended_metrics.measure_memory
kernel launches / fusiontorch.profiler
fusion ratio (ref / kernel)
22.80x
SOL (speed-of-light)nsight_hardware_counters
roofline / occupancynsight_hardware_counters
warp stalls (top 5)nsight
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 0215.5s turn · 26.0s LLM · 189.5s tools · tokens in=2,725 · out=1,145 · total=4,870
reasoning
The task is to replace the forward computation of this PyTorch model with a custom CUDA kernel. The model is a fused 2D bidirectional cross-scan block.
I need to replace the forward computation of the PyTorch model's forward method with a custom CUDA kernel. The model is a fused 2D bidirectional cross-scan block, performing four scans: left-to-right row-major, right-to-left (with flip), top-to-bottom on permuted columns, and bottom-to-top (with 翻). Each scan applies a depthwise Conv1d with kernel_size=3 and padding=1, then maps back and sums them up, adding a residual connection.
To implement this in ModelNew, it must accept the same inputs as the original Model, implement forward with the same signature, and produce the same output numerically. The forward involves preparing tensors for each direction, applying Conv1d, permuting back, summing, and adding residual.
For the CUDA part, I'll use torch.utils.cpp_extension.load_inline to compile the CUDA code inline.
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#include <cuda.h>\n#include <cuda_runtime.h>\n\n__global__ void cross_scan_conv_kernel(\n const float* __restrict__ x,\n const float* __restrict__ weight,\n const float* __restrict__ bias,\n float* __restrict__ out,\n int B, int C, int H, int W, int K\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 float sum = 0.0f;\n\n // LR: row-major flatten\n int pos_lr = h * W + w;\n int L = H * W;\n for (int k = 0; k < K; ++k) {\n int p = pos_lr + (k - K/2);\n if (p >= 0 && p < L) {\n int hh = p / W;\n int ww = p % W;\n sum += x[((b * C + c) * H + hh) * W + ww] * weight[c * K + k];\n }\n }\n sum += bias[c];\n\n // RL: flip W\n int pos_rl = h * W + (W - 1 - w);\n for (int k = 0; k < K; ++k) {\n int p = pos_rl + (k - K/2);\n if (p >= 0 && p < L) {\n int hh = p / W;\n int ww = W - 1 - (p % W);\n sum += x[((b * C + c) * H + hh) * W + ww] * weight[c * K + k];\n }\n }\n sum += bias[c];\n\n // TB: permute to (B,C,W,H), flatten -> vertical scans\n int pos_tb = w * H + h;\n int Lt = W * H;\n for (int k = 0; k < K; ++k) {\n int p = pos_tb + (k - K/2);\n if (p >= 0 && p < Lt) {\n int ww = p / H;\n int hh = p % H;\n sum += x[((b * C + c) * H + hh) * W + ww] * weight[c * K + k];\n }\n }\n sum += bias[c];\n\n // BT: flip H then TB\n int pos_bt = w * H + (H - 1 - h);\n for (int k = 0; k < K; ++k) {\n int p = pos_bt + (k - K/2);\n if (p >= 0 && p < Lt) {\n int ww = p / H;\n int hh = H - 1 - (p % H);\n sum += x[((b * C + c) * H + hh) * W + ww] * weight[c * K + k];\n }\n }\n sum += bias[c];\n\n // residual\n float xval = x[idx];\n out[idx] = xval + 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 int B = x.size(0);\n int C = x.size(1);\n int H = x.size(2);\n int W = x.size(3);\n auto out = torch::zeros_like(x);\n\n int total = B * C * H * W;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n\n cross_scan_conv_kernel<<<blocks, threads>>>(\n x.data_ptr<float>(),\n weight.data_ptr<float>(),\n bias.data_ptr<float>(),\n out.data_ptr<float>(),\n B, C, H, W, kernel_size\n );\n\n return out;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor cross_scan_conv_cuda(torch::Tensor x, torch::Tensor weight, torch::Tensor bias, int kernel_size);\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=True,\n extra_cuda_cflags=['-O3']\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, channels: int, kernel_size: int = 3):\n super().__init__()\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 self.module = module\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n weight = self.dw_conv.weight\n bias = self.dw_conv.bias\n return self.module.cross_scan_conv_cuda(x, weight, bias, self.kernel_size)\n"
}
result ← compile_kernel (ok)
okcompile_kernel
compile_kernel PASSED: kernel compiled without errors.
turn 111.3s turn · 9.6s LLM · 1.7s tools · tokens in=4,887 · out=1,146 · total=6,040
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#include <cuda.h>\n#include <cuda_runtime.h>\n\n__global__ void cross_scan_conv_kernel(\n const float* __restrict__ x,\n const float* __restrict__ weight,\n const float* __restrict__ bias,\n float* __restrict__ out,\n int B, int C, int H, int W, int K\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 float sum = 0.0f;\n\n // LR: row-major flatten\n int pos_lr = h * W + w;\n int L = H * W;\n for (int k = 0; k < K; ++k) {\n int p = pos_lr + (k - K/2);\n if (p >= 0 && p < L) {\n int hh = p / W;\n int ww = p % W;\n sum += x[((b * C + c) * H + hh) * W + ww] * weight[c * K + k];\n }\n }\n sum += bias[c];\n\n // RL: flip W\n int pos_rl = h * W + (W - 1 - w);\n for (int k = 0; k < K; ++k) {\n int p = pos_rl + (k - K/2);\n if (p >= 0 && p < L) {\n int hh = p / W;\n int ww = W - 1 - (p % W);\n sum += x[((b * C + c) * H + hh) * W + ww] * weight[c * K + k];\n }\n }\n sum += bias[c];\n\n // TB: permute to (B,C,W,H), flatten -> vertical scans\n int pos_tb = w * H + h;\n int Lt = W * H;\n for (int k = 0; k < K; ++k) {\n int p = pos_tb + (k - K/2);\n if (p >= 0 && p < Lt) {\n int ww = p / H;\n int hh = p % H;\n sum += x[((b * C + c) * H + hh) * W + ww] * weight[c * K + k];\n }\n }\n sum += bias[c];\n\n // BT: flip H then TB\n int pos_bt = w * H + (H - 1 - h);\n for (int k = 0; k < K; ++k) {\n int p = pos_bt + (k - K/2);\n if (p >= 0 && p < Lt) {\n int ww = p / H;\n int hh = H - 1 - (p % H);\n sum += x[((b * C + c) * H + hh) * W + ww] * weight[c * K + k];\n }\n }\n sum += bias[c];\n\n // residual\n float xval = x[idx];\n out[idx] = xval + 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 int B = x.size(0);\n int C = x.size(1);\n int H = x.size(2);\n int W = x.size(3);\n auto out = torch::zeros_like(x);\n\n int total = B * C * H * W;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n\n cross_scan_conv_kernel<<<blocks, threads>>>(\n x.data_ptr<float>(),\n weight.data_ptr<float>(),\n bias.data_ptr<float>(),\n out.data_ptr<float>(),\n B, C, H, W, kernel_size\n );\n\n return out;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor cross_scan_conv_cuda(torch::Tensor x, torch::Tensor weight, torch::Tensor bias, int kernel_size);\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=True,\n extra_cuda_cflags=['-O3']\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, channels: int, kernel_size: int = 3):\n super().__init__()\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 self.module = module\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n weight = self.dw_conv.weight\n bias = self.dw_conv.bias\n return self.module.cross_scan_conv_cuda(x, weight, bias, self.kernel_size)\n"
}
result ← run_correctness (ok)
okrun_correctness
run_correctness PASSED: (3 / 3) trials all matched the reference.
Numerical precision: max_abs_err=1.91e-06 mean_abs_err=9.86e-08 max_rel_err=3.70e-02
turn 251.8s turn · 9.2s LLM · 42.6s tools · tokens in=6,100 · out=1,145 · total=7,287 · FINAL
reasoning
Correct! Now, to check if faster, but since we need to submit, and it should be faster as it's a custom kernel avoiding all the permutes/flattens etc.
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#include <cuda.h>\n#include <cuda_runtime.h>\n\n__global__ void cross_scan_conv_kernel(\n const float* __restrict__ x,\n const float* __restrict__ weight,\n const float* __restrict__ bias,\n float* __restrict__ out,\n int B, int C, int H, int W, int K\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 float sum = 0.0f;\n\n // LR: row-major flatten\n int pos_lr = h * W + w;\n int L = H * W;\n for (int k = 0; k < K; ++k) {\n int p = pos_lr + (k - K/2);\n if (p >= 0 && p < L) {\n int hh = p / W;\n int ww = p % W;\n sum += x[((b * C + c) * H + hh) * W + ww] * weight[c * K + k];\n }\n }\n sum += bias[c];\n\n // RL: flip W\n int pos_rl = h * W + (W - 1 - w);\n for (int k = 0; k < K; ++k) {\n int p = pos_rl + (k - K/2);\n if (p >= 0 && p < L) {\n int hh = p / W;\n int ww = W - 1 - (p % W);\n sum += x[((b * C + c) * H + hh) * W + ww] * weight[c * K + k];\n }\n }\n sum += bias[c];\n\n // TB: permute to (B,C,W,H), flatten -> vertical scans\n int pos_tb = w * H + h;\n int Lt = W * H;\n for (int k = 0; k < K; ++k) {\n int p = pos_tb + (k - K/2);\n if (p >= 0 && p < Lt) {\n int ww = p / H;\n int hh = p % H;\n sum += x[((b * C + c) * H + hh) * W + ww] * weight[c * K + k];\n }\n }\n sum += bias[c];\n\n // BT: flip H then TB\n int pos_bt = w * H + (H - 1 - h);\n for (int k = 0; k < K; ++k) {\n int p = pos_bt + (k - K/2);\n if (p >= 0 && p < Lt) {\n int ww = p / H;\n int hh = H - 1 - (p % H);\n sum += x[((b * C + c) * H + hh) * W + ww] * weight[c * K + k];\n }\n }\n sum += bias[c];\n\n // residual\n float xval = x[idx];\n out[idx] = xval + 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 int B = x.size(0);\n int C = x.size(1);\n int H = x.size(2);\n int W = x.size(3);\n auto out = torch::zeros_like(x);\n\n int total = B * C * H * W;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n\n cross_scan_conv_kernel<<<blocks, threads>>>(\n x.data_ptr<float>(),\n weight.data_ptr<float>(),\n bias.data_ptr<float>(),\n out.data_ptr<float>(),\n B, C, H, W, kernel_size\n );\n\n return out;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor cross_scan_conv_cuda(torch::Tensor x, torch::Tensor weight, torch::Tensor bias, int kernel_size);\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=True,\n extra_cuda_cflags=['-O3']\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, channels: int, kernel_size: int = 3):\n super().__init__()\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 self.module = module\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n weight = self.dw_conv.weight\n bias = self.dw_conv.bias\n return self.module.cross_scan_conv_cuda(x, weight, bias, self.kernel_size)\n"
}
result ← submit_kernel (ok)
oksubmit_kernel
submit_kernel PASSED: (3 / 3) correctness trials all passed.
Kernel runtime: 0.03 μs
Runtime stats: mean=0.03μs median=?μs std=0.03μs
Numerical precision: max_abs_err=1.91e-06 mean_abs_err=9.86e-08 max_rel_err=3.70e-02
Memory: 35.0 MB peak (ref: 42.0 MB, ratio: 0.8333x)
Kernel launches: 5 (ref: 114, fusion ratio: 22.8)
SOL score: 0.3066
final submitted kernel (144 lines, 3,703 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>
#include <cuda.h>
#include <cuda_runtime.h>
__global__ void cross_scan_conv_kernel(
const float* __restrict__ x,
const float* __restrict__ weight,
const float* __restrict__ bias,
float* __restrict__ out,
int B, int C, int H, int W, int K
) {
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);
float sum = 0.0f;
// LR: row-major flatten
int pos_lr = h * W + w;
int L = H * W;
for (int k = 0; k < K; ++k) {
int p = pos_lr + (k - K/2);
if (p >= 0 && p < L) {
int hh = p / W;
int ww = p % W;
sum += x[((b * C + c) * H + hh) * W + ww] * weight[c * K + k];
}
}
sum += bias[c];
// RL: flip W
int pos_rl = h * W + (W - 1 - w);
for (int k = 0; k < K; ++k) {
int p = pos_rl + (k - K/2);
if (p >= 0 && p < L) {
int hh = p / W;
int ww = W - 1 - (p % W);
sum += x[((b * C + c) * H + hh) * W + ww] * weight[c * K + k];
}
}
sum += bias[c];
// TB: permute to (B,C,W,H), flatten -> vertical scans
int pos_tb = w * H + h;
int Lt = W * H;
for (int k = 0; k < K; ++k) {
int p = pos_tb + (k - K/2);
if (p >= 0 && p < Lt) {
int ww = p / H;
int hh = p % H;
sum += x[((b * C + c) * H + hh) * W + ww] * weight[c * K + k];
}
}
sum += bias[c];
// BT: flip H then TB
int pos_bt = w * H + (H - 1 - h);
for (int k = 0; k < K; ++k) {
int p = pos_bt + (k - K/2);
if (p >= 0 && p < Lt) {
int ww = p / H;
int hh = H - 1 - (p % H);
sum += x[((b * C + c) * H + hh) * W + ww] * weight[c * K + k];
}
}
sum += bias[c];
// residual
float xval = x[idx];
out[idx] = xval + sum;
}
torch::Tensor cross_scan_conv_cuda(
torch::Tensor x,
torch::Tensor weight,
torch::Tensor bias,
int kernel_size
) {
int B = x.size(0);
int C = x.size(1);
int H = x.size(2);
int W = x.size(3);
auto out = torch::zeros_like(x);
int total = B * C * H * W;
int threads = 256;
int blocks = (total + threads - 1) / threads;
cross_scan_conv_kernel<<<blocks, threads>>>(
x.data_ptr<float>(),
weight.data_ptr<float>(),
bias.data_ptr<float>(),
out.data_ptr<float>(),
B, C, H, W, kernel_size
);
return out;
}
'''
cpp_source = '''
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=True,
extra_cuda_cflags=['-O3']
)
class ModelNew(nn.Module):
def __init__(self, channels: int, kernel_size: int = 3):
super().__init__()
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,
)
self.module = module
def forward(self, x: torch.Tensor) -> torch.Tensor:
weight = self.dw_conv.weight
bias = self.dw_conv.bias
return self.module.cross_scan_conv_cuda(x, weight, bias, self.kernel_size)