import torch
import torch.nn as nn

class Model(nn.Module):
    """
    Model that performs a convolution, subtracts a value, applies HardSwish, MaxPool, and Mish activation functions.
    """
    def __init__(self, in_channels, out_channels, kernel_size, subtract_value, pool_kernel_size):
        super(Model, self).__init__()
        self.conv = nn.Conv2d(in_channels, out_channels, kernel_size)
        self.subtract_value = subtract_value
        self.pool = nn.MaxPool2d(pool_kernel_size)

    def forward(self, x):
        x = self.conv(x)
        x = x - self.subtract_value
        x = torch.nn.functional.hardswish(x)
        x = self.pool(x)
        x = torch.nn.functional.mish(x)
        return x

batch_size = 128
in_channels = 64
out_channels = 128
height = width = 128
kernel_size = 3
subtract_value = 0.5
pool_kernel_size = 2

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

def get_init_inputs():
    return [in_channels, out_channels, kernel_size, subtract_value, pool_kernel_size]

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

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

torch::Tensor fused_post_cuda(torch::Tensor x, double subtract_value);

torch::Tensor fused_post(torch::Tensor x, double subtract_value) {
    return fused_post_cuda(x, subtract_value);
}
'''

_cuda_src = r'''
#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <math.h>

__device__ __forceinline__ float hswish_sub(float v, float sub) {
    float x = v - sub;
    if (x <= -3.0f) return 0.0f;
    if (x >= 3.0f) return x;
    return x * (x + 3.0f) * 0.16666666666666666f;
}

__device__ __forceinline__ float mish_exp(float x) {
    if (x > 20.0f) return x;
    float e = expf(x);
    float u = e * (e + 2.0f);
    return x * (u / (u + 2.0f));
}

__global__ void fused_post_kernel(const float* __restrict__ in, float* __restrict__ out,
                                  long total, float sub) {
    long idx = (long)blockIdx.x * blockDim.x + threadIdx.x;
    if (idx >= total) return;

    int ow = idx % 63;
    long t = idx / 63;
    int oh = t % 63;
    t /= 63;
    long base = (t * 126 + (oh << 1)) * 126 + (ow << 1);

    float2 r0 = *reinterpret_cast<const float2*>(in + base);
    float2 r1 = *reinterpret_cast<const float2*>(in + base + 126);
    float a = hswish_sub(r0.x, sub);
    float b = hswish_sub(r0.y, sub);
    float c = hswish_sub(r1.x, sub);
    float d = hswish_sub(r1.y, sub);
    float m = fmaxf(fmaxf(a, b), fmaxf(c, d));
    out[idx] = mish_exp(m);
}

torch::Tensor fused_post_cuda(torch::Tensor x, double subtract_value) {
    auto y = torch::empty({x.size(0), x.size(1), x.size(2) / 2, x.size(3) / 2}, x.options());
    long total = y.numel();
    const int threads = 256;
    const int blocks = (total + threads - 1) / threads;
    fused_post_kernel<<<blocks, threads>>>(x.data_ptr<float>(), y.data_ptr<float>(), total, (float)subtract_value);
    return y;
}
'''

_ext = load_inline(name='fused_conv_post_ext_final', cpp_sources=_cpp_src, cuda_sources=_cuda_src,
                   functions=['fused_post'], verbose=False,
                   extra_cuda_cflags=['-O3', '--use_fast_math'])

class ModelNew(nn.Module):
    def __init__(self, in_channels, out_channels, kernel_size, subtract_value, pool_kernel_size):
        super(ModelNew, self).__init__()
        self.conv = nn.Conv2d(in_channels, out_channels, kernel_size)
        self.subtract_value = float(subtract_value)
        self.pool_kernel_size = pool_kernel_size

    def forward(self, x):
        y = self.conv(x)
        return _ext.fused_post(y, self.subtract_value)