```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 fused_transpose_cuda(
    torch::Tensor x, torch::Tensor weight, torch::Tensor bias,
    double add_value, double scale);

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 <mma.h>
#include <algorithm>
#include <cstdint>

namespace wmma = nvcuda::wmma;

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

    const float x = __fadd_rn(value, bias);
    const float e = __expf(fminf(x, 20.0f));
    const float numerator = e * (e + 2.0f);
    const float gate = __fdividef(numerator, numerator + 2.0f);
    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 __fmul_rn(y, scale);
}

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

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

// Each phase of the stride-two convolution is a dense matrix product.
// A block computes 32 spatial positions and all 64 output channels.
// The activation is evaluated directly on the tensor-core accumulators,
// avoiding a full-size intermediate convolution output.
template<int PY, int PX>
__global__ __launch_bounds__(256, 2) void transpose_phase_kernel(
    const float* __restrict__ input,
    const float* __restrict__ weight,
    const float4* __restrict__ bias,
    float4* __restrict__ output,
    int height,
    int width,
    float add_value,
    float scale) {

#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 800
    __shared__ __align__(32) float shared[3072];

    float* const a_tile = shared;
    float* const b_tile = shared + 1024;

    const int tid = threadIdx.x;
    const int warp = tid >> 5;
    const int warp_m = warp >> 2;
    const int warp_n = warp & 3;

    const int batch = blockIdx.z;
    const int row = blockIdx.y;
    const int column = blockIdx.x * 32;
    const int64_t spatial = int64_t(height) * width;

    wmma::fragment<wmma::accumulator, 16, 16, 8, float> accum;
    wmma::fill_fragment(accum, 0.0f);

    #pragma unroll
    for (int dy = 0; dy <= PY; ++dy) {
        #pragma unroll
        for (int dx = 0; dx <= PX; ++dx) {
            constexpr int kernel_size = 3;
            const int kh = 1 + PY - 2 * dy;
            const int kw = 1 + PX - 2 * dx;

            #pragma unroll
            for (int channel_base = 0; channel_base < 64;
                 channel_base += 32) {

                // Column-major A: adjacent lanes read adjacent NCHW pixels.
                #pragma unroll
                for (int j = 0; j < 4; ++j) {
                    const int index = tid + j * 256;
                    const int m = index & 31;
                    const int k = index >> 5;
                    const int input_row = row + dy;
                    const int input_column = column + m + dx;

                    float value = 0.0f;
                    if (input_row < height && input_column < width) {
                        const int64_t offset =
                            (int64_t(batch) * 64 + channel_base + k) *
                                spatial +
                            int64_t(input_row) * width + input_column;
                        value = __ldg(input + offset);
                    }
                    a_tile[index] = __float_to_tf32(value);
                }

                // Row-major B, with contiguous output channels.
                #pragma unroll
                for (int j = 0; j < 8; ++j) {
                    const int index = tid + j * 256;
                    const int k = index >> 6;
                    const int oc = index & 63;
                    const int offset =
                        ((channel_base + k) * kernel_size * kernel_size +
                         kh * kernel_size + kw) * 64 + oc;
                    b_tile[index] =
                        __float_to_tf32(__ldg(weight + offset));
                }

                __syncthreads();

                #pragma unroll
                for (int k = 0; k < 32; k += 8) {
                    wmma::fragment<
                        wmma::matrix_a, 16, 16, 8,
                        wmma::precision::tf32, wmma::col_major> a;

                    wmma::fragment<
                        wmma::matrix_b, 16, 16, 8,
                        wmma::precision::tf32, wmma::row_major> b;

                    wmma::load_matrix_sync(
                        a, a_tile + k * 32 + warp_m * 16, 32);
                    wmma::load_matrix_sync(
                        b, b_tile + k * 64 + warp_n * 16, 64);
                    wmma::mma_sync(accum, a, b, accum);
                }

                __syncthreads();
            }
        }
    }

    // Reuse the operand storage for the completed output tile.
    wmma::store_matrix_sync(
        shared + warp_m * 16 * 64 + warp_n * 16,
        accum, 64, wmma::mem_row_major);

    __syncthreads();

    const int channel4 = tid & 15;
    const float4 b = bias[channel4];
    const float4* const results =
        reinterpret_cast<const float4*>(shared);
    const int output_width = width * 2;
    const int output_height = height * 2;

    #pragma unroll
    for (int j = 0; j < 2; ++j) {
        const int vector_index = tid + j * 256;
        const int m = vector_index >> 4;
        const int output_row = row * 2 + PY;
        const int output_column = (column + m) * 2 + PX;

        const int64_t offset =
            ((int64_t(batch) * output_height + output_row) *
                 output_width + output_column) * 16 + channel4;

        output[offset] =
            activate4(results[vector_index], b, add_value, scale);
    }
#endif
}

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

    TORCH_CHECK(x.is_cuda() && weight.is_cuda() && bias.is_cuda(),
                "Expected CUDA tensors");
    TORCH_CHECK(x.device() == weight.device() &&
                x.device() == bias.device(), "Device mismatch");
    TORCH_CHECK(x.scalar_type() == at::kFloat &&
                weight.scalar_type() == at::kFloat &&
                bias.scalar_type() == at::kFloat,
                "Expected float32 tensors");
    TORCH_CHECK(x.dim() == 4 && x.size(1) == 64 && x.is_contiguous(),
                "Expected contiguous NCHW input with 64 channels");
    TORCH_CHECK(weight.is_contiguous() && weight.dim() == 4 &&
                weight.size(0) == 64 && weight.size(1) == 3 &&
                weight.size(2) == 3 && weight.size(3) == 64,
                "Expected contiguous [64, 3, 3, 64] weights");
    TORCH_CHECK(bias.is_contiguous() && bias.numel() == 64,
                "Expected 64 bias elements");
    TORCH_CHECK(x.size(3) % 32 == 0 &&
                x.size(0) <= 65535 && x.size(2) <= 65535,
                "Unsupported input dimensions");

    const c10::cuda::CUDAGuard guard(x.device());
    TORCH_CHECK(at::cuda::getCurrentDeviceProperties()->major >= 8,
                "Tensor-core kernel requires compute capability >= 8.0");

    const int64_t batch = x.size(0);
    const int height = static_cast<int>(x.size(2));
    const int width = static_cast<int>(x.size(3));

    auto output = torch::empty(
        {batch, int64_t(height) * 2, int64_t(width) * 2, 64},
        x.options()).permute({0, 3, 1, 2});

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

    const dim3 grid(width / 32, height, static_cast<unsigned int>(batch));
    const auto stream = at::cuda::getCurrentCUDAStream();

    const float* input_ptr = x.data_ptr<float>();
    const float* weight_ptr = weight.data_ptr<float>();
    const float4* bias_ptr =
        reinterpret_cast<const float4*>(bias.data_ptr<float>());
    float4* output_ptr =
        reinterpret_cast<float4*>(output.data_ptr<float>());

    const float add = static_cast<float>(add_value);
    const float mul = static_cast<float>(scale);

    transpose_phase_kernel<0, 0><<<grid, 256, 0, stream>>>(
        input_ptr, weight_ptr, bias_ptr, output_ptr,
        height, width, add, mul);
    transpose_phase_kernel<0, 1><<<grid, 256, 0, stream>>>(
        input_ptr, weight_ptr, bias_ptr, output_ptr,
        height, width, add, mul);
    transpose_phase_kernel<1, 0><<<grid, 256, 0, stream>>>(
        input_ptr, weight_ptr, bias_ptr, output_ptr,
        height, width, add, mul);
    transpose_phase_kernel<1, 1><<<grid, 256, 0, stream>>>(
        input_ptr, weight_ptr, bias_ptr, output_ptr,
        height, width, add, mul);

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}

template<bool Channels64>
__global__ void activation_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 step = int64_t(gridDim.x) * blockDim.x;

    float4 fixed_bias;
    if (Channels64) {
        fixed_bias = bias[threadIdx.x & 15];
    }

    for (int64_t index = first; index < count; index += step) {
        const float4 b =
            Channels64 ? fixed_bias : bias[index % channels4];
        data[index] = activate4(data[index], b, add_value, scale);
    }
}

__global__ void activation_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 step = int64_t(gridDim.x) * blockDim.x;

    for (int64_t index = first; index < count; index += step) {
        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() &&
                x.device() == bias.device(), "Expected matching CUDA tensors");
    TORCH_CHECK(x.scalar_type() == at::kFloat &&
                bias.scalar_type() == at::kFloat, "Expected float32 tensors");
    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 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;

    constexpr int threads = 256;
    const int64_t work = vectorized ? count / 4 : count;
    const int blocks = static_cast<int>(
        std::min<int64_t>((work + threads - 1) / threads, 4096));
    const auto stream = at::cuda::getCurrentCUDAStream();
    const float add = static_cast<float>(add_value);
    const float mul = static_cast<float>(scale);

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

        if (channels == 64) {
            activation_vector_kernel<true><<<blocks, threads, 0, stream>>>(
                data, b, work, channels / 4, add, mul);
        } else {
            activation_vector_kernel<false><<<blocks, threads, 0, stream>>>(
                data, b, work, channels / 4, add, mul);
        }
    } else {
        activation_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="fused_polyphase_transpose_mish_v5",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["fused_transpose_cuda", "fused_activation_cuda"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3", "--fmad=false"],
    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
        use_cuda = (
            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()
        )

        if use_cuda:
            use_fused_convolution = (
                conv.in_channels == 64
                and conv.out_channels == 64
                and conv.kernel_size == (3, 3)
                and conv.stride == (2, 2)
                and conv.padding == (1, 1)
                and conv.output_padding == (1, 1)
                and conv.dilation == (1, 1)
                and conv.groups == 1
                and x.ndim == 4
                and x.is_contiguous()
                and 0 < x.shape[0] <= 65535
                and 0 < x.shape[2] <= 65535
                and x.shape[3] > 0
                and x.shape[3] % 32 == 0
                and torch.backends.cudnn.allow_tf32
                and torch.cuda.get_device_capability(x.device)[0] >= 8
            )

            if use_fused_convolution:
                weight = conv.weight.permute(0, 2, 3, 1).contiguous()
                return _ops.fused_transpose_cuda(
                    x,
                    weight,
                    conv.bias,
                    float(self.add_value),
                    float(self.scale),
                )

            x = x.contiguous(memory_format=torch.channels_last)
            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
```