FW-GLM-5-1 · popcorn · pop_l123_st_gpt
error 12_BatchNorm_GELU_DilatedConv1d_ResAdd.py
model
FW-GLM-5-1
turns / max
1 / 1
tool calls
1
tokens (total)
7042
wall clock
350.5s
truncated
yes
backend
cuda fp32
variant
popcorn
runtime
—
ref runtime
0.00 μs
speedup
—
started
2026-05-11T09:35:13Z
reference PyTorch source (the prompt the agent sees)
import torch
import torch.nn as nn

class Model(nn.Module):
    """
    Pre-norm 1D dilated convolution block with BatchNorm: BatchNorm1d, ReLU activation, dilated Conv1d, and residual addition.
    At high dilation rates, memory access becomes non-contiguous across large strides in the sequence dimension.
    """
    def __init__(self, channels, kernel_size, dilation):
        super(Model, self).__init__()
        self.norm = nn.BatchNorm1d(channels)
        self.activation = nn.ReLU()
        self.conv = nn.Conv1d(channels, channels, kernel_size, padding=dilation, dilation=dilation)

    def forward(self, x):
        residual = x
        x = self.norm(x)
        x = self.activation(x)
        x = self.conv(x)
        x = x + residual
        return x

batch_size = 32
channels = 256
seq_len = 4096
kernel_size = 3
dilation = 128

def get_inputs():
    return [torch.randn(popcorn_pri.jitter_int(batch_size), channels, popcorn_pri.jitter_int(seq_len))]

def get_init_inputs():
    return [channels, kernel_size, dilation]
turn 0 · 1 tool call(s) · 350.5s turn · 72.8s LLM · tokens in=1,059 · out=5,983 · total=7,042 · 25,208 chars (click to expand)
turn 0350.5s turn · 72.8s LLM · 277.7s tools · tokens in=1,059 · out=5,983 · total=7,042 · reasoning clipped
reasoning
reasoning (truncated 4616 chars): Let me analyze this task. I need to replace a PyTorch model that does: 1. BatchNorm1d 2. ReLU activation 3. Dilated Conv1d with padding=dilation 4. Residual addition (x + residual) The input dimensions are: - batch_size = 32 - channels = 256 - seq_len = 4096 - kernel_size = 3 - dilation = 128 This is a pre-norm block, meaning normalization happens before the convolution. Let me think about how to optimize this: 1. **BatchNorm1d**: For 1D data, BatchNorm normalizes across the batch and sequence dimensions for each channel. It has learnable parameters (weight and bias) and maintains running statistics (running_mean and running_var). In inference mode (eval), it uses the running statistics. 2. **ReLU**: Simple element-wise operation. 3. **Dilated Conv1d**: This is the tricky part. With dilation=128, the kernel accesses elements with stride 128 in the sequence dimension. This is highly non-contiguous memory access. 4. **Residual add**: Element-wise addition. For a fused kernel, I should: 1. Fuse BatchNorm + ReLU into a single kernel (this is common optimization) 2. For the dilated conv, I need to be careful about memory access patterns 3. The final addition can potentially be fused Let me write a custom CUDA kernel that fuses these operations: The dilated convolution with kernel_size=3 and dilation=128 means: - For each output position `out_pos`, we access: - `input[out_pos - 128]` (left kernel element) - `input[out_pos]` (center kernel element) - `input[out_pos + 128]` (right kernel element) With padding=dilation=128, the input sequence gets padded, and the output sequence length equals the input sequence length. The key challenge is optimizing memory access patterns. With dilation=128, each thread reads 3 elements separated by 128 positions in the sequence dimension. Since BatchNorm performs element-wise operations with running statistics, I can fuse it with ReLU and the residual addition into a single kernel. The dilated convolution remains the primary computational bottleneck due to its non-contiguous access pattern, requiring careful thread block and memory access design. For the implementation, I'll initialize BatchNorm and Conv1d layers with their respective parameters, then in the forward pass apply normalization, activation, convolution with dilation, and add the residual connection. fused operation. I can load the input element for the residual connection first, then apply BatchNorm and ReLU on it, reuse that same element for convolution, and finally add back the residual. This minimizes memory accesses by combining the operations sequentially. To implement this efficiently, I'll create a CUDA kernel that handles BatchNorm + ReLU fusion first, followed by dilated convolution with residual addition. The kernel will be compiled using `torch.utils.cpp_extension.load_inline` to bind the CUDA source at module load time, then invoked from `ModelNew.forward`. The kernel processes each element in the output tensor, computing global indices from the thread and block coordinates. It iterates through input channels, applying BatchNorm parameters (scale and shift) before accessing the 3-element kernel weights for each output channel. For each input channel, it reads three input values with dilation offsets and multiplies them by the corresponding kernel weights to compute the convolution output. Since this is a pointwise convolution where input and output channels match (256 each), the nested loop structure is unnecessary. The direct mapping allows optimization: each thread handles one output element, loading the BatchNorm-normalized input element, multiplying by the center kernel weight, and adding contributions from the left and right neighbors based on the dilation factor. The BatchNorm operation normalizes across the batch and sequence dimensions, computing per-channel statistics. The learnable gamma and beta parameters then scale and shift the normalized values. For the dilated convolution with kernel size 3, each position needs access to elements at positions offset by the dilation rate. When a single kernel fuses all operations, the challenge is that convolution requires access to multiple input positions that may fall outside a thread's assigned output position. For instance, computing the convolution output at position `p` requires normalized values at positions `p-128`, `p`, and `p+128`. This creates a dependency where threads need neighboring values, making a straightforward fusion difficult without careful handling of the dilation gaps. One approach is to separate the operations: first compute BatchNorm + ReLU across the entire tensor, then perform the convolution with residual addition. This avoids the data sharing problem between threads. Alternatively, I could use cooperative groups or shared memory to handle the overlapping access patterns, but the dilation of 128 makes this impractical since the gaps are too large to efficiently load into shared memory. For a fused kernel, each thread would load its own input element, apply the normalization and activation, store the intermediate result, then fetch the three required positions for the convolution. However, this introduces a race condition if threads write intermediate values that others need to read. The cleanest approach is to process the batch norm and ReLU first, then compute the convolution with residual addition in a single pass, or create a multi-stage kernel where threads compute their normalized values, synchronize, and then proceed with the convolution. For the implementation, I'll use a two-phase strategy: the first kernel applies batch normalization and ReLU in-place on the input tensor, while the second handles the dilated convolution and residual addition, reading from the normalized buffer and accumulating directly into the output. For the convolution phase, each thread processes one element, reading the normalized values at the three kernel positions determined by the dilation factor—left, center, and right—and computing the weighted sum with the channel-specific kernel weights and bias before adding the residual connection. I need to declare the variables before the conditional block so they're accessible afterward. Let me restructure the kernel to handle this properly by moving the declarations outside the bounds check. I need to declare variables outside the if statement so they're accessible when computing the output. The `in_left`, `in_center`, and `in_right` values along with `val` and `residual` should be declared at function scope, then assigned conditionally. If the index is in bounds, I apply BatchNorm and ReLU to compute the center value, then check bounds again for the left and right dilated positions, loading zeros if they're out of range. The residual and output are computed outside the conditional block since they're needed regardless. I see the issue - when out_idx is out of bounds, I'm setting values to 0.0f, but for positions within the padding region, I should actually access the input array. The structure needs adjustment: threads with valid out_idx process their output position, while others do nothing. For valid positions, I need to load the center input element, apply BatchNorm and ReLU, then access left and right neighbors based on dilation. The convolution itself applies kernel weights to these three values - left, center, and right - using the appropriate weight indices. Let me complete the computation with the bias and residual: the output combines the weighted sum of dilated neighbors plus the bias term, then adds the residual connection from the original input at that position. Finally, I store the result to the output array. The kernel launch parameters suggest a two-stage approach: first applying BatchNorm and ReLU, then computing the dilated convolution with residual addition. I need to reconsider the overall structure - the first kernel should transform the input through normalization and activation before the ...[ omitted 4616 characters ]... loat in_right = (s_right < seq_len) ? input[n_input * channels * seq_len + ic * seq_len + s_right] : 0.0f; out_val += w_left * in_left + w_center * in_center + w_right * in_right; } output[n_output * channels + oc] = out_val; } ``` The kernel now handles all three input channels by loading the appropriate input values based on the dilation offsets and multiplying them with their corresponding kernel weights. The bias is added at the start, and each input channel contributes to the final output value through the weighted sum. are declaring the CUDA kernel for the fused operation that combines batch normalization, ReLU activation, dilated convolution, and residual addition in a single pass. The kernel takes the input tensor, output tensor, and all the necessary parameters like channel dimensions, sequence length, dilation rate, and the learned weights and biases from both the batch norm and convolution layers. The forward pass function then sets up the kernel launch configuration by calculating the total number of elements and determining the appropriate block and grid dimensions before invoking the fused kernel on the GPU. For the weight tensor handling, I need to access the convolution's learnable parameters correctly. The weight is a 3D tensor with shape [out_channels, in_channels, kernel_size], and since this is a grouped convolution where groups=1 and in_channels equals out_channels, I access elements using the flattened indexing scheme: weight[out_c * in_channels * kernel_size + in_c * kernel_size + k]. The same indexing applies to the bias tensor, which is 1D with shape [out_channels]. For batch normalization in inference mode, I apply the transformation: (x - running_mean) / sqrt(running_var + eps) * weight + bias, where the running statistics and learnable parameters are stored as 1D tensors with shape [channels]. I need to verify the tensor dimensions match my indexing assumptions. The input has shape [batch, channels, seq_len] = [32, 256, 4096]. For the weight tensor of shape [out_channels, in_channels, kernel_size] = [256, 256, 3], I'll need to flatten the indexing appropriately using row-major ordering. The weight access formula becomes `weight[oc * C * K + ic * K + k]` where C=256 channels and K=3 kernel size. This gives me the correct flattened index for each weight element. For the input tensor with shape [N, C, L], the flattened access is `input[n * C * L + c * L + s]`. The output tensor with shape [N, C, L] follows the same pattern: `output[n * C * L + c * L + s]`. I need to check the convolution configuration carefully. The Conv1d has padding set to dilation (128), which means the output length equals the input length (4096). The effective kernel spans from -128 to +128 around each position, but with padding added only on the boundaries, positions near the edges will partially sample from padded regions. For instance, at position 0 the kernel tries to access -128 (padded, treated as 0) and 128 (valid), while at position 127 it accesses -1 (padded) and 255 (valid). The dilation of 128 creates this asymmetry where the kernel reaches far in both directions, but the symmetric padding only covers the true boundaries. When the kernel's dilated positions fall outside the valid range, those values should be treated as zero. For a given output position `s_out`, the input position accessed is `s_out - dilation + k * dilation` for each kernel element, and I need to check whether these fall within [0, seq_len). Since dilation=128 and kernel_size=3, the left kernel element accesses `s_out - 128`, the center accesses `s_out`, and the right accesses `s_out + 128`. The padding ensures that even when `s_out` is small, the computation handles out-of-bounds positions correctly by zeroing them out. -128) - Right: s + 128 - s = 127: Left=0 (just inside), Center=127, Right=255 - s = 128: Left=0, Center=128, Right=256 - ... - s = seq_len-1: Left=seq_len-129, Center=seq_len-1, Right=seq_len+127 (out of bounds) So I need to check both left and right positions for out-of-bounds and handle them appropriately. For optimization, the BatchNorm and ReLU operations can be fused with the convolution since both are element-wise. BatchNorm doesn't require inter-thread synchronization when computed per-position using pre-computed statistics, so I can integrate these operations together. The residual addition also fuses naturally. However, the dilated convolution introduces a memory access challenge—with dilation=128, the non-contiguous pattern could create scattered reads, though the center position access remains sequential. The key insight is that the center kernel position accesses memory contiguously when threads are laid out sequentially, while the left and right positions have fixed 128-stride offsets that actually maintain good coalescing within warps. If thread 0 accesses position s and thread 1 accesses position s+1, their left accesses hit s-128 and s-129 respectively, keeping the memory pattern relatively intact across the warp. Looking at the weight tensor layout in PyTorch—[out_channels, in_channels, kernel_size] with kernel_size indexed as [left, center, right] based on the implementation—I need to verify this ordering matches my kernel indexing. The standard convention for a 1D convolution kernel with dilation places the leftmost element first, so for dilation=d, the left element accesses input at s-dilation, the center at s, and the right at s+dilation, which aligns with my weight access pattern of [ic, 0], [ic, 1], [ic, 2]. Now I'm ready to construct the complete implementation using `torch.utils.cpp_extension.load_inline` to compile the CUDA kernel. I'll need to handle the function declarations carefully—`load_inline` automatically generates pybind11 bindings from the functions list, so I shouldn't write my own PYBIND11_MODULE. For the kernel itself, the `__global__` function doesn't need to be in the functions list since only the host wrapper gets exposed to Python. Looking at how `load_inline` works with the `functions` parameter, I see that the string entries in that list should match the exact function names I want to expose. I'll implement a fused kernel that combines batch normalization, ReLU, and convolution with residual addition, where the batch norm is applied per-channel using running statistics and learnable affine parameters, followed by ReLU activation and then the convolution operation with proper boundary handling for the dilation. For each output position, I iterate through all output channels and compute the convolution result by accumulating contributions from all input channels. The kernel handles the dilated convolution by computing the left, center, and right positions based on the dilation factor, with boundary checks to zero out contributions from positions outside the input sequence. After summing the weighted contributions from all input channels, I add the bias and residual term, then store the final result. The kernel launch is configured with a fixed block size of 256 threads, and the grid dimensions are calculated to cover all batch elements, channels, and sequence positions. The forward function takes the input tensor along with the necessary parameters and launches the fused kernel. Inside the forward pass, I need to verify the input dimensions match the expected shape, then extract the batch size, channel count, and sequence length. After validating that the weight tensor has the correct dimensions for the convolution, I set up the launch configuration with the number of threads per block and the total number of elements to process. The grid size is calculated to ensure we have enough blocks to cover all elements, then the kernel is launched with the input and output tensors along with the learned parameters like weight and bias, batch normalization statistics, and the convolution configuration details like dilation and kernel size. Finally, the output is returned.
call → submit_kernel
callsubmit_kernel { "kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\n# Define the CUDA kernel\ncuda_source = \"\"\"\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <cstdio>\n\n#define BLOCK_SIZE 256\n\n__global__ void fused_bn_relu_conv_residual_kernel(\n const float* __restrict__ input,\n float* __restrict__ output,\n const float* __restrict__ conv_weight,\n const float* __restrict__ conv_bias,\n const float* __restrict__ bn_weight,\n const float* __restrict__ bn_bias,\n const float* __restrict__ bn_running_mean,\n const float* __restrict__ bn_running_var,\n int batch_size,\n int channels,\n int seq_len,\n int dilation,\n int kernel_size,\n float eps\n) {\n int total_elements = batch_size * channels * seq_len;\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n \n if (idx >= total_elements) return;\n \n // Decode indices: input has shape [N, C, L]\n int s = idx % seq_len;\n int c = (idx / seq_len) % channels;\n int n = idx / (seq_len * channels);\n \n // Load input element\n float in_val = input[idx];\n \n // Apply BatchNorm (using running statistics for inference)\n float mean = bn_running_mean[c];\n float var = bn_running_var[c];\n float gamma = bn_weight[c];\n float beta = bn_bias[c];\n float normalized = (in_val - mean) * rsqrtf(var + eps);\n float bn_out = gamma * normalized + beta;\n \n // Apply ReLU\n float relu_out = fmaxf(0.0f, bn_out);\n \n // For dilated convolution with kernel_size=3, compute the three input positions\n // Weight layout: [out_channels, in_channels, kernel_size]\n // Kernel indices: 0 = left (s - dilation), 1 = center (s), 2 = right (s + dilation)\n \n // For this Conv1d, in_channels == out_channels == channels\n // Convolution: out[c_out] = sum over c_in of sum over k of weight[c_out, c_in, k] * relu_out[c_in, s_offset]\n \n // First, compute the contribution from the center kernel element for channel c\n // Since in_channels == out_channels, the output at channel c gets contribution from input channel c\n // But actually, a standard Conv1d sums over ALL input channels for each output channel\n \n // Let's compute the full convolution for output channel c at position s\n float conv_out = conv_bias[c];\n \n // Accumulate over all input channels\n for (int ic = 0; ic < channels; ++ic) {\n // Weight for output channel c, input channel ic\n // Shape: [out_channels, in_channels, kernel_size] = [channels, channels, 3]\n const float* w_ptr = conv_weight + c * channels * kernel_size + ic * kernel_size;\n float w_left = w_ptr[0];\n float w_center = w_ptr[1];\n float w_right = w_ptr[2];\n \n // Compute dilated positions\n int s_left = s - dilation;\n int s_center = s;\n int s_right = s + dilation;\n \n // Get input values for input channel ic at the three positions\n // Need to apply BatchNorm + ReLU to these values\n \n // For center position (always valid)\n int idx_center = n * channels * seq_len + ic * seq_len + s_center;\n float in_center = input[idx_center];\n float ic_mean = bn_running_mean[ic];\n float ic_var = bn_running_var[ic];\n float ic_gamma = bn_weight[ic];\n float ic_beta = bn_bias[ic];\n float ic_normalized = (in_center - ic_mean) * rsqrtf(ic_var + eps);\n float ic_bn_out = ic_gamma * ic_normalized + ic_beta;\n float ic_relu_center = fmaxf(0.0f, ic_bn_out);\n \n float ic_relu_left = 0.0f;\n if (s_left >= 0) {\n int idx_left = n * channels * seq_len + ic * seq_len + s_left;\n float in_left = input[idx_left];\n float left_normalized = (in_left - ic_mean) * rsqrtf(ic_var + eps);\n float left_bn_out = ic_gamma * left_normalized + ic_beta;\n ic_relu_left = fmaxf(0.0f, left_bn_out);\n }\n \n float ic_relu_right = 0.0f;\n if (s_right < seq_len) {\n int idx_right = n * channels * seq_len + ic * seq_len + s_right;\n float in_right = input[idx_right];\n float right_normalized = (in_right - ic_mean) * rsqrtf(ic_var + eps);\n float right_bn_out = ic_gamma * right_normalized + ic_beta;\n ic_relu_right = fmaxf(0.0f, right_bn_out);\n }\n \n conv_out += w_left * ic_relu_left + w_center * ic_relu_center + w_right * ic_relu_right;\n }\n \n // Add residual (the original input value at this position)\n float residual = in_val;\n float result = conv_out + residual;\n \n // Store output\n output[idx] = result;\n}\n\nvoid fused_bn_relu_conv_residual(\n const float* input,\n float* output,\n const float* conv_weight,\n const float* conv_bias,\n const float* bn_weight,\n const float* bn_bias,\n const float* bn_running_mean,\n const float* bn_running_var,\n int batch_size,\n int channels,\n int seq_len,\n int dilation,\n int kernel_size,\n float eps\n) {\n int total_elements = batch_size * channels * seq_len;\n int blocks = (total_elements + BLOCK_SIZE - 1) / BLOCK_SIZE;\n \n fused_bn_relu_conv_residual_kernel<<<blocks, BLOCK_SIZE>>>(\n input, output, conv_weight, conv_bias, bn_weight, bn_bias,\n bn_running_mean, bn_running_var, batch_size, channels, seq_len,\n dilation, kernel_size, eps\n );\n}\n\"\"\"\n\ncpp_source = \"\"\"\nvoid fused_bn_relu_conv_residual(\n const float* input,\n float* output,\n const float* conv_weight,\n const float* conv_bias,\n const float* bn_weight,\n const float* bn_bias,\n const float* bn_running_mean,\n const float* bn_running_var,\n int batch_size,\n int channels,\n int seq_len,\n int dilation,\n int kernel_size,\n float eps\n);\n\"\"\"\n\n# Compile the extension\nfused_ext = load_inline(\n name='fused_bn_relu_conv_residual_ext',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['fused_bn_relu_conv_residual'],\n verbose=True\n)\n\n\nclass ModelNew(nn.Module):\n \"\"\"\n Fused implementation of: BatchNorm1d -> ReLU -> Dilated Conv1d -> Residual Add\n \"\"\"\n def __init__(self, channels, kernel_size, dilation):\n super(ModelNew, self).__init__()\n self.channels = channels\n self.kernel_size = kernel_size\n self.dilation = dilation\n \n # Initialize layers with same parameters as reference\n self.norm = nn.BatchNorm1d(channels)\n self.conv = nn.Conv1d(channels, channels, kernel_size, padding=dilation, dilation=dilation)\n \n def forward(self, x):\n # Ensure we're in eval mode for consistent batch norm behavior\n if self.training:\n self.eval()\n \n batch_size = x.size(0)\n seq_len = x.size(2)\n \n # Allocate output tensor\n output = torch.empty_like(x)\n \n # Call fused kernel\n fused_ext.fused_bn_relu_conv_residual(\n x.data_ptr(),\n output.data_ptr(),\n self.conv.weight.data_ptr(),\n self.conv.bias.data_ptr(),\n self.norm.weight.data_ptr(),\n self.norm.bias.data_ptr(),\n self.norm.running_mean.data_ptr(),\n self.norm.running_var.data_ptr(),\n batch_size,\n self.channels,\n seq_len,\n self.dilation,\n self.kernel_size,\n self.norm.eps\n )\n \n return output" }
result ← submit_kernel (fail)
failsubmit_kernel submit_kernel FAILED: AcceleratorError: CUDA error: an illegal memory access was encountered Search for `cudaErrorIllegalAddress' in https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__TYPES.html for more information. CUDA kernel errors might be asynchronously reported at some other API call, so the stacktrace below might be incorrect. For debugging consider passing CUDA_LAUNCH_BLOCKING=1 Compile with `TORCH_USE_CUDA_DSA` to enable device-side assertions.
final submitted kernel (223 lines, 7,447 chars) — full source, untruncated
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline

# Define the CUDA kernel
cuda_source = """
#include <cuda.h>
#include <cuda_runtime.h>
#include <cstdio>

#define BLOCK_SIZE 256

__global__ void fused_bn_relu_conv_residual_kernel(
    const float* __restrict__ input,
    float* __restrict__ output,
    const float* __restrict__ conv_weight,
    const float* __restrict__ conv_bias,
    const float* __restrict__ bn_weight,
    const float* __restrict__ bn_bias,
    const float* __restrict__ bn_running_mean,
    const float* __restrict__ bn_running_var,
    int batch_size,
    int channels,
    int seq_len,
    int dilation,
    int kernel_size,
    float eps
) {
    int total_elements = batch_size * channels * seq_len;
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (idx >= total_elements) return;
    
    // Decode indices: input has shape [N, C, L]
    int s = idx % seq_len;
    int c = (idx / seq_len) % channels;
    int n = idx / (seq_len * channels);
    
    // Load input element
    float in_val = input[idx];
    
    // Apply BatchNorm (using running statistics for inference)
    float mean = bn_running_mean[c];
    float var = bn_running_var[c];
    float gamma = bn_weight[c];
    float beta = bn_bias[c];
    float normalized = (in_val - mean) * rsqrtf(var + eps);
    float bn_out = gamma * normalized + beta;
    
    // Apply ReLU
    float relu_out = fmaxf(0.0f, bn_out);
    
    // For dilated convolution with kernel_size=3, compute the three input positions
    // Weight layout: [out_channels, in_channels, kernel_size]
    // Kernel indices: 0 = left (s - dilation), 1 = center (s), 2 = right (s + dilation)
    
    // For this Conv1d, in_channels == out_channels == channels
    // Convolution: out[c_out] = sum over c_in of sum over k of weight[c_out, c_in, k] * relu_out[c_in, s_offset]
    
    // First, compute the contribution from the center kernel element for channel c
    // Since in_channels == out_channels, the output at channel c gets contribution from input channel c
    // But actually, a standard Conv1d sums over ALL input channels for each output channel
    
    // Let's compute the full convolution for output channel c at position s
    float conv_out = conv_bias[c];
    
    // Accumulate over all input channels
    for (int ic = 0; ic < channels; ++ic) {
        // Weight for output channel c, input channel ic
        // Shape: [out_channels, in_channels, kernel_size] = [channels, channels, 3]
        const float* w_ptr = conv_weight + c * channels * kernel_size + ic * kernel_size;
        float w_left = w_ptr[0];
        float w_center = w_ptr[1];
        float w_right = w_ptr[2];
        
        // Compute dilated positions
        int s_left = s - dilation;
        int s_center = s;
        int s_right = s + dilation;
        
        // Get input values for input channel ic at the three positions
        // Need to apply BatchNorm + ReLU to these values
        
        // For center position (always valid)
        int idx_center = n * channels * seq_len + ic * seq_len + s_center;
        float in_center = input[idx_center];
        float ic_mean = bn_running_mean[ic];
        float ic_var = bn_running_var[ic];
        float ic_gamma = bn_weight[ic];
        float ic_beta = bn_bias[ic];
        float ic_normalized = (in_center - ic_mean) * rsqrtf(ic_var + eps);
        float ic_bn_out = ic_gamma * ic_normalized + ic_beta;
        float ic_relu_center = fmaxf(0.0f, ic_bn_out);
        
        float ic_relu_left = 0.0f;
        if (s_left >= 0) {
            int idx_left = n * channels * seq_len + ic * seq_len + s_left;
            float in_left = input[idx_left];
            float left_normalized = (in_left - ic_mean) * rsqrtf(ic_var + eps);
            float left_bn_out = ic_gamma * left_normalized + ic_beta;
            ic_relu_left = fmaxf(0.0f, left_bn_out);
        }
        
        float ic_relu_right = 0.0f;
        if (s_right < seq_len) {
            int idx_right = n * channels * seq_len + ic * seq_len + s_right;
            float in_right = input[idx_right];
            float right_normalized = (in_right - ic_mean) * rsqrtf(ic_var + eps);
            float right_bn_out = ic_gamma * right_normalized + ic_beta;
            ic_relu_right = fmaxf(0.0f, right_bn_out);
        }
        
        conv_out += w_left * ic_relu_left + w_center * ic_relu_center + w_right * ic_relu_right;
    }
    
    // Add residual (the original input value at this position)
    float residual = in_val;
    float result = conv_out + residual;
    
    // Store output
    output[idx] = result;
}

void fused_bn_relu_conv_residual(
    const float* input,
    float* output,
    const float* conv_weight,
    const float* conv_bias,
    const float* bn_weight,
    const float* bn_bias,
    const float* bn_running_mean,
    const float* bn_running_var,
    int batch_size,
    int channels,
    int seq_len,
    int dilation,
    int kernel_size,
    float eps
) {
    int total_elements = batch_size * channels * seq_len;
    int blocks = (total_elements + BLOCK_SIZE - 1) / BLOCK_SIZE;
    
    fused_bn_relu_conv_residual_kernel<<<blocks, BLOCK_SIZE>>>(
        input, output, conv_weight, conv_bias, bn_weight, bn_bias,
        bn_running_mean, bn_running_var, batch_size, channels, seq_len,
        dilation, kernel_size, eps
    );
}
"""

cpp_source = """
void fused_bn_relu_conv_residual(
    const float* input,
    float* output,
    const float* conv_weight,
    const float* conv_bias,
    const float* bn_weight,
    const float* bn_bias,
    const float* bn_running_mean,
    const float* bn_running_var,
    int batch_size,
    int channels,
    int seq_len,
    int dilation,
    int kernel_size,
    float eps
);
"""

# Compile the extension
fused_ext = load_inline(
    name='fused_bn_relu_conv_residual_ext',
    cpp_sources=cpp_source,
    cuda_sources=cuda_source,
    functions=['fused_bn_relu_conv_residual'],
    verbose=True
)


class ModelNew(nn.Module):
    """
    Fused implementation of: BatchNorm1d -> ReLU -> Dilated Conv1d -> Residual Add
    """
    def __init__(self, channels, kernel_size, dilation):
        super(ModelNew, self).__init__()
        self.channels = channels
        self.kernel_size = kernel_size
        self.dilation = dilation
        
        # Initialize layers with same parameters as reference
        self.norm = nn.BatchNorm1d(channels)
        self.conv = nn.Conv1d(channels, channels, kernel_size, padding=dilation, dilation=dilation)
        
    def forward(self, x):
        # Ensure we're in eval mode for consistent batch norm behavior
        if self.training:
            self.eval()
        
        batch_size = x.size(0)
        seq_len = x.size(2)
        
        # Allocate output tensor
        output = torch.empty_like(x)
        
        # Call fused kernel
        fused_ext.fused_bn_relu_conv_residual(
            x.data_ptr(),
            output.data_ptr(),
            self.conv.weight.data_ptr(),
            self.conv.bias.data_ptr(),
            self.norm.weight.data_ptr(),
            self.norm.bias.data_ptr(),
            self.norm.running_mean.data_ptr(),
            self.norm.running_var.data_ptr(),
            batch_size,
            self.channels,
            seq_len,
            self.dilation,
            self.kernel_size,
            self.norm.eps
        )
        
        return output