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

torch::Tensor pool_sum_cuda(
    torch::Tensor even, torch::Tensor odd,
    torch::Tensor bias, int64_t depth);
"""

_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>

// Fold the depth stencil into the channel dimension. Both convolutions
// then use native channels-last 2D convolution layouts.
__global__ void pack_input_kernel(
    const float* __restrict__ x,
    float* __restrict__ even,
    float* __restrict__ odd,
    int channels, int depth, int height, int width,
    int output_depth, 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 batch_depth = blockIdx.z;
    const int n = batch_depth / output_depth;
    const int z = batch_depth % output_depth;
    const int first_c = blockIdx.y * 32;
    const int first_s = blockIdx.x * 32;
    const int spatial = ph * pw;
    const int s = first_s + tx;
    const int ih = s / pw - 1;
    const int iw = s % pw - 1;

    const bool valid_hw =
        s < spatial && ih >= 0 && ih < height &&
        iw >= 0 && iw < width;

    #pragma unroll
    for (int j = 0; j < 32; j += 8) {
        const int expanded_c = first_c + ty + j;
        float value = 0.0f;
        if (expanded_c < 5 * channels && valid_hw) {
            const int stencil = expanded_c / channels;
            const int ic = expanded_c % channels;
            const int id = z + (stencil < 3 ? stencil - 1 : stencil - 3);
            if (id >= 0 && id < depth) {
                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 expanded_c = first_c + tx;

    #pragma unroll
    for (int j = 0; j < 32; j += 8) {
        const int os = first_s + ty + j;
        if (os < spatial && expanded_c < 5 * channels) {
            const int64_t row = int64_t(batch_depth) * spatial + os;
            const float value = tile[tx][ty + j];

            if (expanded_c < 3 * channels) {
                even[row * (3 * channels) + expanded_c] = value;
            } else {
                odd[row * (2 * channels) + expanded_c - 3 * channels] = value;
            }
        }
    }
}

__global__ void pack_weight_kernel(
    const float* __restrict__ weight,
    float* __restrict__ even,
    float* __restrict__ odd,
    int channels,
    int64_t even_count, int64_t total_count,
    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 % channels;
        t /= channels;
        const int rd = t % (is_odd ? 2 : 3);
        t /= is_odd ? 2 : 3;
        const int rw = t % 3;
        t /= 3;
        const int rh = t % 3;
        t /= 3;
        const int oc = t % 64;
        const int phase = t / 64;

        const int kd = (is_odd ? 3 : 4) - 2 * rd;
        const int kh = 4 + (phase >> 1) - 2 * rh;
        const int kw = 4 + (phase & 1) - 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_conv2d_cuda(
    torch::Tensor x, torch::Tensor weight
) {
    TORCH_CHECK(x.is_cuda() && weight.is_cuda(), "Expected CUDA tensors");
    const c10::cuda::CUDAGuard guard(x.device());

    TORCH_CHECK(x.scalar_type() == at::kFloat &&
                weight.scalar_type() == at::kFloat, "Expected float32");
    TORCH_CHECK(x.device() == weight.device(), "Device mismatch");
    TORCH_CHECK(x.dim() == 5 && weight.dim() == 5, "Expected 5D tensors");
    TORCH_CHECK(weight.size(0) == x.size(1) &&
                weight.size(1) == 64, "Channel 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 od = ((2 * depth - 1) / 6) * 3;
    const int ph = ((2 * height - 1) / 6) * 3 + 2;
    const int pw = ((2 * width - 1) / 6) * 3 + 2;

    auto even_x = torch::empty(
        {int64_t(n) * od, ph, pw, 3 * ic}, x.options()
    ).permute({0, 3, 1, 2});

    auto odd_x = torch::empty(
        {int64_t(n) * od, ph, pw, 2 * ic}, x.options()
    ).permute({0, 3, 1, 2});

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

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

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

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

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

    pack_weight_kernel<<<blocks, 256, 0, stream>>>(
        weight.data_ptr<float>(),
        even_weight.data_ptr<float>(), odd_weight.data_ptr<float>(),
        ic, even_count, total_count,
        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)
    );
}

// The two non-overlapping pooling layers combine into a 6x6x6
// maximum, or a 3x3x3 maximum over all eight convolution phases.
__global__ void 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 volume = od * oh * ow;
    const int n = index / volume;

    int s = index % volume;
    const int x = s % ow;
    s /= ow;
    const int y = s % oh;
    const int z = s / oh;

    const int64_t origin =
        ((int64_t(n) * depth + 3 * z) * height + 3 * y) * width + 3 * x;

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

    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) * 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 (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 b = bias[lane];
            value.x += b.x;
            value.y += b.y;
            value.z += b.z;
            value.w += b.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_sum_cuda(
    torch::Tensor even, torch::Tensor odd,
    torch::Tensor bias, int64_t depth
) {
    const c10::cuda::CUDAGuard guard(even.device());
    TORCH_CHECK(even.is_cuda() && odd.is_cuda() && bias.is_cuda(),
                "Expected CUDA tensors");
    TORCH_CHECK(even.device() == odd.device() &&
                even.device() == bias.device(), "Device mismatch");
    TORCH_CHECK(even.scalar_type() == at::kFloat &&
                odd.scalar_type() == at::kFloat &&
                bias.scalar_type() == at::kFloat, "Expected float32");
    TORCH_CHECK(even.sizes() == odd.sizes(), "Phase shape mismatch");
    TORCH_CHECK(even.dim() == 4 && even.size(1) == 256 &&
                bias.numel() == 64, "Invalid phase channels");
    TORCH_CHECK(depth > 0 && even.size(0) % depth == 0,
                "Invalid folded depth");

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

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

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

    if (output.numel() > 0) {
        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>(),
                static_cast<int>(depth), h, w, od, oh, ow
            );
        C10_CUDA_KERNEL_LAUNCH_CHECK();
    }

    return output;
}
"""

_ops = load_inline(
    name="polyphase_depth_folded_conv2d_v6",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["pack_conv2d_cuda", "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_custom = (
            x.is_cuda
            and x.ndim == 5
            and x.dtype == torch.float32
            and conv.weight.dtype == torch.float32
            and conv.weight.device == x.device
            and conv.bias is not None
            and conv.bias.dtype == torch.float32
            and x.size(1) == conv.in_channels
            and x.size(0) > 0
            and min(x.shape[2:]) >= 4
            and x.size(0) * (3 * ((2 * x.size(2) - 1) // 6)) <= 65535
            and (5 * x.size(1) + 31) // 32 <= 65535
            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_custom:
            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_conv2d_cuda(
            x, conv.weight
        )
        even = F.conv2d(even_x, even_weight)
        odd = F.conv2d(odd_x, odd_weight)
        folded_depth = 3 * ((2 * x.size(2) - 1) // 6)
        return _ops.pool_sum_cuda(even, odd, conv.bias, folded_depth)
```