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


_cpp_source = r"""
#include <torch/extension.h>

torch::Tensor channels_last_cuda(torch::Tensor x);
torch::Tensor fused_activation_cuda(
    torch::Tensor x, torch::Tensor bias, double add_value, double scale);
"""

_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <c10/cuda/CUDAGuard.h>
#include <c10/cuda/CUDAException.h>
#include <cuda_runtime.h>
#include <algorithm>
#include <cstdint>

__global__ void transpose_layout_kernel(
    const float* __restrict__ input,
    float* __restrict__ output,
    int channels,
    int64_t spatial) {

    __shared__ float tile[32][33];

    const int tx = threadIdx.x;
    const int ty = threadIdx.y;
    const int channel_start = blockIdx.y * 32;
    const int64_t spatial_start = int64_t(blockIdx.x) * 32;
    const int64_t batch_offset =
        int64_t(blockIdx.z) * channels * spatial;

    #pragma unroll
    for (int j = 0; j < 32; j += 8) {
        const int channel = channel_start + ty + j;
        const int64_t position = spatial_start + tx;
        if (channel < channels && position < spatial) {
            tile[ty + j][tx] =
                input[batch_offset + int64_t(channel) * spatial + position];
        }
    }

    __syncthreads();

    #pragma unroll
    for (int j = 0; j < 32; j += 8) {
        const int channel = channel_start + tx;
        const int64_t position = spatial_start + ty + j;
        if (channel < channels && position < spatial) {
            output[batch_offset + position * channels + channel] =
                tile[tx][ty + j];
        }
    }
}

torch::Tensor channels_last_cuda(torch::Tensor x) {
    TORCH_CHECK(x.is_cuda() && x.scalar_type() == at::kFloat,
                "Expected CUDA float32 input");
    TORCH_CHECK(x.dim() == 4, "Expected four-dimensional input");

    if (x.is_contiguous(at::MemoryFormat::ChannelsLast)) {
        return x;
    }

    const c10::cuda::CUDAGuard guard(x.device());

    if (!x.is_contiguous() ||
        x.size(0) > 65535 ||
        x.size(1) > 65535LL * 32) {
        return x.contiguous(at::MemoryFormat::ChannelsLast);
    }

    const int64_t n = x.size(0);
    const int64_t c = x.size(1);
    const int64_t h = x.size(2);
    const int64_t w = x.size(3);
    const int64_t spatial = h * w;

    auto output = torch::empty({n, h, w, c}, x.options())
                      .permute({0, 3, 1, 2});

    if (x.numel() == 0) {
        return output;
    }

    const dim3 block(32, 8);
    const dim3 grid(
        static_cast<unsigned int>((spatial + 31) / 32),
        static_cast<unsigned int>((c + 31) / 32),
        static_cast<unsigned int>(n));

    transpose_layout_kernel<<<
        grid, block, 0, at::cuda::getCurrentCUDAStream()>>>(
        x.data_ptr<float>(), output.data_ptr<float>(),
        static_cast<int>(c), spatial);

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}

__device__ __forceinline__ float4 streaming_load(const float4* address) {
    float4 value;
    asm volatile(
        "ld.global.cs.v4.f32 {%0, %1, %2, %3}, [%4];"
        : "=f"(value.x), "=f"(value.y), "=f"(value.z), "=f"(value.w)
        : "l"(address)
        : "memory");
    return value;
}

__device__ __forceinline__ void streaming_store(
    float4* address, const float4 value) {

    asm volatile(
        "st.global.cs.v4.f32 [%0], {%1, %2, %3, %4};"
        :
        : "l"(address), "f"(value.x), "f"(value.y),
          "f"(value.z), "f"(value.w)
        : "memory");
}

__device__ __forceinline__ float activate(
    float value, float bias, float add_value, float scale) {

    const float x = value + bias;

    // tanh(softplus(x)) = e^x(e^x + 2) / (e^x(e^x + 2) + 2).
    // Capping the exponential also handles the positive softplus branch
    // without overflow or a divergent branch.
    const float e = __expf(fminf(x, 20.0f));
    const float numerator = e * (e + 2.0f);
    const float gate = __fdividef(numerator, numerator + 2.0f);

    // Preserve the separate float32 rounding of Mish and scalar addition.
    float y = __fadd_rn(__fmul_rn(x, gate), add_value);
    y = y < -1.0f ? -1.0f : y;
    y = y > 1.0f ? 1.0f : y;
    return y * scale;
}

__device__ __forceinline__ float4 activate_vector(
    float4 value, float4 bias, float add_value, float scale) {

    value.x = activate(value.x, bias.x, add_value, scale);
    value.y = activate(value.y, bias.y, add_value, scale);
    value.z = activate(value.z, bias.z, add_value, scale);
    value.w = activate(value.w, bias.w, add_value, scale);
    return value;
}

template<int Threads>
__global__ __launch_bounds__(Threads) void fused_64_kernel(
    float4* __restrict__ data,
    const float4* __restrict__ bias,
    int64_t count,
    float add_value,
    float scale) {

    constexpr int Items = 4;
    constexpr int Tile = Threads * Items;

    const int lane = threadIdx.x;
    const float4 b = bias[lane & 15];
    const int64_t stride = int64_t(gridDim.x) * Tile;

    for (int64_t base = int64_t(blockIdx.x) * Tile;
         base < count;
         base += stride) {

        if (base + Tile <= count) {
            // Issue independent vector loads before evaluating the activation.
            float4 v0 = streaming_load(data + base + lane);
            float4 v1 = streaming_load(data + base + lane + Threads);
            float4 v2 = streaming_load(data + base + lane + 2 * Threads);
            float4 v3 = streaming_load(data + base + lane + 3 * Threads);

            v0 = activate_vector(v0, b, add_value, scale);
            v1 = activate_vector(v1, b, add_value, scale);
            v2 = activate_vector(v2, b, add_value, scale);
            v3 = activate_vector(v3, b, add_value, scale);

            streaming_store(data + base + lane, v0);
            streaming_store(data + base + lane + Threads, v1);
            streaming_store(data + base + lane + 2 * Threads, v2);
            streaming_store(data + base + lane + 3 * Threads, v3);
        } else {
            #pragma unroll
            for (int j = 0; j < Items; ++j) {
                const int64_t index = base + lane + j * Threads;
                if (index < count) {
                    float4 value = streaming_load(data + index);
                    value = activate_vector(value, b, add_value, scale);
                    streaming_store(data + index, value);
                }
            }
        }
    }
}

__global__ void fused_vector_kernel(
    float4* __restrict__ data,
    const float4* __restrict__ bias,
    int64_t count,
    int channels4,
    float add_value,
    float scale) {

    const int64_t first =
        int64_t(blockIdx.x) * blockDim.x + threadIdx.x;
    const int64_t stride = int64_t(gridDim.x) * blockDim.x;

    for (int64_t index = first; index < count; index += stride) {
        const float4 b = bias[index % channels4];
        float4 value = streaming_load(data + index);
        value = activate_vector(value, b, add_value, scale);
        streaming_store(data + index, value);
    }
}

__global__ void fused_scalar_kernel(
    float* __restrict__ data,
    const float* __restrict__ bias,
    int64_t count,
    int channels,
    float add_value,
    float scale) {

    const int64_t first =
        int64_t(blockIdx.x) * blockDim.x + threadIdx.x;
    const int64_t stride = int64_t(gridDim.x) * blockDim.x;

    for (int64_t index = first; index < count; index += stride) {
        data[index] = activate(
            data[index], bias[index % channels], add_value, scale);
    }
}

torch::Tensor fused_activation_cuda(
    torch::Tensor x,
    torch::Tensor bias,
    double add_value,
    double scale) {

    TORCH_CHECK(x.is_cuda() && bias.is_cuda(), "Expected CUDA tensors");
    TORCH_CHECK(x.scalar_type() == at::kFloat &&
                bias.scalar_type() == at::kFloat,
                "Expected float32 tensors");
    TORCH_CHECK(x.device() == bias.device(), "Device mismatch");
    TORCH_CHECK(x.dim() == 4 &&
                x.is_contiguous(at::MemoryFormat::ChannelsLast),
                "Expected channels-last input");
    TORCH_CHECK(bias.is_contiguous() && bias.numel() == x.size(1),
                "Invalid bias");

    const c10::cuda::CUDAGuard guard(x.device());
    const int64_t count = x.numel();

    if (count == 0) {
        return x;
    }

    const int channels = static_cast<int>(x.size(1));
    const float add = static_cast<float>(add_value);
    const float mul = static_cast<float>(scale);
    const auto stream = at::cuda::getCurrentCUDAStream();

    const bool vectorized =
        channels % 4 == 0 &&
        reinterpret_cast<uintptr_t>(x.data_ptr<float>()) % 16 == 0 &&
        reinterpret_cast<uintptr_t>(bias.data_ptr<float>()) % 16 == 0;

    if (vectorized) {
        auto* data = reinterpret_cast<float4*>(x.data_ptr<float>());
        const auto* b =
            reinterpret_cast<const float4*>(bias.data_ptr<float>());
        const int64_t vectors = count / 4;

        if (channels == 64) {
            constexpr int threads = 128;
            constexpr int tile = threads * 4;
            const int blocks = static_cast<int>(
                std::min<int64_t>((vectors + tile - 1) / tile, 4096));

            fused_64_kernel<threads><<<blocks, threads, 0, stream>>>(
                data, b, vectors, add, mul);
        } else {
            constexpr int threads = 256;
            const int blocks = static_cast<int>(
                std::min<int64_t>((vectors + threads - 1) / threads, 4096));

            fused_vector_kernel<<<blocks, threads, 0, stream>>>(
                data, b, vectors, channels / 4, add, mul);
        }
    } else {
        constexpr int threads = 256;
        const int blocks = static_cast<int>(
            std::min<int64_t>((count + threads - 1) / threads, 4096));

        fused_scalar_kernel<<<blocks, threads, 0, stream>>>(
            x.data_ptr<float>(), bias.data_ptr<float>(),
            count, channels, add, mul);
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return x;
}
"""

_ops = load_inline(
    name="channels_last_mish_hardtanh_streaming_v4",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["channels_last_cuda", "fused_activation_cuda"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    verbose=False,
)


class ModelNew(nn.Module):
    def __init__(
        self,
        in_channels,
        out_channels,
        kernel_size,
        stride,
        padding,
        output_padding,
        add_value,
        scale,
    ):
        super().__init__()
        self.conv_transpose = nn.ConvTranspose2d(
            in_channels,
            out_channels,
            kernel_size,
            stride,
            padding,
            output_padding,
        )
        self.conv_transpose.to(memory_format=torch.channels_last)
        self.add_value = add_value
        self.scale = scale

    def forward(self, x):
        conv = self.conv_transpose

        if (
            x.is_cuda
            and x.dtype == torch.float32
            and conv.weight.dtype == torch.float32
            and conv.bias is not None
            and conv.bias.dtype == torch.float32
            and not torch.is_autocast_enabled()
        ):
            x = _ops.channels_last_cuda(x)
            x = F.conv_transpose2d(
                x,
                conv.weight,
                bias=None,
                stride=conv.stride,
                padding=conv.padding,
                output_padding=conv.output_padding,
                groups=conv.groups,
                dilation=conv.dilation,
            )
            x = x.contiguous(memory_format=torch.channels_last)
            return _ops.fused_activation_cuda(
                x, conv.bias, float(self.add_value), float(self.scale)
            )

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