```
# 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

# Let cuDNN autotune the best algorithm for the transposed convolution
torch.backends.cudnn.benchmark = True
torch.backends.cudnn.allow_tf32 = True

# ----------------------------------------------------------------------------
# Custom fused CUDA kernel:
#   out = scale * hardtanh(mish(x) + add_value, -1, 1)
# This replaces the 4 separate elementwise PyTorch kernels
# (mish, add, hardtanh, mul) with a single memory-bound kernel.
# Vectorized with float4 for maximum memory throughput.
# ----------------------------------------------------------------------------
fused_epilogue_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>

__device__ __forceinline__ float mish_f(float x) {
    // mish(x) = x * tanh(softplus(x)), softplus threshold = 20 (matches PyTorch)
    float sp = (x > 20.0f) ? x : log1pf(expf(x));
    return x * tanhf(sp);
}

__device__ __forceinline__ float epilogue_f(float x, float add_value, float scale) {
    float v = mish_f(x) + add_value;
    v = fminf(fmaxf(v, -1.0f), 1.0f);   // hardtanh(-1, 1)
    return v * scale;
}

__global__ void fused_epilogue_vec4_kernel(
    const float4* __restrict__ in,
    float4* __restrict__ out,
    const long n4,
    const float add_value,
    const float scale)
{
    long idx = (long)blockIdx.x * blockDim.x + threadIdx.x;
    const long stride = (long)gridDim.x * blockDim.x;
    for (; idx < n4; idx += stride) {
        float4 v = in[idx];
        float4 r;
        r.x = epilogue_f(v.x, add_value, scale);
        r.y = epilogue_f(v.y, add_value, scale);
        r.z = epilogue_f(v.z, add_value, scale);
        r.w = epilogue_f(v.w, add_value, scale);
        out[idx] = r;
    }
}

__global__ void fused_epilogue_scalar_kernel(
    const float* __restrict__ in,
    float* __restrict__ out,
    const long n,
    const float add_value,
    const float scale)
{
    long idx = (long)blockIdx.x * blockDim.x + threadIdx.x;
    const long stride = (long)gridDim.x * blockDim.x;
    for (; idx < n; idx += stride) {
        out[idx] = epilogue_f(in[idx], add_value, scale);
    }
}

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

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

    const float av = static_cast<float>(add_value);
    const float sc = static_cast<float>(scale);
    const int threads = 256;
    cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    if ((n & 3L) == 0) {
        const long n4 = n >> 2;
        long blocks = (n4 + threads - 1) / threads;
        fused_epilogue_vec4_kernel<<<(unsigned int)blocks, threads, 0, stream>>>(
            reinterpret_cast<const float4*>(xc.data_ptr<float>()),
            reinterpret_cast<float4*>(out.data_ptr<float>()),
            n4, av, sc);
    } else {
        long blocks = (n + threads - 1) / threads;
        fused_epilogue_scalar_kernel<<<(unsigned int)blocks, threads, 0, stream>>>(
            xc.data_ptr<float>(), out.data_ptr<float>(), n, av, sc);
    }
    return out;
}
"""

fused_epilogue_cpp_source = (
    "torch::Tensor fused_epilogue_cuda(torch::Tensor x, double add_value, double scale);"
)

# Compile the inline CUDA extension for the fused activation epilogue
fused_epilogue = load_inline(
    name="fused_mish_add_hardtanh_scale",
    cpp_sources=fused_epilogue_cpp_source,
    cuda_sources=fused_epilogue_cuda_source,
    functions=["fused_epilogue_cuda"],
    verbose=False,
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Optimized model:
      - ConvTranspose2d is delegated to cuDNN (with benchmarking / TF32 enabled),
        which is heavily optimized for this compute-bound GEMM-like operation.
      - The entire pointwise chain (Mish -> add -> Hardtanh -> scale) is replaced
        by a single fused custom CUDA kernel, eliminating 3 kernel launches and
        ~3x redundant global-memory traffic (float4-vectorized).
    """
    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)
        x = fused_epilogue.fused_epilogue_cuda(x, self.add_value, self.scale)
        return x
```