import torch
import torch.nn as nn
import torch.nn.functional as F

class Model(nn.Module):
    def __init__(self, num_classes=1000):
        """
        EfficientNetB0 architecture implementation in PyTorch.

        :param num_classes: The number of output classes (default is 1000 for ImageNet).
        """
        super(Model, self).__init__()
        
        # Initial convolutional layer
        self.conv1 = nn.Conv2d(3, 32, kernel_size=3, stride=2, padding=1, bias=False)
        self.bn1 = nn.BatchNorm2d(32)
        
        # MBConv blocks
        self.blocks = nn.Sequential(
            # MBConv1 (32, 16, 1, 1)
            MBConv(32, 16, kernel_size=3, stride=1, expand_ratio=1),
            # MBConv6 (16, 24, 2, 6)
            MBConv(16, 24, kernel_size=3, stride=2, expand_ratio=6),
            # MBConv6 (24, 24, 1, 6)
            MBConv(24, 24, kernel_size=3, stride=1, expand_ratio=6),
            # MBConv6 (24, 40, 2, 6)
            MBConv(24, 40, kernel_size=5, stride=2, expand_ratio=6),
            # MBConv6 (40, 40, 1, 6)
            MBConv(40, 40, kernel_size=5, stride=1, expand_ratio=6),
            # MBConv6 (40, 80, 2, 6)
            MBConv(40, 80, kernel_size=3, stride=2, expand_ratio=6),
            # MBConv6 (80, 80, 1, 6)
            MBConv(80, 80, kernel_size=3, stride=1, expand_ratio=6),
            # MBConv6 (80, 112, 1, 6)
            MBConv(80, 112, kernel_size=5, stride=1, expand_ratio=6),
            # MBConv6 (112, 112, 1, 6)
            MBConv(112, 112, kernel_size=5, stride=1, expand_ratio=6),
            # MBConv6 (112, 192, 2, 6)
            MBConv(112, 192, kernel_size=5, stride=2, expand_ratio=6),
            # MBConv6 (192, 192, 1, 6)
            MBConv(192, 192, kernel_size=5, stride=1, expand_ratio=6),
            # MBConv6 (192, 192, 1, 6)
            MBConv(192, 192, kernel_size=5, stride=1, expand_ratio=6),
            # MBConv6 (192, 320, 1, 6)
            MBConv(192, 320, kernel_size=3, stride=1, expand_ratio=6)
        )
        
        # Final convolutional layer
        self.conv2 = nn.Conv2d(320, 1280, kernel_size=1, stride=1, padding=0, bias=False)
        self.bn2 = nn.BatchNorm2d(1280)
        
        # Fully connected layer
        self.fc = nn.Linear(1280, num_classes)
    
    def forward(self, x):
        """
        Forward pass of the EfficientNetB0 model.

        :param x: The input tensor, shape (batch_size, 3, 224, 224)
        :return: The output tensor, shape (batch_size, num_classes)
        """
        x = F.relu(self.bn1(self.conv1(x)))
        x = self.blocks(x)
        x = F.relu(self.bn2(self.conv2(x)))
        x = F.adaptive_avg_pool2d(x, (1, 1))
        x = x.view(x.size(0), -1)
        x = self.fc(x)
        return x

class MBConv(nn.Module):
    def __init__(self, in_channels, out_channels, kernel_size, stride, expand_ratio):
        """
        MBConv block implementation.

        :param in_channels: Number of input channels.
        :param out_channels: Number of output channels.
        :param kernel_size: Kernel size for the depthwise convolution.
        :param stride: Stride for the depthwise convolution.
        :param expand_ratio: Expansion ratio for the intermediate channels.
        """
        super(MBConv, self).__init__()
        
        self.use_residual = (stride == 1 and in_channels == out_channels)
        hidden_dim = in_channels * expand_ratio
        
        if expand_ratio != 1:
            self.expand_conv = nn.Sequential(
                nn.Conv2d(in_channels, hidden_dim, kernel_size=1, stride=1, padding=0, bias=False),
                nn.BatchNorm2d(hidden_dim),
                nn.ReLU6(inplace=True)
            )
        
        self.depthwise_conv = nn.Sequential(
            nn.Conv2d(hidden_dim, hidden_dim, kernel_size=kernel_size, stride=stride, padding=(kernel_size-1)//2, groups=hidden_dim, bias=False),
            nn.BatchNorm2d(hidden_dim),
            nn.ReLU6(inplace=True)
        )
        
        self.project_conv = nn.Sequential(
            nn.Conv2d(hidden_dim, out_channels, kernel_size=1, stride=1, padding=0, bias=False),
            nn.BatchNorm2d(out_channels)
        )
    
    def forward(self, x):
        """
        Forward pass of the MBConv block.

        :param x: The input tensor, shape (batch_size, in_channels, H, W)
        :return: The output tensor, shape (batch_size, out_channels, H', W')
        """
        identity = x
        
        if hasattr(self, 'expand_conv'):
            x = self.expand_conv(x)
        
        x = self.depthwise_conv(x)
        x = self.project_conv(x)
        
        if self.use_residual:
            x += identity
        
        return x

# Test code
batch_size = 10
num_classes = 1000

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

def get_init_inputs():
    return [num_classes]

import torch
import torch.nn as nn
import torch.nn.functional as F
from torch.utils.cpp_extension import load_inline

_cpp_src = r"""
#include <torch/extension.h>

torch::Tensor bn_act_(torch::Tensor x, torch::Tensor weight, torch::Tensor bias,
                      torch::Tensor mean, torch::Tensor var, double eps, int act);
torch::Tensor bn_add_(torch::Tensor x, torch::Tensor weight, torch::Tensor bias,
                      torch::Tensor mean, torch::Tensor var, torch::Tensor residual, double eps);
torch::Tensor avgpool2d(torch::Tensor x);
"""

_cuda_src = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>

__global__ void bn_act_ch_kernel(float* __restrict__ x,
                                 const float* __restrict__ weight,
                                 const float* __restrict__ bias,
                                 const float* __restrict__ mean,
                                 const float* __restrict__ var,
                                 int C, int HW, float eps, int act) {
    int nc = blockIdx.x;
    int c = nc % C;
    long base = (long)nc * (long)HW;
    float scale = weight[c] * rsqrtf(var[c] + eps);
    float shift = bias[c] - mean[c] * scale;
    for (int i = threadIdx.x; i < HW; i += blockDim.x) {
        float v = fmaf(x[base + i], scale, shift);
        if (act == 1) {
            v = v > 0.0f ? v : 0.0f;
        } else if (act == 2) {
            v = v > 0.0f ? v : 0.0f;
            v = v < 6.0f ? v : 6.0f;
        }
        x[base + i] = v;
    }
}

__global__ void bn_add_ch_kernel(float* __restrict__ x,
                                 const float* __restrict__ weight,
                                 const float* __restrict__ bias,
                                 const float* __restrict__ mean,
                                 const float* __restrict__ var,
                                 const float* __restrict__ residual,
                                 int C, int HW, float eps) {
    int nc = blockIdx.x;
    int c = nc % C;
    long base = (long)nc * (long)HW;
    float scale = weight[c] * rsqrtf(var[c] + eps);
    float shift = bias[c] - mean[c] * scale;
    for (int i = threadIdx.x; i < HW; i += blockDim.x) {
        float v = fmaf(x[base + i], scale, shift);
        x[base + i] = v + residual[base + i];
    }
}

__global__ void avgpool_ch_kernel(const float* __restrict__ x,
                                  float* __restrict__ out,
                                  int C, int HW) {
    int nc = blockIdx.x;
    long base = (long)nc * (long)HW;
    float sum = 0.0f;
    for (int i = threadIdx.x; i < HW; i += blockDim.x) {
        sum += x[base + i];
    }
    __shared__ float smem[256];
    smem[threadIdx.x] = sum;
    __syncthreads();
    for (int stride = blockDim.x >> 1; stride > 0; stride >>= 1) {
        if (threadIdx.x < stride) smem[threadIdx.x] += smem[threadIdx.x + stride];
        __syncthreads();
    }
    if (threadIdx.x == 0) out[nc] = smem[0] / (float)HW;
}

torch::Tensor bn_act_(torch::Tensor x, torch::Tensor weight, torch::Tensor bias,
                      torch::Tensor mean, torch::Tensor var, double eps, int act) {
    const int N = (int)x.size(0);
    const int C = (int)x.size(1);
    const int H = (int)x.size(2);
    const int W = (int)x.size(3);
    const int HW = H * W;
    const int blocks = N * C;
    bn_act_ch_kernel<<<blocks, 256, 0, at::cuda::getCurrentCUDAStream()>>>(
        x.data_ptr<float>(), weight.data_ptr<float>(), bias.data_ptr<float>(),
        mean.data_ptr<float>(), var.data_ptr<float>(), C, HW, (float)eps, act);
    return x;
}

torch::Tensor bn_add_(torch::Tensor x, torch::Tensor weight, torch::Tensor bias,
                      torch::Tensor mean, torch::Tensor var, torch::Tensor residual, double eps) {
    const int N = (int)x.size(0);
    const int C = (int)x.size(1);
    const int H = (int)x.size(2);
    const int W = (int)x.size(3);
    const int HW = H * W;
    const int blocks = N * C;
    bn_add_ch_kernel<<<blocks, 256, 0, at::cuda::getCurrentCUDAStream()>>>(
        x.data_ptr<float>(), weight.data_ptr<float>(), bias.data_ptr<float>(),
        mean.data_ptr<float>(), var.data_ptr<float>(), residual.data_ptr<float>(),
        C, HW, (float)eps);
    return x;
}

torch::Tensor avgpool2d(torch::Tensor x) {
    const int N = (int)x.size(0);
    const int C = (int)x.size(1);
    const int H = (int)x.size(2);
    const int W = (int)x.size(3);
    const int HW = H * W;
    auto out = torch::empty({N, C}, x.options());
    avgpool_ch_kernel<<<N * C, 256, 0, at::cuda::getCurrentCUDAStream()>>>(
        x.data_ptr<float>(), out.data_ptr<float>(), C, HW);
    return out;
}
"""

_effnet_fused = load_inline(
    name="effnet_b0_fused_bn_relu_ext",
    cpp_sources=_cpp_src,
    cuda_sources=_cuda_src,
    functions=["bn_act_", "bn_add_", "avgpool2d"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    verbose=False,
)


class _MBConvNew(nn.Module):
    def __init__(self, in_channels, out_channels, kernel_size, stride, expand_ratio):
        super(_MBConvNew, self).__init__()
        self.use_residual = (stride == 1 and in_channels == out_channels)
        hidden_dim = in_channels * expand_ratio
        if expand_ratio != 1:
            self.expand_conv = nn.Sequential(
                nn.Conv2d(in_channels, hidden_dim, kernel_size=1, stride=1, padding=0, bias=False),
                nn.BatchNorm2d(hidden_dim),
                nn.ReLU6(inplace=True),
            )
        self.depthwise_conv = nn.Sequential(
            nn.Conv2d(hidden_dim, hidden_dim, kernel_size=kernel_size, stride=stride,
                      padding=(kernel_size - 1) // 2, groups=hidden_dim, bias=False),
            nn.BatchNorm2d(hidden_dim),
            nn.ReLU6(inplace=True),
        )
        self.project_conv = nn.Sequential(
            nn.Conv2d(hidden_dim, out_channels, kernel_size=1, stride=1, padding=0, bias=False),
            nn.BatchNorm2d(out_channels),
        )


class ModelNew(nn.Module):
    def __init__(self, num_classes=1000):
        super(ModelNew, self).__init__()
        self.conv1 = nn.Conv2d(3, 32, kernel_size=3, stride=2, padding=1, bias=False)
        self.bn1 = nn.BatchNorm2d(32)
        self.blocks = nn.Sequential(
            _MBConvNew(32, 16, kernel_size=3, stride=1, expand_ratio=1),
            _MBConvNew(16, 24, kernel_size=3, stride=2, expand_ratio=6),
            _MBConvNew(24, 24, kernel_size=3, stride=1, expand_ratio=6),
            _MBConvNew(24, 40, kernel_size=5, stride=2, expand_ratio=6),
            _MBConvNew(40, 40, kernel_size=5, stride=1, expand_ratio=6),
            _MBConvNew(40, 80, kernel_size=3, stride=2, expand_ratio=6),
            _MBConvNew(80, 80, kernel_size=3, stride=1, expand_ratio=6),
            _MBConvNew(80, 112, kernel_size=5, stride=1, expand_ratio=6),
            _MBConvNew(112, 112, kernel_size=5, stride=1, expand_ratio=6),
            _MBConvNew(112, 192, kernel_size=5, stride=2, expand_ratio=6),
            _MBConvNew(192, 192, kernel_size=5, stride=1, expand_ratio=6),
            _MBConvNew(192, 192, kernel_size=5, stride=1, expand_ratio=6),
            _MBConvNew(192, 320, kernel_size=3, stride=1, expand_ratio=6),
        )
        self.conv2 = nn.Conv2d(320, 1280, kernel_size=1, stride=1, padding=0, bias=False)
        self.bn2 = nn.BatchNorm2d(1280)
        self.fc = nn.Linear(1280, num_classes)

    @staticmethod
    def _bn_act_inplace(x, bn, act):
        return _effnet_fused.bn_act_(x, bn.weight, bn.bias, bn.running_mean, bn.running_var, float(bn.eps), int(act))

    @staticmethod
    def _bn_add_inplace(x, bn, residual):
        return _effnet_fused.bn_add_(x, bn.weight, bn.bias, bn.running_mean, bn.running_var, residual, float(bn.eps))

    def _forward_block(self, block, x):
        identity = x
        if hasattr(block, 'expand_conv'):
            conv = block.expand_conv[0]
            x = F.conv2d(x, conv.weight, None, conv.stride, conv.padding, conv.dilation, conv.groups)
            x = self._bn_act_inplace(x, block.expand_conv[1], 2)  # ReLU6
        conv = block.depthwise_conv[0]
        x = F.conv2d(x, conv.weight, None, conv.stride, conv.padding, conv.dilation, conv.groups)
        x = self._bn_act_inplace(x, block.depthwise_conv[1], 2)  # ReLU6
        conv = block.project_conv[0]
        x = F.conv2d(x, conv.weight, None, conv.stride, conv.padding, conv.dilation, conv.groups)
        if block.use_residual:
            x = self._bn_add_inplace(x, block.project_conv[1], identity)
        else:
            x = self._bn_act_inplace(x, block.project_conv[1], 0)
        return x

    def forward(self, x):
        # Optimized inference path: cuDNN convolutions plus fused CUDA BN/activation kernels.
        x = F.conv2d(x, self.conv1.weight, None, self.conv1.stride, self.conv1.padding,
                     self.conv1.dilation, self.conv1.groups)
        x = self._bn_act_inplace(x, self.bn1, 1)  # ReLU
        for block in self.blocks:
            x = self._forward_block(block, x)
        x = F.conv2d(x, self.conv2.weight, None, self.conv2.stride, self.conv2.padding,
                     self.conv2.dilation, self.conv2.groups)
        x = self._bn_act_inplace(x, self.bn2, 1)  # ReLU
        x = _effnet_fused.avgpool2d(x)
        x = F.linear(x, self.fc.weight, self.fc.bias)
        return x