```
# Copyright (c) Meta Platforms, Inc. and affiliates.
# All rights reserved.
#
# This source code is licensed under the license found in the
# LICENSE file in the root directory of this source tree.

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

torch.backends.cudnn.benchmark = True

# Fused epilogue: y = hardtanh(mish(x) + add_value, min_val, max_val) * scale
# This replaces 4 separate memory-bound pointwise kernels (mish, add, hardtanh, mul)
# with a single vectorized (float4) pass over the tensor.
_fused_epilogue_cuda_src = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>

__device__ __forceinline__ float mish_fwd(float x) {
    // softplus with the same numerical-stability threshold (20) used by PyTorch
    float sp = (x > 20.0f) ? x : log1pf(expf(x));
    return x * tanhf(sp);
}

__device__ __forceinline__ float epilogue_op(float x, float add_value,
                                             float min_val, float max_val,
                                             float scale) {
    float z = mish_fwd(x) + add_value;
    z = fminf(fmaxf(z, min_val), max_val);   // hardtanh
    return z * scale;
}

__global__ void epilogue_vec4_kernel(const float4* __restrict__ in,
                                     float4* __restrict__ out,
                                     long n4,
                                     float add_value, float min_val,
                                     float max_val, float scale) {
    long i = blockIdx.x * (long)blockDim.x + threadIdx.x;
    if (i < n4) {
        float4 v = in[i];
        v.x = epilogue_op(v.x, add_value, min_val, max_val, scale);
        v.y = epilogue_op(v.y, add_value, min_val, max_val, scale);
        v.z = epilogue_op(v.z, add_value, min_val, max_val, scale);
        v.w = epilogue_op(v.w, add_value, min_val, max_val, scale);
        out[i] = v;
    }
}

__global__ void epilogue_tail_kernel(const float* __restrict__ in,
                                     float* __restrict__ out,
                                     long start, long n,
                                     float add_value, float min_val,
                                     float max_val, float scale) {
    long i = start + blockIdx.x * (long)blockDim.x + threadIdx.x;
    if (i < n) {
        out[i] = epilogue_op(in[i], add_value, min_val, max_val, scale);
    }
}

torch::Tensor fused_epilogue(torch::Tensor x, double add_value, double min_val,
                             double max_val, double scale) {
    TORCH_CHECK(x.is_cuda(), "input must be a CUDA tensor");
    TORCH_CHECK(x.scalar_type() == at::kFloat, "input must be float32");

    auto xc = x.contiguous();
    auto out = torch::empty_like(xc);
    long n = xc.numel();
    if (n == 0) return out;

    const float* in_ptr = xc.data_ptr<float>();
    float* out_ptr = out.data_ptr<float>();
    auto stream = at::cuda::getCurrentCUDAStream();

    long n4 = n >> 2;
    if (n4 > 0) {
        int block = 256;
        long grid = (n4 + block - 1) / block;
        epilogue_vec4_kernel<<<grid, block, 0, stream>>>(
            reinterpret_cast<const float4*>(in_ptr),
            reinterpret_cast<float4*>(out_ptr),
            n4,
            (float)add_value, (float)min_val, (float)max_val, (float)scale);
    }
    long start = n4 << 2;
    long rem = n - start;
    if (rem > 0) {
        epilogue_tail_kernel<<<1, (int)rem, 0, stream>>>(
            in_ptr, out_ptr, start, n,
            (float)add_value, (float)min_val, (float)max_val, (float)scale);
    }
    return out;
}
"""

_fused_epilogue_cpp_src = (
    "torch::Tensor fused_epilogue(torch::Tensor x, double add_value, double min_val, double max_val, double scale);"
)

_fused_epilogue = load_inline(
    name="fused_mish_add_hardtanh_scale",
    cpp_sources=_fused_epilogue_cpp_src,
    cuda_sources=_fused_epilogue_cuda_src,
    functions=["fused_epilogue"],
    verbose=False,
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Optimized model: ConvTranspose2d runs on cuDNN, while the Mish + add +
    Hardtanh + scale chain is fused into a single custom CUDA kernel,
    eliminating 3 extra full-tensor read/write passes and kernel launches.
    """
    def __init__(self, in_channels, out_channels, kernel_size, stride, padding, output_padding, add_value, scale):
        super(ModelNew, self).__init__()
        self.conv_transpose = nn.ConvTranspose2d(in_channels, out_channels, kernel_size, stride, padding, output_padding)
        self.add_value = float(add_value)
        self.scale = float(scale)

    def forward(self, x):
        x = self.conv_transpose(x)
        if x.is_cuda and x.dtype == torch.float32:
            return _fused_epilogue.fused_epilogue(x, self.add_value, -1.0, 1.0, self.scale)
        # Fallback for non-CUDA / non-float32 inputs
        x = torch.nn.functional.mish(x)
        x = x + self.add_value
        x = torch.nn.functional.hardtanh(x, min_val=-1, max_val=1)
        x = x * self.scale
        return x
```