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

# Enable cuDNN autotuning for the fixed-shape transposed convolution
torch.backends.cudnn.benchmark = True

# Fused CUDA kernel implementing:
#   out = scale * hardtanh(mish(x) + add_value, min_val, max_val)
# This fuses 4 element-wise passes over the (large) conv output into a
# single memory-bound pass, eliminating ~6 GB of intermediate global
# memory traffic for this workload.
fused_act_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>

__device__ __forceinline__ float mish_add_hardtanh_scale_op(
    float v, float add_value, float min_val, float max_val, float scale) {
    // mish(x) = x * tanh(softplus(x)), softplus(x) = log1p(exp(x))
    // numerically stable: for v > 20, softplus(v) == v and tanh(v) == 1 in fp32
    float sp = (v > 20.0f) ? v : log1pf(expf(v));
    float m = v * tanhf(sp);
    float z = m + add_value;
    z = fminf(fmaxf(z, min_val), max_val);   // hardtanh clamp
    return z * scale;                        // scale
}

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

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

torch::Tensor fused_mish_add_hardtanh_scale(
    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 = (long)xc.numel();
    if (n == 0) {
        return out;
    }

    cudaStream_t stream = at::cuda::getCurrentCUDAStream();
    const int block = 256;
    const long max_blocks = 1 << 20;  // plenty; grid-stride loop covers the rest

    if (n % 4 == 0) {
        long n4 = n >> 2;
        long blocks = (n4 + block - 1) / block;
        if (blocks > max_blocks) blocks = max_blocks;
        fused_act_kernel_vec4<<<(unsigned int)blocks, block, 0, stream>>>(
            reinterpret_cast<const float4*>(xc.data_ptr<float>()),
            reinterpret_cast<float4*>(out.data_ptr<float>()),
            n4,
            (float)add_value, (float)min_val, (float)max_val, (float)scale);
    } else {
        long blocks = (n + block - 1) / block;
        if (blocks > max_blocks) blocks = max_blocks;
        fused_act_kernel_scalar<<<(unsigned int)blocks, block, 0, stream>>>(
            xc.data_ptr<float>(),
            out.data_ptr<float>(),
            n,
            (float)add_value, (float)min_val, (float)max_val, (float)scale);
    }
    return out;
}
"""

fused_act_cpp_source = (
    "torch::Tensor fused_mish_add_hardtanh_scale(torch::Tensor x, double add_value, "
    "double min_val, double max_val, double scale);"
)

fused_act = load_inline(
    name="fused_mish_add_hardtanh_scale",
    cpp_sources=fused_act_cpp_source,
    cuda_sources=fused_act_cuda_source,
    functions=["fused_mish_add_hardtanh_scale"],
    verbose=False,
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Optimized model: cuDNN transposed convolution followed by a single fused
    CUDA kernel that applies Mish, scalar add, Hardtanh, and scaling in one
    memory-bound pass.
    """
    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_act.fused_mish_add_hardtanh_scale(x, self.add_value, -1.0, 1.0, self.scale)
        return x
```