```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_polyphase_cuda(
    torch::Tensor x, torch::Tensor weight);

torch::Tensor polyphase_pool_sum_cuda(
    torch::Tensor even, torch::Tensor odd, 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__ even,
    float* __restrict__ odd,
    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 / (pw * ph) - 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 c = first_c + ty + j;
        float value = 0.0f;
        if (valid && c < channels) {
            value = x[int64_t(n) * sn + int64_t(c) * sc +
                      int64_t(id) * sd + int64_t(ih) * sh +
                      int64_t(iw) * sw];
        }
        tile[ty + j][tx] = value;
    }

    __syncthreads();

    const int c = first_c + tx;
    #pragma unroll
    for (int j = 0; j < 32; j += 8) {
        const int os = first_s + ty + j;
        if (os < volume && c < channels) {
            const float value = tile[tx][ty + j];
            even[(int64_t(n) * volume + os) * channels + c] = value;

            if (os >= ph * pw) {
                const int odd_volume = (pd - 1) * ph * pw;
                odd[(int64_t(n) * odd_volume + os - ph * pw) *
                    channels + c] = value;
            }
        }
    }
}

__global__ void pack_weight_kernel(
    const float* __restrict__ weight,
    float* __restrict__ even,
    float* __restrict__ odd,
    int64_t even_count,
    int64_t total_count,
    int in_channels,
    int64_t si, int64_t so,
    int64_t sd, int64_t sh, int64_t sw
) {
    for (int64_t index = int64_t(blockIdx.x) * blockDim.x + threadIdx.x;
         index < total_count;
         index += int64_t(blockDim.x) * gridDim.x) {
        const bool is_odd = index >= even_count;
        const int64_t local = is_odd ? index - even_count : index;
        int64_t t = local;

        const int ic = t % in_channels;
        t /= in_channels;
        const int rw = t % 3;
        t /= 3;
        const int rh = t % 3;
        t /= 3;
        const int kd_size = is_odd ? 2 : 3;
        const int rd = t % kd_size;
        t /= kd_size;
        const int oc = t % 64;
        const int phase = t / 64;

        const int kd = (is_odd ? 3 : 4) - 2 * rd;
        const int kh = ((phase >> 1) & 1) + 4 - 2 * rh;
        const int kw = (phase & 1) + 4 - 2 * rw;

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

        if (is_odd) {
            odd[local] = value;
        } else {
            even[local] = value;
        }
    }
}

std::vector<torch::Tensor> pack_polyphase_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 tensors");
    TORCH_CHECK(x.device() == weight.device(), "Device mismatch");
    TORCH_CHECK(weight.size(1) == 64, "Expected 64 output channels");

    const int n = x.size(0);
    const int ic = x.size(1);
    const int d = x.size(2);
    const int h = x.size(3);
    const int w = x.size(4);

    const int pd = ((2 * d - 1) / 6) * 3 + 2;
    const int ph = ((2 * h - 1) / 6) * 3 + 2;
    const int pw = ((2 * w - 1) / 6) * 3 + 2;

    auto even_x = torch::empty({n, pd, ph, pw, ic}, x.options())
                      .permute({0, 4, 1, 2, 3});
    auto odd_x = torch::empty({n, pd - 1, ph, pw, ic}, x.options())
                     .permute({0, 4, 1, 2, 3});

    auto even_weight = torch::empty({256, 3, 3, 3, ic}, weight.options())
                           .permute({0, 4, 1, 2, 3});
    auto odd_weight = torch::empty({256, 2, 3, 3, ic}, weight.options())
                          .permute({0, 4, 1, 2, 3});

    const auto stream = at::cuda::getCurrentCUDAStream();
    const dim3 threads(32, 8);
    const dim3 blocks(
        (pd * ph * pw + 31) / 32,
        (ic + 31) / 32,
        n
    );

    pack_input_kernel<<<blocks, threads, 0, stream>>>(
        x.data_ptr<float>(),
        even_x.data_ptr<float>(), odd_x.data_ptr<float>(),
        ic, d, h, w, pd, ph, pw,
        x.stride(0), x.stride(1), x.stride(2),
        x.stride(3), x.stride(4)
    );

    const int64_t total_count = even_weight.numel() + odd_weight.numel();
    const int weight_blocks = static_cast<int>(
        std::min<int64_t>((total_count + 255) / 256, 65535)
    );

    pack_weight_kernel<<<weight_blocks, 256, 0, stream>>>(
        weight.data_ptr<float>(),
        even_weight.data_ptr<float>(), odd_weight.data_ptr<float>(),
        even_weight.numel(), total_count, ic,
        weight.stride(0), weight.stride(1), weight.stride(2),
        weight.stride(3), weight.stride(4)
    );

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return {even_x, odd_x, even_weight, odd_weight};
}

__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 polyphase_pool_sum_kernel(
    const float4* __restrict__ even,
    const float4* __restrict__ odd,
    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 d = s / oh;

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

    const float4* base =
        (tid < 64 ? even : odd) + origin * 64 + (tid & 63);

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

    #pragma unroll
    for (int z = 0; z < 3; ++z) {
        #pragma unroll
        for (int y = 0; y < 3; ++y) {
            #pragma unroll
            for (int x = 0; x < 3; ++x) {
                const int64_t offset =
                    (int64_t(z) * height * width + y * width + x) * 64;
                maximum = max4(maximum, base[offset]);
            }
        }
    }

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

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

    if (tid < 32) {
        float sum = 0.0f;

        if (tid < 16) {
            float4 pooled = max4(
                max4(partial[0][tid], partial[1][tid]),
                max4(partial[2][tid], partial[3][tid])
            );
            const float4 b = bias[tid];

            pooled.x += b.x;
            pooled.y += b.y;
            pooled.z += b.z;
            pooled.w += b.w;
            sum = (pooled.x + pooled.y) + (pooled.z + pooled.w);
        }

        #pragma unroll
        for (int offset = 8; offset > 0; offset >>= 1) {
            sum += __shfl_down_sync(0xffffffff, sum, offset);
        }

        if (tid == 0) {
            output[index] = sum;
        }
    }
}

torch::Tensor polyphase_pool_sum_cuda(
    torch::Tensor even, torch::Tensor odd, torch::Tensor bias
) {
    const c10::cuda::CUDAGuard guard(even.device());

    TORCH_CHECK(even.scalar_type() == at::kFloat &&
                odd.scalar_type() == at::kFloat &&
                bias.scalar_type() == at::kFloat,
                "Expected float32 tensors");
    TORCH_CHECK(even.device() == odd.device() &&
                even.device() == bias.device(), "Device mismatch");
    TORCH_CHECK(even.sizes() == odd.sizes(), "Phase shape mismatch");

    even = even.contiguous(at::MemoryFormat::ChannelsLast3d);
    odd = odd.contiguous(at::MemoryFormat::ChannelsLast3d);
    bias = bias.contiguous();

    const int d = even.size(2);
    const int h = even.size(3);
    const int w = even.size(4);
    const int od = d / 3;
    const int oh = h / 3;
    const int ow = w / 3;

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

    if (output.numel() != 0) {
        polyphase_pool_sum_kernel
            <<<static_cast<unsigned int>(output.numel()), 128, 0,
               at::cuda::getCurrentCUDAStream()>>>(
                reinterpret_cast<const float4*>(even.data_ptr<float>()),
                reinterpret_cast<const float4*>(odd.data_ptr<float>()),
                reinterpret_cast<const float4*>(bias.data_ptr<float>()),
                output.data_ptr<float>(),
                d, h, w, od, oh, ow
            );
        C10_CUDA_KERNEL_LAUNCH_CHECK();
    }

    return output;
}
"""

_ops = load_inline(
    name="polyphase_conv3d_split_depth_v4",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["pack_polyphase_cuda", "polyphase_pool_sum_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 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
            )

        even_x, odd_x, even_weight, odd_weight = _ops.pack_polyphase_cuda(
            x, conv.weight
        )
        even = F.conv3d(even_x, even_weight)
        odd = F.conv3d(odd_x, odd_weight)
        return _ops.polyphase_pool_sum_cuda(even, odd, conv.bias)
```