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

_fused_epilogue_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>
#include <algorithm>
#include <cstdint>

namespace {
__device__ __forceinline__ float softplus_mish(float x) {
    return x > 20.0f ? x : log1pf(expf(x));
}

__device__ __forceinline__ double softplus_mish(double x) {
    return x > 20.0 ? x : log1p(exp(x));
}

__device__ __forceinline__ float tanh_value(float x) {
    return tanhf(x);
}

__device__ __forceinline__ double tanh_value(double x) {
    return tanh(x);
}

template <typename scalar_t>
__device__ __forceinline__ scalar_t fused_epilogue_value(
    scalar_t v,
    scalar_t add_value,
    scalar_t scale) {
    scalar_t sp = softplus_mish(v);
    scalar_t h = v * tanh_value(sp) + add_value;
    h = h < static_cast<scalar_t>(-1) ? static_cast<scalar_t>(-1)
        : (h > static_cast<scalar_t>(1) ? static_cast<scalar_t>(1) : h);
    return h * scale;
}

__global__ void fused_epilogue_float4_kernel(
    const float4* __restrict__ x,
    float4* __restrict__ out,
    int64_t n4,
    float add_value,
    float scale) {
    int64_t idx = static_cast<int64_t>(blockIdx.x) * blockDim.x + threadIdx.x;
    const int64_t stride = static_cast<int64_t>(gridDim.x) * blockDim.x;

    for (; idx < n4; idx += stride) {
        float4 v = x[idx];
        float4 r;
        r.x = fused_epilogue_value<float>(v.x, add_value, scale);
        r.y = fused_epilogue_value<float>(v.y, add_value, scale);
        r.z = fused_epilogue_value<float>(v.z, add_value, scale);
        r.w = fused_epilogue_value<float>(v.w, add_value, scale);
        out[idx] = r;
    }
}

template <typename scalar_t>
__global__ void fused_epilogue_kernel(
    const scalar_t* __restrict__ x,
    scalar_t* __restrict__ out,
    int64_t n,
    scalar_t add_value,
    scalar_t scale) {
    int64_t idx = static_cast<int64_t>(blockIdx.x) * blockDim.x + threadIdx.x;
    const int64_t stride = static_cast<int64_t>(gridDim.x) * blockDim.x;

    for (; idx < n; idx += stride) {
        out[idx] = fused_epilogue_value<scalar_t>(x[idx], add_value, scale);
    }
}
}  // namespace

torch::Tensor fused_epilogue_cuda(torch::Tensor x, double add_value, double scale) {
    TORCH_CHECK(x.is_cuda(), "fused_epilogue_cuda expects a CUDA tensor");

    auto x_contig = x.contiguous();
    auto out = torch::empty_like(x_contig);
    const int64_t n = x_contig.numel();
    if (n == 0) {
        return out;
    }

    constexpr int threads = 256;
    auto stream = at::cuda::getCurrentCUDAStream();

    if (x_contig.scalar_type() == at::kFloat && ((n & 3) == 0)) {
        const float* x_ptr = x_contig.data_ptr<float>();
        float* out_ptr = out.data_ptr<float>();
        const bool aligned =
            ((reinterpret_cast<std::uintptr_t>(x_ptr) & 15) == 0) &&
            ((reinterpret_cast<std::uintptr_t>(out_ptr) & 15) == 0);

        if (aligned) {
            const int64_t n4 = n >> 2;
            const int64_t blocks64 = (n4 + threads - 1) / threads;
            const unsigned int blocks =
                static_cast<unsigned int>(std::min<int64_t>(blocks64, 65535));
            fused_epilogue_float4_kernel<<<blocks, threads, 0, stream>>>(
                reinterpret_cast<const float4*>(x_ptr),
                reinterpret_cast<float4*>(out_ptr),
                n4,
                static_cast<float>(add_value),
                static_cast<float>(scale));
            return out;
        }
    }

    const int64_t blocks64 = (n + threads - 1) / threads;
    const unsigned int blocks =
        static_cast<unsigned int>(std::min<int64_t>(blocks64, 65535));

    AT_DISPATCH_FLOATING_TYPES(x_contig.scalar_type(), "fused_epilogue_cuda", [&] {
        fused_epilogue_kernel<scalar_t><<<blocks, threads, 0, stream>>>(
            x_contig.data_ptr<scalar_t>(),
            out.data_ptr<scalar_t>(),
            n,
            static_cast<scalar_t>(add_value),
            static_cast<scalar_t>(scale));
    });

    return out;
}
"""

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

_fused_epilogue = load_inline(
    name="fused_mish_add_hardtanh_scale_epilogue",
    cpp_sources=_fused_epilogue_cpp_source,
    cuda_sources=_fused_epilogue_cuda_source,
    functions=["fused_epilogue_cuda"],
    verbose=False,
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Model that performs a transposed convolution, applies Mish activation, adds a value,
    applies Hardtanh activation, and scales the output.
    """
    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 in (torch.float32, torch.float64):
            return _fused_epilogue.fused_epilogue_cuda(x, self.add_value, self.scale)

        x = F.mish(x)
        x = x + self.add_value
        x = F.hardtanh(x, min_val=-1, max_val=1)
        x = x * self.scale
        return x
```