grok-4.3 · popcorn · pop_l123_all_gpt
correct 3_SSM2D_FusedBidirectionalCrossScan.py
model
grok-4.3
turns / max
3 / 15
tool calls
3
tokens (total)
18197
wall clock
278.7s
truncated
no
backend
cuda fp32
variant
popcorn
runtime
0.03 μs
ref runtime
0.17 μs
speedup
6.41x
started
2026-05-11T21:03:35Z
timingkernelbench.timing
kernel runtime
0.03 μs
ref runtime
0.17 μs
speedup
6.41x
kernel mean / median / std
0.03 / — / 0.03 μs
ref mean / median / std
0.17 / — / 0.09 μs
numerical precisioneval_kernel_against_ref
max abs error
1.91e-06
mean abs error
9.86e-08
max rel error
0.04
mean rel error
2.94e-07
gpu memoryextended_metrics.measure_memory
kernel peak
35.00 MB
ref peak
42.00 MB
kernel / ref
0.83x
kernel launches / fusiontorch.profiler
kernel launches
5
ref launches
114
fusion ratio (ref / kernel)
22.80x
total cuda time
173 μs
SOL (speed-of-light)nsight_hardware_counters
SOL score (0–1)
0.307
dram util %
2.8%
compute util %
30.7%
bottleneck
compute-bound
dominant pipe
FP32
arithmetic intensity
—
ridge point
20.00
roofline / occupancynsight_hardware_counters
dram bandwidth
92.43 GB/s
dram util %
2.8%
fp32 tflops
—
fp32 util %
—
fp16 tflops
—
occupancy %
81.4%
bottleneck
compute-bound
dominant pipe %
30.7%
arithmetic intensity
—
ridge point
20.00
L1 hit %
—
L2 hit %
170.8%
ld sectors / req
2.65
st sectors / req
—
regs / thread
48
smem / block
2048 B
block size
384
peak bw
3350.00 GB/s
peak fp32 tflops
67.00
peak fp16 tflops
1979.00
warp stalls (top 5)nsight
short_scoreboard
498.3%
mio_throttle
185.4%
long_scoreboard
171.2%
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)