```
# 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
from torch.utils.cpp_extension import load_inline

# Fused CUDA kernel implementing: Mish -> Add -> Hardtanh(-1, 1) -> Scale
# This fuses 4 element-wise passes over the large ConvTranspose2d output into
# a single memory-bound pass, and uses float4 vectorized loads/stores.
fused_act_source = """
#include <torch/extension.h>
#include <cuda_runtime.h>

__device__ __forceinline__ float fused_act_op(float v, float add_value, float scale) {
    // Mish: v * tanh(softplus(v)), where softplus(v) = log(1 + exp(v)).
    // Numerically stable: for v > 20, softplus(v) ~ v (and avoids exp overflow).
    float sp = (v > 20.0f) ? v : log1pf(expf(v));
    float y = v * tanhf(sp);
    // Add constant
    y += add_value;
    // Hardtanh clamp to [-1, 1]
    y = fminf(fmaxf(y, -1.0f), 1.0f);
    // Scale
    y *= scale;
    return y;
}

__global__ void fused_act_vec4_kernel(const float4* __restrict__ x,
                                      float4* __restrict__ out,
                                      long n4,
                                      float add_value,
                                      float scale) {
    long idx = blockIdx.x * (long)blockDim.x + threadIdx.x;
    if (idx < n4) {
        float4 v = x[idx];
        float4 r;
        r.x = fused_act_op(v.x, add_value, scale);
        r.y = fused_act_op(v.y, add_value, scale);
        r.z = fused_act_op(v.z, add_value, scale);
        r.w = fused_act_op(v.w, add_value, scale);
        out[idx] = r;
    }
}

__global__ void fused_act_tail_kernel(const float* __restrict__ x,
                                      float* __restrict__ out,
                                      int n,
                                      float add_value,
                                      float scale) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        out[idx] = fused_act_op(x[idx], add_value, scale);
    }
}

torch::Tensor fused_act_cuda(torch::Tensor x, double add_value, 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;
    }

    float av = static_cast<float>(add_value);
    float sc = static_cast<float>(scale);
    const int block_size = 256;

    const float* x_ptr = xc.data_ptr<float>();
    float* out_ptr = out.data_ptr<float>();

    long n4 = n / 4;
    if (n4 > 0) {
        long num_blocks = (n4 + block_size - 1) / block_size;
        fused_act_vec4_kernel<<<num_blocks, block_size>>>(
            reinterpret_cast<const float4*>(x_ptr),
            reinterpret_cast<float4*>(out_ptr),
            n4, av, sc);
    }

    int rem = static_cast<int>(n % 4);
    if (rem > 0) {
        fused_act_tail_kernel<<<1, 32>>>(x_ptr + n4 * 4, out_ptr + n4 * 4, rem, av, sc);
    }

    return out;
}
"""

fused_act_cpp_source = (
    "torch::Tensor fused_act_cuda(torch::Tensor x, double add_value, double scale);"
)

# Compile the inline fused CUDA operator
fused_act = load_inline(
    name="fused_mish_add_hardtanh_scale",
    cpp_sources=fused_act_cpp_source,
    cuda_sources=fused_act_source,
    functions=["fused_act_cuda"],
    verbose=True,
    extra_cflags=[""],
    extra_ldflags=[""],
)


class ModelNew(nn.Module):
    """
    Optimized model: ConvTranspose2d (cuDNN) followed by a single fused custom
    CUDA kernel that applies Mish, constant addition, Hardtanh(-1, 1), and
    scaling in one pass over the data.
    """
    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)
        self.fused_act = fused_act

    def forward(self, x):
        x = self.conv_transpose(x)
        x = self.fused_act.fused_act_cuda(x, self.add_value, self.scale)
        return x
```