```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 to_channels_last_cuda(torch::Tensor x);
torch::Tensor fused_epilogue_cuda(
    torch::Tensor x, torch::Tensor conv_bias, torch::Tensor bias);
"""

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

__device__ __forceinline__ float fast_tanh(float x) {
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 750
    float result;
    asm("tanh.approx.f32 %0, %1;" : "=f"(result) : "f"(x));
    return result;
#else
    return tanhf(x);
#endif
}

__global__ void nchw_to_nhwc_64(
    const float* __restrict__ input,
    float* __restrict__ output,
    int64_t spatial)
{
    __shared__ float tile[64][65];

    const int t = threadIdx.x;
    const int lane = t & 63;
    const int row = t >> 6;
    const int64_t start = int64_t(blockIdx.x) * 64;
    const int64_t base = int64_t(blockIdx.y) * 64 * spatial;

    #pragma unroll
    for (int j = 0; j < 64; j += 4) {
        const int c = row + j;
        const int64_t p = start + lane;
        if (p < spatial)
            tile[c][lane] = input[base + int64_t(c) * spatial + p];
    }

    __syncthreads();

    #pragma unroll
    for (int j = 0; j < 64; j += 4) {
        const int p = row + j;
        if (start + p < spatial)
            output[base + (start + p) * 64 + lane] = tile[lane][p];
    }
}

torch::Tensor to_channels_last_cuda(torch::Tensor x) {
    TORCH_CHECK(x.is_cuda() && x.scalar_type() == at::kFloat,
                "Expected a CUDA float32 tensor");
    TORCH_CHECK(x.dim() == 4 && x.size(1) == 64 && x.is_contiguous(),
                "Expected contiguous NCHW input with 64 channels");

    const c10::cuda::CUDAGuard guard(x.device());
    auto output = torch::empty(
        x.sizes(), x.options().memory_format(at::MemoryFormat::ChannelsLast));

    if (x.numel() != 0) {
        const int64_t spatial = x.size(2) * x.size(3);
        const auto stream = c10::cuda::getCurrentCUDAStream(x.get_device());
        const dim3 grid(
            static_cast<unsigned int>((spatial + 63) / 64),
            static_cast<unsigned int>(x.size(0)));
        nchw_to_nhwc_64<<<grid, 256, 0, stream.stream()>>>(
            x.data_ptr<float>(), output.data_ptr<float>(), spatial);
        C10_CUDA_KERNEL_LAUNCH_CHECK();
    }
    return output;
}

__global__ void epilogue_64(
    float4* __restrict__ output,
    const float4* __restrict__ conv_bias,
    const float4* __restrict__ bias,
    int64_t count)
{
    const int c = threadIdx.x & 15;
    const float4 a = conv_bias[c];
    const float4 b = bias[c];
    const int64_t step = int64_t(blockDim.x) * gridDim.x;

    for (int64_t i = int64_t(blockIdx.x) * blockDim.x + threadIdx.x;
         i < count; i += step) {
        float4 v = output[i];
        v.x = fast_tanh((v.x + a.x) - b.x);
        v.y = fast_tanh((v.y + a.y) - b.y);
        v.z = fast_tanh((v.z + a.z) - b.z);
        v.w = fast_tanh((v.w + a.w) - b.w);
        output[i] = v;
    }
}

__global__ void epilogue_vectorized(
    float4* __restrict__ output,
    const float4* __restrict__ conv_bias,
    const float4* __restrict__ bias,
    int64_t count,
    int64_t channel_vectors)
{
    for (int64_t i = int64_t(blockIdx.x) * blockDim.x + threadIdx.x;
         i < count; i += int64_t(blockDim.x) * gridDim.x) {
        const int64_t c = i % channel_vectors;
        const float4 a = conv_bias[c];
        const float4 b = bias[c];
        float4 v = output[i];
        v.x = fast_tanh((v.x + a.x) - b.x);
        v.y = fast_tanh((v.y + a.y) - b.y);
        v.z = fast_tanh((v.z + a.z) - b.z);
        v.w = fast_tanh((v.w + a.w) - b.w);
        output[i] = v;
    }
}

__global__ void epilogue_generic(
    float* __restrict__ output,
    const float* __restrict__ conv_bias,
    const float* __restrict__ bias,
    int64_t count,
    int64_t channels,
    int64_t height,
    int64_t width,
    int64_t s0,
    int64_t s1,
    int64_t s2,
    int64_t s3)
{
    for (int64_t i = int64_t(blockIdx.x) * blockDim.x + threadIdx.x;
         i < count; i += int64_t(blockDim.x) * gridDim.x) {
        int64_t q = i;
        const int64_t c = q % channels;
        q /= channels;
        const int64_t w = q % width;
        q /= width;
        const int64_t h = q % height;
        const int64_t n = q / height;
        const float b = bias[n*s0 + c*s1 + h*s2 + w*s3];
        output[i] = fast_tanh((output[i] + conv_bias[c]) - b);
    }
}

torch::Tensor fused_epilogue_cuda(
    torch::Tensor x, torch::Tensor conv_bias, torch::Tensor bias)
{
    TORCH_CHECK(x.is_cuda() && conv_bias.is_cuda() && bias.is_cuda(),
                "Expected CUDA tensors");
    TORCH_CHECK(x.device() == conv_bias.device() &&
                x.device() == bias.device(), "Devices must match");
    TORCH_CHECK(x.scalar_type() == at::kFloat &&
                conv_bias.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 output");
    TORCH_CHECK(conv_bias.is_contiguous() &&
                conv_bias.numel() == x.size(1),
                "Invalid convolution bias");

    const c10::cuda::CUDAGuard guard(x.device());
    const auto expanded = bias.expand(x.sizes());
    const int64_t count = x.numel();
    if (!count)
        return x;

    const int64_t channels = x.size(1);
    const auto stream = c10::cuda::getCurrentCUDAStream(x.get_device());
    constexpr int threads = 256;

    const bool vectorized =
        channels % 4 == 0 &&
        expanded.stride(0) == 0 &&
        expanded.stride(1) == 1 &&
        expanded.stride(2) == 0 &&
        expanded.stride(3) == 0 &&
        reinterpret_cast<uintptr_t>(x.data_ptr<float>()) % 16 == 0 &&
        reinterpret_cast<uintptr_t>(bias.data_ptr<float>()) % 16 == 0 &&
        reinterpret_cast<uintptr_t>(conv_bias.data_ptr<float>()) % 16 == 0;

    if (vectorized) {
        const int64_t vectors = count / 4;
        const int blocks = static_cast<int>(
            std::min<int64_t>((vectors + threads - 1) / threads, 4096));

        auto* values = reinterpret_cast<float4*>(x.data_ptr<float>());
        const auto* a =
            reinterpret_cast<const float4*>(conv_bias.data_ptr<float>());
        const auto* b =
            reinterpret_cast<const float4*>(bias.data_ptr<float>());

        if (channels == 64) {
            epilogue_64<<<blocks, threads, 0, stream.stream()>>>(
                values, a, b, vectors);
        } else {
            epilogue_vectorized<<<blocks, threads, 0, stream.stream()>>>(
                values, a, b, vectors, channels / 4);
        }
    } else {
        const int blocks = static_cast<int>(
            std::min<int64_t>((count + threads - 1) / threads, 4096));
        epilogue_generic<<<blocks, threads, 0, stream.stream()>>>(
            x.data_ptr<float>(),
            conv_bias.data_ptr<float>(),
            bias.data_ptr<float>(),
            count, channels, x.size(2), x.size(3),
            expanded.stride(0), expanded.stride(1),
            expanded.stride(2), expanded.stride(3));
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return x;
}
"""

_fused = load_inline(
    name="transpose_conv_fast_epilogue_v3",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["to_channels_last_cuda", "fused_epilogue_cuda"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    verbose=False,
)


class _ChannelsLast(torch.autograd.Function):
    @staticmethod
    def forward(ctx, x):
        return _fused.to_channels_last_cuda(x)

    @staticmethod
    def backward(ctx, grad_output):
        return grad_output.contiguous()


class _FusedEpilogue(torch.autograd.Function):
    @staticmethod
    def forward(ctx, x, conv_bias, bias):
        ctx.bias_shape = bias.shape
        ctx.mark_dirty(x)
        _fused.fused_epilogue_cuda(x, conv_bias, bias)
        ctx.save_for_backward(x)
        return x

    @staticmethod
    def backward(ctx, grad_output):
        (output,) = ctx.saved_tensors
        grad = grad_output * (1 - output.square())
        grad_x = grad if ctx.needs_input_grad[0] else None
        grad_conv_bias = (
            grad.sum(dim=(0, 2, 3)) if ctx.needs_input_grad[1] else None
        )
        grad_bias = (
            -grad.sum_to_size(ctx.bias_shape) if ctx.needs_input_grad[2] else None
        )
        return grad_x, grad_conv_bias, grad_bias


class ModelNew(nn.Module):
    def __init__(
        self,
        in_channels,
        out_channels,
        kernel_size,
        bias_shape,
        stride=2,
        padding=1,
        output_padding=1,
    ):
        super().__init__()
        self.conv_transpose = nn.ConvTranspose2d(
            in_channels,
            out_channels,
            kernel_size,
            stride=stride,
            padding=padding,
            output_padding=output_padding,
        )
        self.bias = nn.Parameter(torch.randn(bias_shape))
        self.conv_transpose.to(memory_format=torch.channels_last)

    def forward(self, x):
        conv = self.conv_transpose
        if (
            not x.is_cuda
            or x.dtype != torch.float32
            or conv.weight.dtype != torch.float32
            or conv.bias.dtype != torch.float32
            or self.bias.dtype != torch.float32
            or torch.is_autocast_enabled()
            or x.ndim != 4
        ):
            return torch.tanh(conv(x) - self.bias)

        if not x.is_contiguous(memory_format=torch.channels_last):
            if x.is_contiguous() and x.shape[1] == 64 and x.shape[0] <= 65535:
                if torch.is_grad_enabled() and x.requires_grad:
                    x = _ChannelsLast.apply(x)
                else:
                    x = _fused.to_channels_last_cuda(x)
            else:
                x = x.contiguous(memory_format=torch.channels_last)

        output = F.conv_transpose2d(
            x,
            conv.weight.contiguous(memory_format=torch.channels_last),
            bias=None,
            stride=conv.stride,
            padding=conv.padding,
            output_padding=conv.output_padding,
            groups=conv.groups,
            dilation=conv.dilation,
        )
        output = output.contiguous(memory_format=torch.channels_last)

        if torch.is_grad_enabled() and (
            output.requires_grad
            or conv.bias.requires_grad
            or self.bias.requires_grad
        ):
            return _FusedEpilogue.apply(output, conv.bias, self.bias)

        return _fused.fused_epilogue_cuda(output, conv.bias, self.bias)
```