```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>
#include <vector>

std::vector<torch::Tensor> pack_phases_cuda(
    torch::Tensor x, torch::Tensor weight);

torch::Tensor pool_phases_cuda(
    torch::Tensor a, torch::Tensor b,
    torch::Tensor c, torch::Tensor d,
    torch::Tensor bias);
"""

_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 <cmath>
#include <cstdint>
#include <vector>

__global__ void pack_input_kernel(
    const float* __restrict__ x,
    float* __restrict__ a,
    float* __restrict__ b,
    float* __restrict__ c,
    float* __restrict__ d,
    int channels, int depth, int height, int width,
    int pd, int ph, int pw,
    int64_t sn, int64_t sc,
    int64_t sd, int64_t sh, int64_t sw
) {
    __shared__ float tile[32][33];

    const int tx = threadIdx.x;
    const int ty = threadIdx.y;
    const int n = blockIdx.z;
    const int first_c = blockIdx.y * 32;
    const int first_s = blockIdx.x * 32;
    const int volume = pd * ph * pw;
    const int s = first_s + tx;

    const int iw = s % pw - 1;
    const int ih = (s / pw) % ph - 1;
    const int id = s / (ph * pw) - 1;

    const bool valid = s < volume &&
        id >= 0 && id < depth &&
        ih >= 0 && ih < height &&
        iw >= 0 && iw < width;

    #pragma unroll
    for (int j = 0; j < 32; j += 8) {
        const int ic = first_c + ty + j;
        float value = 0.0f;
        if (valid && ic < channels) {
            value = x[int64_t(n) * sn + int64_t(ic) * sc +
                      int64_t(id) * sd + int64_t(ih) * sh +
                      int64_t(iw) * sw];
        }
        tile[ty + j][tx] = value;
    }
    __syncthreads();

    const int ic = first_c + tx;

    #pragma unroll
    for (int j = 0; j < 32; j += 8) {
        const int os = first_s + ty + j;
        if (os < volume && ic < channels) {
            const float value = tile[tx][ty + j];
            const int z = os / (ph * pw);
            const int y = (os / pw) % ph;
            const int w = os % pw;

            a[(int64_t(n) * volume + os) * channels + ic] = value;

            if (y > 0) {
                const int64_t target =
                    ((int64_t(n) * pd + z) * (ph - 1) + y - 1) * pw + w;
                b[target * channels + ic] = value;
            }
            if (z > 0) {
                const int64_t target =
                    ((int64_t(n) * (pd - 1) + z - 1) * ph + y) * pw + w;
                c[target * channels + ic] = value;
            }
            if (z > 0 && y > 0) {
                const int64_t target =
                    ((int64_t(n) * (pd - 1) + z - 1) * (ph - 1) +
                     y - 1) * pw + w;
                d[target * channels + ic] = value;
            }
        }
    }
}

__global__ void pack_weight_kernel(
    const float* __restrict__ weight,
    float* __restrict__ a,
    float* __restrict__ b,
    float* __restrict__ c,
    float* __restrict__ d,
    int channels,
    int64_t si, int64_t so,
    int64_t sd, int64_t sh, int64_t sw
) {
    const int phase = blockIdx.y;
    const int dp = phase >> 1;
    const int hp = phase & 1;
    const int kd_size = 3 - dp;
    const int kh_size = 3 - hp;
    const int64_t count = int64_t(128) * kd_size * kh_size * 3 * channels;

    float* target = phase == 0 ? a : (phase == 1 ? b : (phase == 2 ? c : d));

    for (int64_t index = int64_t(blockIdx.x) * blockDim.x + threadIdx.x;
         index < count;
         index += int64_t(blockDim.x) * gridDim.x) {
        int64_t t = index;
        const int ic = t % channels;
        t /= channels;
        const int rw = t % 3;
        t /= 3;
        const int rh = t % kh_size;
        t /= kh_size;
        const int rd = t % kd_size;
        t /= kd_size;
        const int oc = t & 63;
        const int wp = t >> 6;

        const int kd = 4 - dp - 2 * rd;
        const int kh = 4 - hp - 2 * rh;
        const int kw = 4 + wp - 2 * rw;

        float value = 0.0f;
        if (kw < 5) {
            value = weight[int64_t(ic) * si + int64_t(oc) * so +
                           int64_t(kd) * sd + int64_t(kh) * sh +
                           int64_t(kw) * sw];
        }
        target[index] = value;
    }
}

std::vector<torch::Tensor> pack_phases_cuda(
    torch::Tensor x, torch::Tensor weight
) {
    const c10::cuda::CUDAGuard guard(x.device());

    TORCH_CHECK(x.is_cuda() && weight.is_cuda(), "Expected CUDA tensors");
    TORCH_CHECK(x.scalar_type() == at::kFloat &&
                weight.scalar_type() == at::kFloat, "Expected float32");
    TORCH_CHECK(x.device() == weight.device(), "Device mismatch");

    const int n = x.size(0);
    const int ic = x.size(1);
    const int depth = x.size(2);
    const int height = x.size(3);
    const int width = x.size(4);

    const int pd = ((2 * depth - 1) / 6) * 3 + 2;
    const int ph = ((2 * height - 1) / 6) * 3 + 2;
    const int pw = ((2 * width - 1) / 6) * 3 + 2;

    std::vector<torch::Tensor> packed;
    packed.reserve(8);

    for (int p = 0; p < 4; ++p) {
        packed.push_back(
            torch::empty({n, pd - (p >> 1), ph - (p & 1), pw, ic},
                         x.options()).permute({0, 4, 1, 2, 3})
        );
    }
    for (int p = 0; p < 4; ++p) {
        packed.push_back(
            torch::empty({128, 3 - (p >> 1), 3 - (p & 1), 3, ic},
                         weight.options()).permute({0, 4, 1, 2, 3})
        );
    }

    const auto stream = at::cuda::getCurrentCUDAStream();

    pack_input_kernel<<<
        dim3((pd * ph * pw + 31) / 32, (ic + 31) / 32, n),
        dim3(32, 8), 0, stream>>>(
            x.data_ptr<float>(),
            packed[0].data_ptr<float>(),
            packed[1].data_ptr<float>(),
            packed[2].data_ptr<float>(),
            packed[3].data_ptr<float>(),
            ic, depth, height, width, pd, ph, pw,
            x.stride(0), x.stride(1), x.stride(2),
            x.stride(3), x.stride(4)
        );

    const int blocks = static_cast<int>(
        std::min<int64_t>((packed[4].numel() + 255) / 256, 65535)
    );

    pack_weight_kernel<<<dim3(blocks, 4), 256, 0, stream>>>(
        weight.data_ptr<float>(),
        packed[4].data_ptr<float>(),
        packed[5].data_ptr<float>(),
        packed[6].data_ptr<float>(),
        packed[7].data_ptr<float>(),
        ic,
        weight.stride(0), weight.stride(1), weight.stride(2),
        weight.stride(3), weight.stride(4)
    );

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return packed;
}

__device__ __forceinline__ float max_nan(float a, float b) {
    return isnan(a) ? a : (isnan(b) ? b : fmaxf(a, b));
}

__device__ __forceinline__ float4 max4(float4 a, float4 b) {
    return make_float4(
        max_nan(a.x, b.x), max_nan(a.y, b.y),
        max_nan(a.z, b.z), max_nan(a.w, b.w)
    );
}

__global__ void pool_phases_kernel(
    const float4* __restrict__ a,
    const float4* __restrict__ b,
    const float4* __restrict__ c,
    const float4* __restrict__ d,
    const float4* __restrict__ bias,
    float* __restrict__ output,
    int depth, int height, int width,
    int od, int oh, int ow
) {
    __shared__ float4 partial[4][16];

    const int tid = threadIdx.x;
    const int lane = tid & 31;
    const int warp = tid >> 5;
    const int index = blockIdx.x;
    const int output_volume = od * oh * ow;
    const int n = index / output_volume;

    int s = index % output_volume;
    const int w = s % ow;
    s /= ow;
    const int h = s % oh;
    const int z = s / oh;

    const int64_t origin =
        ((int64_t(n) * depth + 3 * z) * height + 3 * h) * width + 3 * w;

    const float4* phase = warp == 0 ? a : (warp == 1 ? b : (warp == 2 ? c : d));
    const float4* base = phase + origin * 32 + lane;

    float4 maximum = make_float4(
        -INFINITY, -INFINITY, -INFINITY, -INFINITY
    );

    #pragma unroll
    for (int dz = 0; dz < 3; ++dz) {
        #pragma unroll
        for (int dy = 0; dy < 3; ++dy) {
            #pragma unroll
            for (int dx = 0; dx < 3; ++dx) {
                const int64_t offset =
                    (int64_t(dz) * height * width + dy * width + dx) * 32;
                maximum = max4(maximum, base[offset]);
            }
        }
    }

    float4 other;
    other.x = __shfl_xor_sync(0xffffffff, maximum.x, 16);
    other.y = __shfl_xor_sync(0xffffffff, maximum.y, 16);
    other.z = __shfl_xor_sync(0xffffffff, maximum.z, 16);
    other.w = __shfl_xor_sync(0xffffffff, maximum.w, 16);
    maximum = max4(maximum, other);

    if (lane < 16) {
        partial[warp][lane] = maximum;
    }
    __syncthreads();

    if (warp == 0) {
        float sum = 0.0f;
        if (lane < 16) {
            float4 value = max4(
                max4(partial[0][lane], partial[1][lane]),
                max4(partial[2][lane], partial[3][lane])
            );
            const float4 bv = bias[lane];
            value.x += bv.x;
            value.y += bv.y;
            value.z += bv.z;
            value.w += bv.w;
            sum = (value.x + value.y) + (value.z + value.w);
        }

        #pragma unroll
        for (int offset = 8; offset > 0; offset >>= 1) {
            sum += __shfl_down_sync(0xffffffff, sum, offset);
        }
        if (lane == 0) {
            output[index] = sum;
        }
    }
}

torch::Tensor pool_phases_cuda(
    torch::Tensor a, torch::Tensor b,
    torch::Tensor c, torch::Tensor d,
    torch::Tensor bias
) {
    const c10::cuda::CUDAGuard guard(a.device());

    a = a.contiguous(at::MemoryFormat::ChannelsLast3d);
    b = b.contiguous(at::MemoryFormat::ChannelsLast3d);
    c = c.contiguous(at::MemoryFormat::ChannelsLast3d);
    d = d.contiguous(at::MemoryFormat::ChannelsLast3d);
    bias = bias.contiguous();

    const int depth = a.size(2);
    const int height = a.size(3);
    const int width = a.size(4);
    const int od = depth / 3;
    const int oh = height / 3;
    const int ow = width / 3;

    auto output = torch::empty({a.size(0), 1, od, oh, ow}, a.options());

    if (output.numel() > 0) {
        pool_phases_kernel<<<
            static_cast<unsigned int>(output.numel()), 128, 0,
            at::cuda::getCurrentCUDAStream()>>>(
                reinterpret_cast<const float4*>(a.data_ptr<float>()),
                reinterpret_cast<const float4*>(b.data_ptr<float>()),
                reinterpret_cast<const float4*>(c.data_ptr<float>()),
                reinterpret_cast<const float4*>(d.data_ptr<float>()),
                reinterpret_cast<const float4*>(bias.data_ptr<float>()),
                output.data_ptr<float>(),
                depth, height, width, od, oh, ow
            );
        C10_CUDA_KERNEL_LAUNCH_CHECK();
    }
    return output;
}
"""

_ops = load_inline(
    name="polyphase_conv3d_split_depth_height_v5",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["pack_phases_cuda", "pool_phases_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):
        super().__init__()
        self.conv_transpose = nn.ConvTranspose3d(
            in_channels,
            out_channels,
            kernel_size,
            stride=stride,
            padding=padding,
        )
        self.max_pool1 = nn.MaxPool3d(kernel_size=2)
        self.max_pool2 = nn.MaxPool3d(kernel_size=3)

    def forward(self, x):
        conv = self.conv_transpose
        use_polyphase = (
            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 x.ndim == 5
            and 0 < x.size(0) <= 65535
            and x.size(1) == conv.in_channels
            and min(x.shape[2:]) >= 4
            and conv.out_channels == 64
            and conv.kernel_size == (5, 5, 5)
            and conv.stride == (2, 2, 2)
            and conv.padding == (2, 2, 2)
            and conv.output_padding == (0, 0, 0)
            and conv.dilation == (1, 1, 1)
            and conv.groups == 1
            and not torch.is_autocast_enabled()
        )

        if not use_polyphase:
            y = conv(x)
            return self.max_pool2(self.max_pool1(y)).sum(
                dim=1, keepdim=True
            )

        packed = _ops.pack_phases_cuda(x, conv.weight)
        a = F.conv3d(packed[0], packed[4])
        b = F.conv3d(packed[1], packed[5])
        c = F.conv3d(packed[2], packed[6])
        d = F.conv3d(packed[3], packed[7])
        return _ops.pool_phases_cuda(a, b, c, d, conv.bias)
```