```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 input_channels_last(torch::Tensor x);
torch::Tensor pack_polyphase_weights(torch::Tensor weight);
torch::Tensor polyphase_epilogue(
    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>

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

__device__ __forceinline__ float4 activate4(
    float4 v, const float4 a, const float4 b)
{
    v.x = activate((v.x + a.x) - b.x);
    v.y = activate((v.y + a.y) - b.y);
    v.z = activate((v.z + a.z) - b.z);
    v.w = activate((v.w + a.w) - b.w);
    return v;
}

__device__ __forceinline__ void store_streaming(float4* p, float4 v) {
    asm volatile(
        "st.global.cs.v4.f32 [%0], {%1, %2, %3, %4};"
        :
        : "l"(p), "f"(v.x), "f"(v.y), "f"(v.z), "f"(v.w)
        : "memory");
}

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

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

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

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

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

    if (x.numel()) {
        const int spatial = x.size(2) * x.size(3);
        const auto stream = c10::cuda::getCurrentCUDAStream(x.get_device());

        transpose_input<<<
            dim3((spatial + 63) / 64, x.size(0)),
            256, 0, stream.stream()
        >>>(x.data_ptr<float>(), output.data_ptr<float>(), spatial);

        C10_CUDA_KERNEL_LAUNCH_CHECK();
    }
    return output;
}

__global__ void pack_weights_kernel(
    const float* __restrict__ weight,
    float* __restrict__ packed,
    int64_t s0, int64_t s1, int64_t s2, int64_t s3)
{
    const int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= 64 * 64 * 16)
        return;

    const int ic = i & 63;
    const int kw = (i >> 6) & 1;
    const int kh = (i >> 7) & 1;
    const int oc = (i >> 8) & 63;
    const int phase = i >> 14;

    const int original_h = 3 - 2 * kh - (phase >> 1);
    const int original_w = 3 - 2 * kw - (phase & 1);

    packed[i] = weight[
        ic * s0 + oc * s1 + original_h * s2 + original_w * s3];
}

torch::Tensor pack_polyphase_weights(torch::Tensor weight) {
    const c10::cuda::CUDAGuard guard(weight.device());
    auto packed = torch::empty(
        {256, 64, 2, 2},
        weight.options().memory_format(at::MemoryFormat::ChannelsLast));

    const auto stream = c10::cuda::getCurrentCUDAStream(weight.get_device());

    pack_weights_kernel<<<256, 256, 0, stream.stream()>>>(
        weight.data_ptr<float>(), packed.data_ptr<float>(),
        weight.stride(0), weight.stride(1),
        weight.stride(2), weight.stride(3));

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return packed;
}

// Process complete pairs of output columns. The first column is handled
// by the first tile, avoiding a nearly empty tile at the right edge.
template<int FIXED_LOW_H, int FIXED_LOW_W>
__global__ __launch_bounds__(256, 4) void shuffle_activate_kernel(
    const float4* __restrict__ input,
    float4* __restrict__ output,
    const float4* __restrict__ conv_bias,
    const float4* __restrict__ bias,
    int dynamic_low_h, int dynamic_low_w)
{
    const int low_h = FIXED_LOW_H ? FIXED_LOW_H : dynamic_low_h;
    const int low_w = FIXED_LOW_W ? FIXED_LOW_W : dynamic_low_w;
    const int height = 2 * low_h - 1;
    const int width = 2 * low_w - 1;

    const int channel = threadIdx.x & 15;
    const int phase = (threadIdx.x >> 4) & 3;
    const int ph = phase >> 1;
    const int pw = phase & 1;
    const int pixel_lane = threadIdx.x >> 6;

    const int ih = blockIdx.y;
    const int n = blockIdx.z;
    const int oh = 2 * ih - ph;

    if (oh < 0)
        return;

    const float4 a = conv_bias[channel];
    const float4 b = bias[channel];

    const int input_row =
        ((n * low_h + ih) * low_w) * 64 + phase * 16 + channel;
    const int output_row =
        ((n * height + oh) * width) * 16 + channel;

    const int first_col = int(blockIdx.x) * 32 + pixel_lane + 1;

    #pragma unroll
    for (int j = 0; j < 32; j += 4) {
        const int iw = first_col + j;
        if (FIXED_LOW_W == 257 || iw < low_w) {
            float4 v = input[input_row + iw * 64];
            v = activate4(v, a, b);
            store_streaming(
                output + output_row + (2 * iw - pw) * 16, v);
        }
    }

    if (blockIdx.x == 0 && pixel_lane == 0 && pw == 0) {
        float4 v = input[input_row];
        store_streaming(output + output_row, activate4(v, a, b));
    }
}

torch::Tensor polyphase_epilogue(
    torch::Tensor x, torch::Tensor conv_bias, torch::Tensor bias)
{
    const c10::cuda::CUDAGuard guard(x.device());

    const int low_h = x.size(2);
    const int low_w = x.size(3);
    const int height = 2 * low_h - 1;
    const int width = 2 * low_w - 1;

    auto output = torch::empty(
        {x.size(0), 64, height, width},
        x.options().memory_format(at::MemoryFormat::ChannelsLast));

    if (!output.numel())
        return output;

    const auto stream = c10::cuda::getCurrentCUDAStream(x.get_device());
    const dim3 grid((low_w + 30) / 32, low_h, x.size(0));

    const auto* input_ptr =
        reinterpret_cast<const float4*>(x.data_ptr<float>());
    auto* output_ptr =
        reinterpret_cast<float4*>(output.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 (low_h == 257 && low_w == 257) {
        shuffle_activate_kernel<257, 257>
            <<<grid, 256, 0, stream.stream()>>>(
                input_ptr, output_ptr, a, b, low_h, low_w);
    } else {
        shuffle_activate_kernel<0, 0>
            <<<grid, 256, 0, stream.stream()>>>(
                input_ptr, output_ptr, a, b, low_h, low_w);
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_ops = load_inline(
    name="polyphase_transpose_conv_v6",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=[
        "input_channels_last",
        "pack_polyphase_weights",
        "polyphase_epilogue",
    ],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    verbose=False,
)


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

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


class _PackWeights(torch.autograd.Function):
    @staticmethod
    def forward(ctx, weight):
        return _ops.pack_polyphase_weights(weight)

    @staticmethod
    def backward(ctx, grad):
        result = grad.new_empty((64, 64, 4, 4))
        for ph in range(2):
            for pw in range(2):
                phase = 2 * ph + pw
                values = grad[phase * 64:(phase + 1) * 64]
                result[:, :, 1 - ph::2, 1 - pw::2] = (
                    values.permute(1, 0, 2, 3).flip((2, 3))
                )
        return result


class _ShuffleActivate(torch.autograd.Function):
    @staticmethod
    def forward(ctx, x, conv_bias, bias):
        output = _ops.polyphase_epilogue(x, conv_bias, bias)
        ctx.input_shape = x.shape
        ctx.bias_shape = bias.shape
        ctx.save_for_backward(output)
        return output

    @staticmethod
    def backward(ctx, grad_output):
        (output,) = ctx.saved_tensors
        grad = grad_output * (1.0 - output.square())

        grad_x = None
        if ctx.needs_input_grad[0]:
            grad_x = torch.empty(
                ctx.input_shape,
                device=grad.device,
                dtype=grad.dtype,
                memory_format=torch.channels_last,
            )
            grad_x.zero_()
            for ph in range(2):
                for pw in range(2):
                    values = grad[:, :, ph::2, pw::2]
                    phase = 2 * ph + pw
                    grad_x[
                        :,
                        phase * 64:(phase + 1) * 64,
                        ph:ph + values.shape[2],
                        pw:pw + values.shape[3],
                    ] = values

        channel_grad = (
            grad.sum(dim=(0, 2, 3))
            if ctx.needs_input_grad[1] or ctx.needs_input_grad[2]
            else None
        )
        grad_conv_bias = channel_grad if ctx.needs_input_grad[1] else None
        grad_bias = (
            -channel_grad.reshape(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))

    def forward(self, x):
        conv = self.conv_transpose
        fast_path = (
            x.is_cuda
            and x.dtype == torch.float32
            and x.ndim == 4
            and 0 < x.shape[0] <= 65535
            and x.shape[1] == 64
            and 0 < x.shape[2] < 65535
            and x.shape[3] > 0
            and conv.in_channels == 64
            and conv.out_channels == 64
            and conv.kernel_size == (4, 4)
            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 conv.weight.dtype == torch.float32
            and conv.bias is not None
            and conv.bias.dtype == torch.float32
            and self.bias.dtype == torch.float32
            and conv.weight.device == x.device
            and conv.bias.device == x.device
            and self.bias.device == x.device
            and self.bias.shape == (64, 1, 1)
            and self.bias.is_contiguous()
            and conv.bias.is_contiguous()
            and self.bias.data_ptr() % 16 == 0
            and conv.bias.data_ptr() % 16 == 0
            and not torch.is_autocast_enabled()
            and (
                x.shape[0]
                * (x.shape[2] + 1)
                * (x.shape[3] + 1)
                * 64
                < 2**31 - 1048576
            )
        )

        if not fast_path:
            return torch.tanh(conv(x) - self.bias)

        if not x.is_contiguous(memory_format=torch.channels_last):
            if x.is_contiguous():
                x = _ChannelsLast.apply(x)
            else:
                x = x.contiguous(memory_format=torch.channels_last)

        weight = _PackWeights.apply(conv.weight)
        phases = F.conv2d(x, weight, bias=None, stride=1, padding=1)
        phases = phases.contiguous(memory_format=torch.channels_last)
        return _ShuffleActivate.apply(phases, conv.bias, self.bias)
```