correct
3_SSM2D_FusedBidirectionalCrossScan.py
started
2026-05-11T13:21:48Z
timingkernelbench.timing
kernel mean / median / std
0.01 / — / 1.57e-04 μs
ref mean / median / std
0.25 / — / 1.14 μs
numerical precisioneval_kernel_against_ref
gpu memoryextended_metrics.measure_memory
kernel launches / fusiontorch.profiler
fusion ratio (ref / kernel)
38.00x
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 071.8s turn · 2.9s LLM · 68.9s tools · tokens in=2,194 · out=6 · total=2,361
reasoning
The task is to replace the forward computation with a custom CUDA kernel for faster execution while maintaining numerical equivalence.
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 1320.3s turn · 17.7s LLM · 302.6s tools · tokens in=2,594 · out=1,097 · total=3,936
reasoning
H100 with high bandwidth. The operation is memory-bound likely, since conv1d with k=3 on sequences of length H*W=1024, channels=64, batch=4.
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_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\n {\n int pos = h * W + w;\n for (int k = 0; k < K; ++k) {\n int p = pos + k - K/2;\n if (p >= 0 && p < H*W) {\n int ph = p / W;\n int pw = p % W;\n sum += x[((b*C + c)*H + ph)*W + pw] * weight[c*K + k];\n }\n }\n }\n\n // RL\n {\n int pos = h * W + (W - 1 - w);\n for (int k = 0; k < K; ++k) {\n int p = pos + k - K/2;\n if (p >= 0 && p < H*W) {\n int ph = p / W;\n int pw = p % W;\n int orig_w = W-1 - pw;\n sum += x[((b*C + c)*H + ph)*W + orig_w] * weight[c*K + k];\n }\n }\n }\n\n // TB\n {\n int pos = w * H + h;\n for (int k = 0; k < K; ++k) {\n int p = pos + k - K/2;\n if (p >= 0 && p < W*H) {\n int pw = p / H;\n int ph = p % H;\n sum += x[((b*C + c)*H + ph)*W + pw] * weight[c*K + k];\n }\n }\n }\n\n // BT\n {\n int pos = w * H + (H - 1 - h);\n for (int k = 0; k < K; ++k) {\n int p = pos + k - K/2;\n if (p >= 0 && p < W*H) {\n int pw = p / H;\n int ph = p % H;\n int orig_h = H-1 - ph;\n sum += x[((b*C + c)*H + orig_h)*W + pw] * weight[c*K + k];\n }\n }\n }\n\n float res = x[idx] + sum + bias[c];\n out[idx] = res;\n}\n\ntorch::Tensor cross_scan_conv_cuda(torch::Tensor x, torch::Tensor weight, torch::Tensor bias, int kernel_size) {\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 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 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 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 self.module = module\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n weight = self.dw_conv.weight.view(self.channels, self.kernel_size)\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 2166.7s turn · 10.1s LLM · 156.6s tools · tokens in=3,953 · out=1,098 · total=5,058
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_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\n {\n int pos = h * W + w;\n for (int k = 0; k < K; ++k) {\n int p = pos + k - K/2;\n if (p >= 0 && p < H*W) {\n int ph = p / W;\n int pw = p % W;\n sum += x[((b*C + c)*H + ph)*W + pw] * weight[c*K + k];\n }\n }\n }\n\n // RL\n {\n int pos = h * W + (W - 1 - w);\n for (int k = 0; k < K; ++k) {\n int p = pos + k - K/2;\n if (p >= 0 && p < H*W) {\n int ph = p / W;\n int pw = p % W;\n int orig_w = W-1 - pw;\n sum += x[((b*C + c)*H + ph)*W + orig_w] * weight[c*K + k];\n }\n }\n }\n\n // TB\n {\n int pos = w * H + h;\n for (int k = 0; k < K; ++k) {\n int p = pos + k - K/2;\n if (p >= 0 && p < W*H) {\n int pw = p / H;\n int ph = p % H;\n sum += x[((b*C + c)*H + ph)*W + pw] * weight[c*K + k];\n }\n }\n }\n\n // BT\n {\n int pos = w * H + (H - 1 - h);\n for (int k = 0; k < K; ++k) {\n int p = pos + k - K/2;\n if (p >= 0 && p < W*H) {\n int pw = p / H;\n int ph = p % H;\n int orig_h = H-1 - ph;\n sum += x[((b*C + c)*H + orig_h)*W + pw] * weight[c*K + k];\n }\n }\n }\n\n float res = x[idx] + sum + bias[c];\n out[idx] = res;\n}\n\ntorch::Tensor cross_scan_conv_cuda(torch::Tensor x, torch::Tensor weight, torch::Tensor bias, int kernel_size) {\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 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 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 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 self.module = module\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n weight = self.dw_conv.weight.view(self.channels, self.kernel_size)\n bias = self.dw_conv.bias\n return self.module.cross_scan_conv_cuda(x, weight, bias, self.kernel_size)\n"
}
result ← run_correctness (fail)
failrun_correctness
run_correctness FAILED: (0 / 5) trials did not all match.
correctness_issue: Output mismatch
max_difference: ['1.723052', '1.723052', '1.723052', '1.723052', '1.723052']
avg_difference: ['0.888087', '0.888087', '0.888087', '0.888087', '0.888087']
turn 3100.8s turn · 15.6s LLM · 85.2s tools · tokens in=5,155 · out=1,103 · total=6,463
reasoning
There's a bug in the kernel. The sum is missing the bias? No, I added bias at the end, but in reference, the bias is added per conv, so four times?
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_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\n {\n int pos = h * W + w;\n for (int k = 0; k < K; ++k) {\n int p = pos + k - K/2;\n if (p >= 0 && p < H*W) {\n int ph = p / W;\n int pw = p % W;\n sum += x[((b*C + c)*H + ph)*W + pw] * weight[c*K + k];\n }\n }\n }\n\n // RL\n {\n int pos = h * W + (W - 1 - w);\n for (int k = 0; k < K; ++k) {\n int p = pos + k - K/2;\n if (p >= 0 && p < H*W) {\n int ph = p / W;\n int pw = p % W;\n int orig_w = W-1 - pw;\n sum += x[((b*C + c)*H + ph)*W + orig_w] * weight[c*K + k];\n }\n }\n }\n\n // TB\n {\n int pos = w * H + h;\n for (int k = 0; k < K; ++k) {\n int p = pos + k - K/2;\n if (p >= 0 && p < W*H) {\n int pw = p / H;\n int ph = p % H;\n sum += x[((b*C + c)*H + ph)*W + pw] * weight[c*K + k];\n }\n }\n }\n\n // BT\n {\n int pos = w * H + (H - 1 - h);\n for (int k = 0; k < K; ++k) {\n int p = pos + k - K/2;\n if (p >= 0 && p < W*H) {\n int pw = p / H;\n int ph = p % H;\n int orig_h = H-1 - ph;\n sum += x[((b*C + c)*H + orig_h)*W + pw] * weight[c*K + k];\n }\n }\n }\n\n float res = x[idx] + sum + 4.0f * bias[c];\n out[idx] = res;\n}\n\ntorch::Tensor cross_scan_conv_cuda(torch::Tensor x, torch::Tensor weight, torch::Tensor bias, int kernel_size) {\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 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 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 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 self.module = module\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n weight = self.dw_conv.weight.view(self.channels, self.kernel_size)\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 4318.5s turn · 17.3s LLM · 301.2s tools · tokens in=6,480 · out=1,104 · total=7,588
reasoning
Now run 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_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\n {\n int pos = h * W + w;\n for (int k = 0; k < K; ++k) {\n int p = pos + k - K/2;\n if (p >= 0 && p < H*W) {\n int ph = p / W;\n int pw = p % W;\n sum += x[((b*C + c)*H + ph)*W + pw] * weight[c*K + k];\n }\n }\n }\n\n // RL\n {\n int pos = h * W + (W - 1 - w);\n for (int k = 0; k < K; ++k) {\n int p = pos + k - K/2;\n if (p >= 0 && p < H*W) {\n int ph = p / W;\n int pw = p % W;\n int orig_w = W-1 - pw;\n sum += x[((b*C + c)*H + ph)*W + orig_w] * weight[c*K + k];\n }\n }\n }\n\n // TB\n {\n int pos = w * H + h;\n for (int k = 0; k < K; ++k) {\n int p = pos + k - K/2;\n if (p >= 0 && p < W*H) {\n int pw = p / H;\n int ph = p % H;\n sum += x[((b*C + c)*H + ph)*W + pw] * weight[c*K + k];\n }\n }\n }\n\n // BT\n {\n int pos = w * H + (H - 1 - h);\n for (int k = 0; k < K; ++k) {\n int p = pos + k - K/2;\n if (p >= 0 && p < W*H) {\n int pw = p / H;\n int ph = p % H;\n int orig_h = H-1 - ph;\n sum += x[((b*C + c)*H + orig_h)*W + pw] * weight[c*K + k];\n }\n }\n }\n\n float res = x[idx] + sum + 4.0f * bias[c];\n out[idx] = res;\n}\n\ntorch::Tensor cross_scan_conv_cuda(torch::Tensor x, torch::Tensor weight, torch::Tensor bias, int kernel_size) {\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 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 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 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 self.module = module\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n weight = self.dw_conv.weight.view(self.channels, self.kernel_size)\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: (5 / 5) trials all matched the reference.
Numerical precision: max_abs_err=2.38e-06 mean_abs_err=1.00e-07 max_rel_err=7.69e-02
turn 5311.6s turn · 10.8s LLM · 300.8s tools · tokens in=7,648 · out=1,103 · total=8,784 · FINAL
reasoning
Correct now! Great, max error small.
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_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\n {\n int pos = h * W + w;\n for (int k = 0; k < K; ++k) {\n int p = pos + k - K/2;\n if (p >= 0 && p < H*W) {\n int ph = p / W;\n int pw = p % W;\n sum += x[((b*C + c)*H + ph)*W + pw] * weight[c*K + k];\n }\n }\n }\n\n // RL\n {\n int pos = h * W + (W - 1 - w);\n for (int k = 0; k < K; ++k) {\n int p = pos + k - K/2;\n if (p >= 0 && p < H*W) {\n int ph = p / W;\n int pw = p % W;\n int orig_w = W-1 - pw;\n sum += x[((b*C + c)*H + ph)*W + orig_w] * weight[c*K + k];\n }\n }\n }\n\n // TB\n {\n int pos = w * H + h;\n for (int k = 0; k < K; ++k) {\n int p = pos + k - K/2;\n if (p >= 0 && p < W*H) {\n int pw = p / H;\n int ph = p % H;\n sum += x[((b*C + c)*H + ph)*W + pw] * weight[c*K + k];\n }\n }\n }\n\n // BT\n {\n int pos = w * H + (H - 1 - h);\n for (int k = 0; k < K; ++k) {\n int p = pos + k - K/2;\n if (p >= 0 && p < W*H) {\n int pw = p / H;\n int ph = p % H;\n int orig_h = H-1 - ph;\n sum += x[((b*C + c)*H + orig_h)*W + pw] * weight[c*K + k];\n }\n }\n }\n\n float res = x[idx] + sum + 4.0f * bias[c];\n out[idx] = res;\n}\n\ntorch::Tensor cross_scan_conv_cuda(torch::Tensor x, torch::Tensor weight, torch::Tensor bias, int kernel_size) {\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 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 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 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 self.module = module\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n weight = self.dw_conv.weight.view(self.channels, self.kernel_size)\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: (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=2.38e-06 mean_abs_err=1.00e-07 max_rel_err=7.69e-02
Memory: 34.0 MB peak (ref: 41.0 MB, ratio: 0.8293x)
Kernel launches: 3 (ref: 114, fusion ratio: 38.0)
SOL score: 0.278
Flagged for excessive speedup — this submission may have been rejected by automated review.
final submitted kernel (143 lines, 3,762 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_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
{
int pos = h * W + w;
for (int k = 0; k < K; ++k) {
int p = pos + k - K/2;
if (p >= 0 && p < H*W) {
int ph = p / W;
int pw = p % W;
sum += x[((b*C + c)*H + ph)*W + pw] * weight[c*K + k];
}
}
}
// RL
{
int pos = h * W + (W - 1 - w);
for (int k = 0; k < K; ++k) {
int p = pos + k - K/2;
if (p >= 0 && p < H*W) {
int ph = p / W;
int pw = p % W;
int orig_w = W-1 - pw;
sum += x[((b*C + c)*H + ph)*W + orig_w] * weight[c*K + k];
}
}
}
// TB
{
int pos = w * H + h;
for (int k = 0; k < K; ++k) {
int p = pos + k - K/2;
if (p >= 0 && p < W*H) {
int pw = p / H;
int ph = p % H;
sum += x[((b*C + c)*H + ph)*W + pw] * weight[c*K + k];
}
}
}
// BT
{
int pos = w * H + (H - 1 - h);
for (int k = 0; k < K; ++k) {
int p = pos + k - K/2;
if (p >= 0 && p < W*H) {
int pw = p / H;
int ph = p % H;
int orig_h = H-1 - ph;
sum += x[((b*C + c)*H + orig_h)*W + pw] * weight[c*K + k];
}
}
}
float res = x[idx] + sum + 4.0f * bias[c];
out[idx] = res;
}
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);
auto out = torch::empty_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__()
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,
)
self.module = module
def forward(self, x: torch.Tensor) -> torch.Tensor:
weight = self.dw_conv.weight.view(self.channels, self.kernel_size)
bias = self.dw_conv.bias
return self.module.cross_scan_conv_cuda(x, weight, bias, self.kernel_size)