gpt-5.5-priority · original · l123_single_gpt
compile_fail 16_DenseNet201.py
model
gpt-5.5-priority
turns / max
1 / 1
tool calls
1
tokens (total)
0
wall clock
—
truncated
no
backend
cuda fp32
variant
original
runtime
—
ref runtime
-1.00 μs
speedup
—
started
2026-05-05T00:51:17Z
compile error: Error building extension 'densenet201_fused_ext_v1': [1/3] /opt/conda/bin/x86_64-conda-linux-gnu-c++ -MMD -MF main.o.d -DTORCH_EXTENSION_NAME=densenet201_fused_ext_v1 -DTORCH_API_INCLUDE_EXTENSION_H -isystem /scratch/abaner/KernelBench/.venv/lib/python3.10/site-packages/torch/include -isystem /scratch/abaner/KernelBench/.venv/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /usr/local/cuda/include -isystem /opt/conda/include/python3.10 -fPIC -std=c++17 -O3 -c /scratch/abaner/KernelBench/runs/l123_single_gpt/original/gpt-5.5-priority/level_3_problem_16_cache/k_7310ca26
timingkernelbench.timing
kernel runtime
—
ref runtime
—
speedup
—
kernel mean / median / std
—
ref mean / median / std
—
reference PyTorch source (the prompt the agent sees)
import torch
import torch.nn as nn
import torch.nn.functional as F

class DenseBlock(nn.Module):
    def __init__(self, num_layers: int, num_input_features: int, growth_rate: int):
        """
        :param num_layers: The number of layers in the dense block
        :param num_input_features: The number of input feature maps
        :param growth_rate: The growth rate for the dense block (new features added per layer)
        """
        super(DenseBlock, self).__init__()
        layers = []
        for i in range(num_layers):
            layers.append(self._make_layer(num_input_features + i * growth_rate, growth_rate))
        self.layers = nn.ModuleList(layers)

    def _make_layer(self, in_features: int, growth_rate: int):
        """
        Creates a single layer with BatchNorm, ReLU, Conv2D, and Dropout.
        """
        return nn.Sequential(
            nn.BatchNorm2d(in_features),
            nn.ReLU(inplace=True),
            nn.Conv2d(in_features, growth_rate, kernel_size=3, padding=1, bias=False),
            nn.Dropout(0.0)
        )

    def forward(self, x):
        """
        :param x: Input tensor of shape (batch_size, num_input_features, height, width)
        :return: Concatenated output tensor with shape (batch_size, num_output_features, height, width)
        """
        features = [x]
        for layer in self.layers:
            new_feature = layer(x)
            features.append(new_feature)
            x = torch.cat(features, 1)  # Concatenate along channel axis
        return x

class TransitionLayer(nn.Module):
    def __init__(self, num_input_features: int, num_output_features: int):
        """
        :param num_input_features: The number of input feature maps
        :param num_output_features: The number of output feature maps
        """
        super(TransitionLayer, self).__init__()
        self.transition = nn.Sequential(
            nn.BatchNorm2d(num_input_features),
            nn.ReLU(inplace=True),
            nn.Conv2d(num_input_features, num_output_features, kernel_size=1, bias=False),
            nn.AvgPool2d(kernel_size=2, stride=2)
        )

    def forward(self, x):
        """
        :param x: Input tensor of shape (batch_size, num_input_features, height, width)
        :return: Downsampled tensor with reduced number of feature maps
        """
        return self.transition(x)

class Model(nn.Module):
    def __init__(self, growth_rate: int = 32, num_classes: int = 1000):
        """
        :param growth_rate: The growth rate of the DenseNet (new features added per layer)
        :param num_classes: The number of output classes for classification
        """
        super(Model, self).__init__()

        # Initial convolution and pooling
        self.features = nn.Sequential(
            nn.Conv2d(3, 64, kernel_size=7, stride=2, padding=3, bias=False),
            nn.BatchNorm2d(64),
            nn.ReLU(inplace=True),
            nn.MaxPool2d(kernel_size=3, stride=2, padding=1)
        )

        # Each dense block is followed by a transition layer, except the last one
        num_features = 64
        block_layers = [6, 12, 48, 32]  # Corresponding layers in DenseNet201

        self.dense_blocks = nn.ModuleList()
        self.transition_layers = nn.ModuleList()

        for i, num_layers in enumerate(block_layers):
            block = DenseBlock(num_layers=num_layers, num_input_features=num_features, growth_rate=growth_rate)
            self.dense_blocks.append(block)
            num_features = num_features + num_layers * growth_rate

            if i != len(block_layers) - 1:
                transition = TransitionLayer(num_input_features=num_features, num_output_features=num_features // 2)
                self.transition_layers.append(transition)
                num_features = num_features // 2

        # Final batch norm and classifier
        self.final_bn = nn.BatchNorm2d(num_features)
        self.classifier = nn.Linear(num_features, num_classes)

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        """
        :param x: Input tensor of shape (batch_size, 3, height, width)
        :return: Output tensor of shape (batch_size, num_classes)
        """
        x = self.features(x)

        for i, block in enumerate(self.dense_blocks):
            x = block(x)
            if i != len(self.dense_blocks) - 1:
                x = self.transition_layers[i](x)

        x = self.final_bn(x)
        x = F.relu(x, inplace=True)
        x = F.adaptive_avg_pool2d(x, (1, 1)).view(x.size(0), -1)
        x = self.classifier(x)
        return x

# Testing the DenseNet201 model
batch_size = 10
num_classes = 10
height, width = 224, 224  # Standard input size for DenseNet

def get_inputs():
    return [torch.rand(batch_size, 3, height, width)]

def get_init_inputs():
    return [32, num_classes]
turn 0 · FINAL · 1 tool call(s) · 36.1s turn · 36.1s LLM · 15,741 chars (click to expand)
turn 036.1s turn · 36.1s LLM · 0.0s tools · FINAL
reasoning
**Evaluating DenseNet implementation** I need to submit a code that implements a faster version of DenseNet. It’s a challenging task, and I’m considering using PyTorch modules or a custom CUDA kernel for the forward pass. I must ensure correctness and that the new model accepts the same arguments, possibly copying weights from the reference model. I might have to define the same parameter structure and ensure BatchNorm is properly managed during training and evaluation, while aiming for speed enhancements with optimized operations. **Exploring BatchNorm and Model Optimization** I'm considering default BatchNorm parameters like running mean as 0, variance as 1, weight as 1, and bias as 0 if it's not trained. In evaluation mode, BatchNorm might act like an identity function. For training, it uses batch statistics, which are a bit nontrivial. If I copy the state dictionary, I’ll need to keep everything aligned. I want to optimize the dense block concatenation in my new model due to inefficiencies in the reference DenseBlock, where full feature concatenation occurs at every layer. **Considering Layer Optimization in PyTorch** I'm looking at how PyTorch concatenates outputs across layers, which can lead to inefficiencies. To avoid repeatedly concatenating, I could preallocate output and compute layers sequentially. However, I need to consider that Conv2D requires contiguous input with growing channels. Maybe I can build the input from output slices and apply the sequential layer on non-contiguous outputs. I’m wondering if a strided slice might work and if that would result in a contiguous input for different channel configurations using the NCHW format. **Considering implementation details** I’m thinking about implementing a max pooling layer with a 3x3 size, stride of 2, and padding of 1 after the initial ReLU. I might use torch maxpool for this. I wonder if it's fast enough—maybe! However, with 98 layers and the many custom kernel launches, there could be overhead. Still, if I replace batch normalization and ReLU with a single launch, it might optimize performance while keeping convolution with CUDNN.
call → submit_kernel
callsubmit_kernel { "kernel_code": "import torch\nimport torch.nn as nn\nimport torch.nn.functional as F\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_src = r'''\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <vector>\n\n__global__ void bn_relu_kernel(const float* __restrict__ x, const float* __restrict__ w, const float* __restrict__ b,\n const float* __restrict__ mean, const float* __restrict__ var,\n float* __restrict__ y, long total, int C, int H, int W, float eps){\n long idx = blockIdx.x * blockDim.x + threadIdx.x;\n if(idx >= total) return;\n int hw = H*W;\n int c = (idx / hw) % C;\n float v = x[idx];\n float a = w[c] * rsqrtf(var[c] + eps);\n float z = (v - mean[c]) * a + b[c];\n y[idx] = z > 0.0f ? z : 0.0f;\n}\n\n__global__ void avgpool2_kernel(const float* __restrict__ x, float* __restrict__ y, int N, int C, int H, int W){\n long total = (long)N*C*(H/2)*(W/2);\n long idx = blockIdx.x * blockDim.x + threadIdx.x;\n if(idx >= total) return;\n int Wo = W/2, Ho = H/2;\n int ow = idx % Wo;\n int oh = (idx / Wo) % Ho;\n int c = (idx / (Wo*Ho)) % C;\n int n = idx / (Wo*Ho*C);\n long base = ((long)n*C + c)*H*W + (2*oh)*W + 2*ow;\n y[idx] = 0.25f*(x[base] + x[base+1] + x[base+W] + x[base+W+1]);\n}\n\n__global__ void final_linear_kernel(const float* __restrict__ x,\n const float* __restrict__ bn_w, const float* __restrict__ bn_b,\n const float* __restrict__ mean, const float* __restrict__ var,\n const float* __restrict__ fc_w, const float* __restrict__ fc_b,\n float* __restrict__ out,\n int N, int C, int H, int W, int K, float eps){\n int nk = blockIdx.x;\n int tid = threadIdx.x;\n int n = nk / K;\n int k = nk - n*K;\n extern __shared__ float s[];\n float sum = 0.0f;\n int HW = H*W;\n int tot = C*HW;\n for(int t = tid; t < tot; t += blockDim.x){\n int c = t / HW;\n int r = t - c*HW;\n float v = x[((long)n*C + c)*HW + r];\n float z = (v - mean[c]) * (bn_w[c] * rsqrtf(var[c] + eps)) + bn_b[c];\n z = z > 0.0f ? z : 0.0f;\n sum += z * (fc_w[(long)k*C + c] / (float)HW);\n }\n s[tid] = sum;\n __syncthreads();\n for(int stride = blockDim.x/2; stride > 0; stride >>= 1){\n if(tid < stride) s[tid] += s[tid+stride];\n __syncthreads();\n }\n if(tid == 0) out[nk] = s[0] + fc_b[k];\n}\n\ntorch::Tensor bn_relu(torch::Tensor x, torch::Tensor w, torch::Tensor b, torch::Tensor mean, torch::Tensor var, double eps){\n auto y = torch::empty_like(x);\n int N=x.size(0), C=x.size(1), H=x.size(2), W=x.size(3);\n long total = (long)N*C*H*W;\n int threads=256;\n int blocks=(total+threads-1)/threads;\n bn_relu_kernel<<<blocks, threads>>>(x.data_ptr<float>(), w.data_ptr<float>(), b.data_ptr<float>(), mean.data_ptr<float>(), var.data_ptr<float>(), y.data_ptr<float>(), total, C, H, W, (float)eps);\n return y;\n}\n\ntorch::Tensor avgpool2(torch::Tensor x){\n int N=x.size(0), C=x.size(1), H=x.size(2), W=x.size(3);\n auto y = torch::empty({N,C,H/2,W/2}, x.options());\n long total=(long)N*C*(H/2)*(W/2);\n int threads=256;\n int blocks=(total+threads-1)/threads;\n avgpool2_kernel<<<blocks, threads>>>(x.data_ptr<float>(), y.data_ptr<float>(), N,C,H,W);\n return y;\n}\n\ntorch::Tensor final_linear(torch::Tensor x, torch::Tensor bn_w, torch::Tensor bn_b, torch::Tensor mean, torch::Tensor var, torch::Tensor fc_w, torch::Tensor fc_b, double eps){\n int N=x.size(0), C=x.size(1), H=x.size(2), W=x.size(3), K=fc_w.size(0);\n auto out=torch::empty({N,K}, x.options());\n int threads=256;\n final_linear_kernel<<<N*K, threads, threads*sizeof(float)>>>(x.data_ptr<float>(), bn_w.data_ptr<float>(), bn_b.data_ptr<float>(), mean.data_ptr<float>(), var.data_ptr<float>(), fc_w.data_ptr<float>(), fc_b.data_ptr<float>(), out.data_ptr<float>(), N,C,H,W,K,(float)eps);\n return out;\n}\n'''\n\n_ext = load_inline(name='densenet201_fused_ext_v1', cpp_sources='', cuda_sources=cuda_src,\n functions=['bn_relu','avgpool2','final_linear'], verbose=False,\n extra_cuda_cflags=['-O3', '--use_fast_math'], extra_cflags=['-O3'])\n\nclass DenseBlock(nn.Module):\n def __init__(self, num_layers: int, num_input_features: int, growth_rate: int):\n super(DenseBlock, self).__init__()\n layers = []\n for i in range(num_layers):\n layers.append(self._make_layer(num_input_features + i * growth_rate, growth_rate))\n self.layers = nn.ModuleList(layers)\n def _make_layer(self, in_features: int, growth_rate: int):\n return nn.Sequential(\n nn.BatchNorm2d(in_features),\n nn.ReLU(inplace=True),\n nn.Conv2d(in_features, growth_rate, kernel_size=3, padding=1, bias=False),\n nn.Dropout(0.0)\n )\n\nclass TransitionLayer(nn.Module):\n def __init__(self, num_input_features: int, num_output_features: int):\n super(TransitionLayer, self).__init__()\n self.transition = nn.Sequential(\n nn.BatchNorm2d(num_input_features),\n nn.ReLU(inplace=True),\n nn.Conv2d(num_input_features, num_output_features, kernel_size=1, bias=False),\n nn.AvgPool2d(kernel_size=2, stride=2)\n )\n\nclass ModelNew(nn.Module):\n def __init__(self, growth_rate: int = 32, num_classes: int = 1000):\n super(ModelNew, self).__init__()\n self.features = nn.Sequential(\n nn.Conv2d(3, 64, kernel_size=7, stride=2, padding=3, bias=False),\n nn.BatchNorm2d(64),\n nn.ReLU(inplace=True),\n nn.MaxPool2d(kernel_size=3, stride=2, padding=1)\n )\n num_features = 64\n block_layers = [6, 12, 48, 32]\n self.dense_blocks = nn.ModuleList()\n self.transition_layers = nn.ModuleList()\n for i, num_layers in enumerate(block_layers):\n block = DenseBlock(num_layers=num_layers, num_input_features=num_features, growth_rate=growth_rate)\n self.dense_blocks.append(block)\n num_features = num_features + num_layers * growth_rate\n if i != len(block_layers) - 1:\n transition = TransitionLayer(num_input_features=num_features, num_output_features=num_features // 2)\n self.transition_layers.append(transition)\n num_features = num_features // 2\n self.final_bn = nn.BatchNorm2d(num_features)\n self.classifier = nn.Linear(num_features, num_classes)\n\n def _bn_relu(self, x, bn):\n # The benchmark uses fp32 CUDA inference. For training or CPU, preserve exact PyTorch semantics.\n if x.is_cuda and x.dtype is torch.float32 and (not self.training):\n return _ext.bn_relu(x.contiguous(), bn.weight, bn.bias, bn.running_mean, bn.running_var, bn.eps)\n return F.relu(bn(x), inplace=True)\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n if not x.is_cuda or x.dtype is not torch.float32 or self.training:\n x = self.features(x)\n for i, block in enumerate(self.dense_blocks):\n features=[x]\n for layer in block.layers:\n new_feature = layer(x)\n features.append(new_feature)\n x = torch.cat(features, 1)\n if i != len(self.dense_blocks)-1:\n x = self.transition_layers[i].transition(x)\n x = self.final_bn(x)\n x = F.relu(x, inplace=True)\n x = F.adaptive_avg_pool2d(x, (1,1)).view(x.size(0), -1)\n return self.classifier(x)\n\n # initial stem: fuse batchnorm+relu, leave highly optimized conv/maxpool to cuDNN/ATen\n x = self.features[0](x)\n x = self._bn_relu(x, self.features[1])\n x = self.features[3](x)\n\n for bi, block in enumerate(self.dense_blocks):\n feats = [x]\n for layer in block.layers:\n z = self._bn_relu(x, layer[0])\n nf = layer[2](z)\n feats.append(nf)\n x = torch.cat(feats, 1)\n if bi != len(self.dense_blocks) - 1:\n tr = self.transition_layers[bi].transition\n z = self._bn_relu(x, tr[0])\n x = tr[2](z)\n x = _ext.avgpool2(x.contiguous())\n\n return _ext.final_linear(x.contiguous(), self.final_bn.weight, self.final_bn.bias,\n self.final_bn.running_mean, self.final_bn.running_var,\n self.classifier.weight, self.classifier.bias, self.final_bn.eps)\n" }
result ← submit_kernel (fail)
failsubmit_kernel submit_kernel FAILED: kernel did not compile. Error building extension 'densenet201_fused_ext_v1': [1/3] /opt/conda/bin/x86_64-conda-linux-gnu-c++ -MMD -MF main.o.d -DTORCH_EXTENSION_NAME=densenet201_fused_ext_v1 -DTORCH_API_INCLUDE_EXTENSION_H -isystem /scratch/abaner/KernelBench/.venv/lib/python3.10/site-packages/torch/include -isystem /scratch/abaner/KernelBench/.venv/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /usr/local/cuda/include -isystem /opt/conda/include/python3.10 -fPIC -std=c++17 -O3 -c /scratch/abaner/KernelBench/runs/l123_single_gpt/original/gpt-5.5-priority/level_3_problem_16_cache/k_7310ca26d0d2/densenet201_fused_ext_v1/main.cpp -o main.o FAILED: [code=1] main.o /opt/conda/bin/x86_64-conda-linux-gnu-c++ -MMD -MF main.o.d -DTORCH_EXTENSION_NAME=densenet201_fused_ext_v1 -DTORCH_API_INCLUDE_EXTENSION_H -isystem /scratch/abaner/KernelBench/.venv/lib/python3.10/site-packages/torch/include -isystem /scratch/abaner/KernelBench/.venv/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /usr/local/cuda/include -isystem /opt/conda/include/python3.10 -fPIC -std=c++17 -O3 -c /scratch/abaner/KernelBench/runs/l123_single_gpt/original/gpt-5.5-priority/level_3_problem_16_cache/k_7310ca26d0d2/densenet201_fused_ext_v1/main.cpp -o main.o /scratch/abaner/KernelBench/runs/l123_single_gpt/original/gpt-5.5-priority/level_3_problem_16_cache/k_7310ca26d0d2/densenet201_fused_ext_v1/main.cpp: In function 'void pybind11_init_densenet201_fused_ext_v1(pybind11::module_&)': /scratch/abaner/KernelBench/runs/l123_single_gpt/original/gpt-5.5-priority/level_3_problem_16_cache/k_7310ca26d0d2/densenet201_fused_ext_v1/main.cpp:4:46: error: 'bn_relu' was not declared in this scope 4 | m.def("bn_relu", torch::wrap_pybind_function(bn_relu), "bn_relu"); | ^~~~~~~ /scratch/abaner/KernelBench/runs/l123_single_gpt/original/gpt-5.5-priority/level_3_problem_16_cache/k_7310ca26d0d2/densenet201_fused_ext_v1/main.cpp:5:47: error: 'avgpool2' was not declared in this scope 5 | m.def("avgpool2", torch::wrap_pybind_function(avgpool2), "avgpool2"); | ^~~~~~~~ /scratch/abaner/KernelBench/runs/l123_single_gpt/original/gpt-5.5-priority/level_3_problem_16_cache/k_7310ca26d0d2/densenet201_fused_ext_v1/main.cpp:6:51: error: 'final_linear' was not declared in this scope 6 | m.def("final_linear", torch::wrap_pybind_function(final_linear), "final_linear"); | ^~~~~~~~~~~~ [2/3] /usr/local/cuda/bin/nvcc -MD -MF cuda.cuda.o.d -ccbin /opt/conda/bin/x86_64-conda-linux-gnu-cc -DTORCH_EXTENSION_NAME=densenet201_fused_ext_v1 -DTORCH_API_INCLUDE_EXTENSION_H -isystem /scratch/abaner/KernelBench/.venv/lib/python3.10/site-packages/torch/include -isystem /scratch/abaner/KernelBench/.venv/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /usr/local/cuda/include -isystem /opt/conda/include/python3.10 -D__CUDA_NO_HALF_OPERATORS__ -D__CUDA_NO_HALF_CONVERSIONS__ -D__CUDA_NO_BFLOAT16_CONVERSIONS__ -D__CUDA_NO_HALF2_OPERATORS__ --expt-relaxed-constexpr -gencode=arch=compute_90,code=compute_90 -gencode=arch=compute_90,code=sm_90 --compiler-options '-fPIC' -O3 --use_fast_math -std=c++17 -c /scratch/abaner/KernelBench/runs/l123_single_gpt/original/gpt-5.5-priority/level_3_problem_16_cache/k_7310ca26d0d2/densenet201_fused_ext_v1/cuda.cu -o cuda.cuda.o nvcc warning : incompatible redefinition for option 'compiler-bindir', the last value of this option was used ninja: build stopped: subcommand failed.
final submitted kernel (194 lines, 8,700 chars) — full source, untruncated
import torch
import torch.nn as nn
import torch.nn.functional as F
from torch.utils.cpp_extension import load_inline

cuda_src = r'''
#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <vector>

__global__ void bn_relu_kernel(const float* __restrict__ x, const float* __restrict__ w, const float* __restrict__ b,
                               const float* __restrict__ mean, const float* __restrict__ var,
                               float* __restrict__ y, long total, int C, int H, int W, float eps){
    long idx = blockIdx.x * blockDim.x + threadIdx.x;
    if(idx >= total) return;
    int hw = H*W;
    int c = (idx / hw) % C;
    float v = x[idx];
    float a = w[c] * rsqrtf(var[c] + eps);
    float z = (v - mean[c]) * a + b[c];
    y[idx] = z > 0.0f ? z : 0.0f;
}

__global__ void avgpool2_kernel(const float* __restrict__ x, float* __restrict__ y, int N, int C, int H, int W){
    long total = (long)N*C*(H/2)*(W/2);
    long idx = blockIdx.x * blockDim.x + threadIdx.x;
    if(idx >= total) return;
    int Wo = W/2, Ho = H/2;
    int ow = idx % Wo;
    int oh = (idx / Wo) % Ho;
    int c = (idx / (Wo*Ho)) % C;
    int n = idx / (Wo*Ho*C);
    long base = ((long)n*C + c)*H*W + (2*oh)*W + 2*ow;
    y[idx] = 0.25f*(x[base] + x[base+1] + x[base+W] + x[base+W+1]);
}

__global__ void final_linear_kernel(const float* __restrict__ x,
                                    const float* __restrict__ bn_w, const float* __restrict__ bn_b,
                                    const float* __restrict__ mean, const float* __restrict__ var,
                                    const float* __restrict__ fc_w, const float* __restrict__ fc_b,
                                    float* __restrict__ out,
                                    int N, int C, int H, int W, int K, float eps){
    int nk = blockIdx.x;
    int tid = threadIdx.x;
    int n = nk / K;
    int k = nk - n*K;
    extern __shared__ float s[];
    float sum = 0.0f;
    int HW = H*W;
    int tot = C*HW;
    for(int t = tid; t < tot; t += blockDim.x){
        int c = t / HW;
        int r = t - c*HW;
        float v = x[((long)n*C + c)*HW + r];
        float z = (v - mean[c]) * (bn_w[c] * rsqrtf(var[c] + eps)) + bn_b[c];
        z = z > 0.0f ? z : 0.0f;
        sum += z * (fc_w[(long)k*C + c] / (float)HW);
    }
    s[tid] = sum;
    __syncthreads();
    for(int stride = blockDim.x/2; stride > 0; stride >>= 1){
        if(tid < stride) s[tid] += s[tid+stride];
        __syncthreads();
    }
    if(tid == 0) out[nk] = s[0] + fc_b[k];
}

torch::Tensor bn_relu(torch::Tensor x, torch::Tensor w, torch::Tensor b, torch::Tensor mean, torch::Tensor var, double eps){
    auto y = torch::empty_like(x);
    int N=x.size(0), C=x.size(1), H=x.size(2), W=x.size(3);
    long total = (long)N*C*H*W;
    int threads=256;
    int blocks=(total+threads-1)/threads;
    bn_relu_kernel<<<blocks, threads>>>(x.data_ptr<float>(), w.data_ptr<float>(), b.data_ptr<float>(), mean.data_ptr<float>(), var.data_ptr<float>(), y.data_ptr<float>(), total, C, H, W, (float)eps);
    return y;
}

torch::Tensor avgpool2(torch::Tensor x){
    int N=x.size(0), C=x.size(1), H=x.size(2), W=x.size(3);
    auto y = torch::empty({N,C,H/2,W/2}, x.options());
    long total=(long)N*C*(H/2)*(W/2);
    int threads=256;
    int blocks=(total+threads-1)/threads;
    avgpool2_kernel<<<blocks, threads>>>(x.data_ptr<float>(), y.data_ptr<float>(), N,C,H,W);
    return y;
}

torch::Tensor final_linear(torch::Tensor x, torch::Tensor bn_w, torch::Tensor bn_b, torch::Tensor mean, torch::Tensor var, torch::Tensor fc_w, torch::Tensor fc_b, double eps){
    int N=x.size(0), C=x.size(1), H=x.size(2), W=x.size(3), K=fc_w.size(0);
    auto out=torch::empty({N,K}, x.options());
    int threads=256;
    final_linear_kernel<<<N*K, threads, threads*sizeof(float)>>>(x.data_ptr<float>(), bn_w.data_ptr<float>(), bn_b.data_ptr<float>(), mean.data_ptr<float>(), var.data_ptr<float>(), fc_w.data_ptr<float>(), fc_b.data_ptr<float>(), out.data_ptr<float>(), N,C,H,W,K,(float)eps);
    return out;
}
'''

_ext = load_inline(name='densenet201_fused_ext_v1', cpp_sources='', cuda_sources=cuda_src,
                   functions=['bn_relu','avgpool2','final_linear'], verbose=False,
                   extra_cuda_cflags=['-O3', '--use_fast_math'], extra_cflags=['-O3'])

class DenseBlock(nn.Module):
    def __init__(self, num_layers: int, num_input_features: int, growth_rate: int):
        super(DenseBlock, self).__init__()
        layers = []
        for i in range(num_layers):
            layers.append(self._make_layer(num_input_features + i * growth_rate, growth_rate))
        self.layers = nn.ModuleList(layers)
    def _make_layer(self, in_features: int, growth_rate: int):
        return nn.Sequential(
            nn.BatchNorm2d(in_features),
            nn.ReLU(inplace=True),
            nn.Conv2d(in_features, growth_rate, kernel_size=3, padding=1, bias=False),
            nn.Dropout(0.0)
        )

class TransitionLayer(nn.Module):
    def __init__(self, num_input_features: int, num_output_features: int):
        super(TransitionLayer, self).__init__()
        self.transition = nn.Sequential(
            nn.BatchNorm2d(num_input_features),
            nn.ReLU(inplace=True),
            nn.Conv2d(num_input_features, num_output_features, kernel_size=1, bias=False),
            nn.AvgPool2d(kernel_size=2, stride=2)
        )

class ModelNew(nn.Module):
    def __init__(self, growth_rate: int = 32, num_classes: int = 1000):
        super(ModelNew, self).__init__()
        self.features = nn.Sequential(
            nn.Conv2d(3, 64, kernel_size=7, stride=2, padding=3, bias=False),
            nn.BatchNorm2d(64),
            nn.ReLU(inplace=True),
            nn.MaxPool2d(kernel_size=3, stride=2, padding=1)
        )
        num_features = 64
        block_layers = [6, 12, 48, 32]
        self.dense_blocks = nn.ModuleList()
        self.transition_layers = nn.ModuleList()
        for i, num_layers in enumerate(block_layers):
            block = DenseBlock(num_layers=num_layers, num_input_features=num_features, growth_rate=growth_rate)
            self.dense_blocks.append(block)
            num_features = num_features + num_layers * growth_rate
            if i != len(block_layers) - 1:
                transition = TransitionLayer(num_input_features=num_features, num_output_features=num_features // 2)
                self.transition_layers.append(transition)
                num_features = num_features // 2
        self.final_bn = nn.BatchNorm2d(num_features)
        self.classifier = nn.Linear(num_features, num_classes)

    def _bn_relu(self, x, bn):
        # The benchmark uses fp32 CUDA inference. For training or CPU, preserve exact PyTorch semantics.
        if x.is_cuda and x.dtype is torch.float32 and (not self.training):
            return _ext.bn_relu(x.contiguous(), bn.weight, bn.bias, bn.running_mean, bn.running_var, bn.eps)
        return F.relu(bn(x), inplace=True)

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        if not x.is_cuda or x.dtype is not torch.float32 or self.training:
            x = self.features(x)
            for i, block in enumerate(self.dense_blocks):
                features=[x]
                for layer in block.layers:
                    new_feature = layer(x)
                    features.append(new_feature)
                    x = torch.cat(features, 1)
                if i != len(self.dense_blocks)-1:
                    x = self.transition_layers[i].transition(x)
            x = self.final_bn(x)
            x = F.relu(x, inplace=True)
            x = F.adaptive_avg_pool2d(x, (1,1)).view(x.size(0), -1)
            return self.classifier(x)

        # initial stem: fuse batchnorm+relu, leave highly optimized conv/maxpool to cuDNN/ATen
        x = self.features[0](x)
        x = self._bn_relu(x, self.features[1])
        x = self.features[3](x)

        for bi, block in enumerate(self.dense_blocks):
            feats = [x]
            for layer in block.layers:
                z = self._bn_relu(x, layer[0])
                nf = layer[2](z)
                feats.append(nf)
                x = torch.cat(feats, 1)
            if bi != len(self.dense_blocks) - 1:
                tr = self.transition_layers[bi].transition
                z = self._bn_relu(x, tr[0])
                x = tr[2](z)
                x = _ext.avgpool2(x.contiguous())

        return _ext.final_linear(x.contiguous(), self.final_bn.weight, self.final_bn.bias,
                                 self.final_bn.running_mean, self.final_bn.running_var,
                                 self.classifier.weight, self.classifier.bias, self.final_bn.eps)