Looking at the previous solution, I found the bug: in `fused_forward_cuda`, `const int OC = wpc.size(2);` reads the wrong dimension. The packed weight `wp` has shape `[IC, kD, kH, kW, OC]`, so `size(2)` is `5` (kernel height), not `64` (output channels). This made the conv kernel compute only 5 output channels with garbage indexing, and the pool kernel then read wildly out of bounds. The fix is to take `OC` from the bias (`bc.size(0)`).

I also rewrote the pooling kernel so warps read coalesced memory (one warp per channel, lanes stride the 6×6×6 window) instead of the previous fully-uncoalesced channel-per-thread mapping.

```python
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline

_cuda_src = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <c10/cuda/CUDAException.h>
#include <c10/cuda/CUDAGuard.h>
#include <cuda_runtime.h>
#include <float.h>

#define TILE 8
#define OC_TILE 8
#define IC_CHUNK 4
#define K 5
#define NTAP 125
#define NTHREADS 256

// ---------------------------------------------------------------------------
// Custom ConvTranspose3d (gather formulation), stride=2, kernel=5, pad=2.
// out[b, oc, D, H, W] = bias[oc]
//     + sum_{ic,kd,kh,kw} X[b,ic,(D+pad-kd)/2,(H+pad-kh)/2,(W+pad-kw)/2]
//                         * W[ic,oc,kd,kh,kw]
// Only parity-matched, in-range taps are evaluated. Weight is pre-packed to
// layout [IC][125][OC] (oc contiguous) for coalesced / vectorized access.
// ---------------------------------------------------------------------------
__global__ void __launch_bounds__(NTHREADS) convt3d_s2k5_kernel(
    const float* __restrict__ X,
    const float* __restrict__ Wp,
    const float* __restrict__ bias,
    float* __restrict__ Y,
    int IC, int Di, int Hi, int Wi,
    int OC, int Do, int Ho, int Wo, int pad)
{
    __shared__ __align__(16) float sx[IC_CHUNK * 216];
    __shared__ __align__(16) float sw[NTAP * IC_CHUNK * OC_TILE];

    const int tid = threadIdx.x;                 // 0..255
    const int tx = tid & 7;
    const int ty = (tid >> 3) & 7;
    const int tz = tid >> 6;                     // 0..3

    const int ntw = (Wo + TILE - 1) / TILE;
    const int nth = (Ho + TILE - 1) / TILE;
    const int tw = blockIdx.x % ntw;
    const int th = (blockIdx.x / ntw) % nth;
    const int td = blockIdx.x / (ntw * nth);

    const int W0 = tw * TILE, H0 = th * TILE, D0 = td * TILE;
    const int oc0 = blockIdx.y * OC_TILE;
    const int b   = blockIdx.z;

    const int W = W0 + tx, H = H0 + ty;
    const int Da = D0 + tz;          // voxel 0 depth
    const int Db = D0 + tz + 4;      // voxel 1 depth

    // input region required by this output tile (at most 6x6x6)
    const int d_lo = max(0, (D0 + pad - (K - 1)) / 2);
    const int d_hi = min(Di - 1, (D0 + TILE - 1 + pad) / 2);
    const int h_lo = max(0, (H0 + pad - (K - 1)) / 2);
    const int h_hi = min(Hi - 1, (H0 + TILE - 1 + pad) / 2);
    const int w_lo = max(0, (W0 + pad - (K - 1)) / 2);
    const int w_hi = min(Wi - 1, (W0 + TILE - 1 + pad) / 2);
    const int region_d = d_hi - d_lo + 1;
    const int region_h = h_hi - h_lo + 1;
    const int region_w = w_hi - w_lo + 1;

    // per-thread local tap indices (-1 = invalid). Validity does NOT depend on
    // whether this thread's output voxel is in range.
    int dloc_a[K], dloc_b[K], hloc[K], wloc[K];
    #pragma unroll
    for (int kd = 0; kd < K; ++kd) {
        int t = Da + pad - kd;
        int d = t >> 1;
        dloc_a[kd] = (t >= 0 && !(t & 1) && d >= d_lo && d <= d_hi) ? (d - d_lo) : -1;
        t = Db + pad - kd;
        d = t >> 1;
        dloc_b[kd] = (t >= 0 && !(t & 1) && d >= d_lo && d <= d_hi) ? (d - d_lo) : -1;
        t = H + pad - kd;
        d = t >> 1;
        hloc[kd] = (t >= 0 && !(t & 1) && d >= h_lo && d <= h_hi) ? (d - h_lo) : -1;
        t = W + pad - kd;
        d = t >> 1;
        wloc[kd] = (t >= 0 && !(t & 1) && d >= w_lo && d <= w_hi) ? (d - w_lo) : -1;
    }

    const bool hw_ok    = (H < Ho) && (W < Wo);
    const bool valid_a  = hw_ok && (Da < Do);
    const bool valid_b  = hw_ok && (Db < Do);

    float acc_a[OC_TILE], acc_b[OC_TILE];
    #pragma unroll
    for (int o = 0; o < OC_TILE; ++o) { acc_a[o] = 0.f; acc_b[o] = 0.f; }

    for (int ic0 = 0; ic0 < IC; ic0 += IC_CHUNK) {
        // stage input tile into shared memory (zeros outside logical region)
        #pragma unroll
        for (int it = 0; it < 4; ++it) {
            const int idx = tid + it * NTHREADS;
            if (idx < IC_CHUNK * 216) {
                const int ic = idx / 216;
                const int r  = idx % 216;
                const int dz = r / 36;
                const int dy = (r / 6) % 6;
                const int dx = r % 6;
                float v = 0.f;
                if (dz < region_d && dy < region_h && dx < region_w && (ic0 + ic) < IC) {
                    const long xb = ((long)(b * IC + ic0 + ic) * Di + (d_lo + dz)) * Hi * Wi;
                    v = X[xb + (h_lo + dy) * Wi + (w_lo + dx)];
                }
                sx[idx] = v;
            }
        }
        // stage packed weights into shared memory
        #pragma unroll
        for (int it = 0; it < 16; ++it) {
            const int idx = tid + it * NTHREADS;
            if (idx < NTAP * IC_CHUNK * OC_TILE) {
                const int oc  = idx & (OC_TILE - 1);
                const int ic  = (idx / OC_TILE) & (IC_CHUNK - 1);
                const int tap = idx / (OC_TILE * IC_CHUNK);
                float v = 0.f;
                if ((ic0 + ic) < IC && (oc0 + oc) < OC) {
                    v = Wp[((long)(ic0 + ic) * NTAP + tap) * OC + (oc0 + oc)];
                }
                sw[idx] = v;
            }
        }
        __syncthreads();

        #pragma unroll
        for (int kd = 0; kd < K; ++kd) {
            const int dla = dloc_a[kd];
            const int dlb = dloc_b[kd];
            if (dla < 0 && dlb < 0) continue;
            #pragma unroll
            for (int kh = 0; kh < K; ++kh) {
                const int hl = hloc[kh];
                if (hl < 0) continue;
                #pragma unroll
                for (int kw = 0; kw < K; ++kw) {
                    const int wl = wloc[kw];
                    if (wl < 0) continue;
                    const int tapid = kd * 25 + kh * 5 + kw;
                    #pragma unroll
                    for (int ic = 0; ic < IC_CHUNK; ++ic) {
                        const float4* wp4 = reinterpret_cast<const float4*>(&sw[(tapid * IC_CHUNK + ic) * OC_TILE]);
                        const float4 w0 = wp4[0];
                        const float4 w1 = wp4[1];
                        if (valid_a && dla >= 0) {
                            const float xv = sx[ic * 216 + (dla * 6 + hl) * 6 + wl];
                            acc_a[0] = fmaf(xv, w0.x, acc_a[0]);
                            acc_a[1] = fmaf(xv, w0.y, acc_a[1]);
                            acc_a[2] = fmaf(xv, w0.z, acc_a[2]);
                            acc_a[3] = fmaf(xv, w0.w, acc_a[3]);
                            acc_a[4] = fmaf(xv, w1.x, acc_a[4]);
                            acc_a[5] = fmaf(xv, w1.y, acc_a[5]);
                            acc_a[6] = fmaf(xv, w1.z, acc_a[6]);
                            acc_a[7] = fmaf(xv, w1.w, acc_a[7]);
                        }
                        if (valid_b && dlb >= 0) {
                            const float xv = sx[ic * 216 + (dlb * 6 + hl) * 6 + wl];
                            acc_b[0] = fmaf(xv, w0.x, acc_b[0]);
                            acc_b[1] = fmaf(xv, w0.y, acc_b[1]);
                            acc_b[2] = fmaf(xv, w0.z, acc_b[2]);
                            acc_b[3] = fmaf(xv, w0.w, acc_b[3]);
                            acc_b[4] = fmaf(xv, w1.x, acc_b[4]);
                            acc_b[5] = fmaf(xv, w1.y, acc_b[5]);
                            acc_b[6] = fmaf(xv, w1.z, acc_b[6]);
                            acc_b[7] = fmaf(xv, w1.w, acc_b[7]);
                        }
                    }
                }
            }
        }
        __syncthreads();
    }

    const long ostride = (long)Do * Ho * Wo;
    const long ybase = ((long)b * OC + oc0) * ostride + (long)H * Wo + W;
    if (valid_a) {
        const long yb = ybase + (long)Da * Ho * Wo;
        #pragma unroll
        for (int o = 0; o < OC_TILE; ++o) {
            if (oc0 + o < OC) Y[yb + (long)o * ostride] = acc_a[o] + bias[oc0 + o];
        }
    }
    if (valid_b) {
        const long yb = ybase + (long)Db * Ho * Wo;
        #pragma unroll
        for (int o = 0; o < OC_TILE; ++o) {
            if (oc0 + o < OC) Y[yb + (long)o * ostride] = acc_b[o] + bias[oc0 + o];
        }
    }
}

// ---------------------------------------------------------------------------
// Fused MaxPool(2,stride2)+MaxPool(3,stride3)+sum over channels.
// Composition of the non-overlapping 2x2x2 then 3x3x3 max-pools equals one
// non-overlapping 6x6x6 max-pool. Per output element, compute each channel's
// max over its 6x6x6 window and sum across the channels (requires C % 8 == 0).
// One block (256 threads = 8 warps) per output element; each warp handles one
// channel at a time so global loads are coalesced (lanes stride the window).
// ---------------------------------------------------------------------------
__global__ void pool6_sum_kernel(
    const float* __restrict__ Y,
    float* __restrict__ O,
    int C, int D, int H, int W,
    int OD, int OH, int OW)
{
    const int oid = blockIdx.x;
    const int ow = oid % OW;
    const int oh = (oid / OW) % OH;
    const int od = oid / (OW * OH);
    const int b  = blockIdx.y;

    const int tid  = threadIdx.x;   // 0..255
    const int warp = tid >> 5;      // 0..7 -> channel slot
    const int lane = tid & 31;      // window-position stride

    const long HW = (long)H * W;

    __shared__ float acc[8];
    if (tid < 8) acc[tid] = 0.f;
    __syncthreads();

    for (int cb = 0; cb < C; cb += 8) {
        const int c = cb + warp;
        const float* yb = Y + (((long)b * C + c) * D + (long)6 * od) * HW
                            + (long)6 * oh * W + 6 * ow;
        float m = -FLT_MAX;
        #pragma unroll
        for (int i = 0; i < 7; ++i) {
            const int v = lane + 32 * i;
            if (v < 216) {
                const int dz = v / 36;
                const int dy = (v / 6) % 6;
                const int dx = v % 6;
                m = fmaxf(m, yb[(long)dz * HW + (long)dy * W + dx]);
            }
        }
        #pragma unroll
        for (int off = 16; off > 0; off >>= 1)
            m = fmaxf(m, __shfl_down_sync(0xffffffffu, m, off));
        if (lane == 0) acc[warp] += m;
    }
    __syncthreads();
    if (tid == 0) {
        float s = 0.f;
        #pragma unroll
        for (int i = 0; i < 8; ++i) s += acc[i];
        O[((long)b * OD + od) * OH * OW + (long)oh * OW + ow] = s;
    }
}

torch::Tensor fused_forward_cuda(torch::Tensor x, torch::Tensor wp, torch::Tensor bias,
                                 int64_t stride, int64_t pad)
{
    TORCH_CHECK(x.is_cuda() && wp.is_cuda() && bias.is_cuda(), "tensors must be CUDA");
    const c10::cuda::OptionalCUDAGuard guard(x.device());
    auto xc  = x.contiguous();
    auto wpc = wp.contiguous();
    auto bc  = bias.contiguous();

    const int B  = xc.size(0);
    const int IC = xc.size(1);
    const int Di = xc.size(2);
    const int Hi = xc.size(3);
    const int Wi = xc.size(4);
    const int OC = bc.size(0);   // packed weight is [IC, 5, 5, 5, OC]
    TORCH_CHECK(wpc.size(0) == IC && wpc.size(4) == OC, "packed weight shape mismatch");

    const int Do = (Di - 1) * (int)stride - 2 * (int)pad + K;
    const int Ho = (Hi - 1) * (int)stride - 2 * (int)pad + K;
    const int Wo = (Wi - 1) * (int)stride - 2 * (int)pad + K;

    auto Y = torch::empty({B, OC, Do, Ho, Wo}, xc.options());

    const int ntw = (Wo + TILE - 1) / TILE;
    const int nth = (Ho + TILE - 1) / TILE;
    const int ntd = (Do + TILE - 1) / TILE;

    cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    dim3 grid(ntw * nth * ntd, (OC + OC_TILE - 1) / OC_TILE, B);
    convt3d_s2k5_kernel<<<grid, NTHREADS, 0, stream>>>(
        xc.data_ptr<float>(), wpc.data_ptr<float>(), bc.data_ptr<float>(),
        Y.data_ptr<float>(), IC, Di, Hi, Wi, OC, Do, Ho, Wo, (int)pad);
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    // fused pool1 (k=2,s=2) + pool2 (k=3,s=3) + channel sum
    const int D1 = (Do - 2) / 2 + 1;
    const int H1 = (Ho - 2) / 2 + 1;
    const int W1 = (Wo - 2) / 2 + 1;
    const int OD = (D1 - 3) / 3 + 1;
    const int OH = (H1 - 3) / 3 + 1;
    const int OW = (W1 - 3) / 3 + 1;

    auto O = torch::empty({B, 1, OD, OH, OW}, xc.options());
    dim3 pgrid(OD * OH * OW, B, 1);
    pool6_sum_kernel<<<pgrid, 256, 0, stream>>>(
        Y.data_ptr<float>(), O.data_ptr<float>(), OC, Do, Ho, Wo, OD, OH, OW);
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    return O;
}
"""

_cpp_src = "torch::Tensor fused_forward_cuda(torch::Tensor x, torch::Tensor wp, torch::Tensor bias, int64_t stride, int64_t pad);"

_fused = load_inline(
    name="convt3d_pool_sum_fused_v6",
    cpp_sources=_cpp_src,
    cuda_sources=_cuda_src,
    functions=["fused_forward_cuda"],
    verbose=False,
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Optimized Model:
      - ConvTranspose3d replaced by a custom tiled CUDA kernel (stride=2, k=5
        gather formulation) that evaluates only parity-matched kernel taps,
        with shared-memory blocking of inputs and pre-packed weights.
      - MaxPool3d(2) + MaxPool3d(3) + channel-sum fused into a single CUDA
        kernel exploiting that the composed non-overlapping pooling equals
        one non-overlapping 6x6x6 max-pool, with warp-per-channel coalesced
        loads and warp-shuffle reductions.
    """

    def __init__(self, in_channels, out_channels, kernel_size, stride, padding):
        super(ModelNew, self).__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):
        ct = self.conv_transpose
        if (
            x.is_cuda
            and x.dtype == torch.float32
            and ct.kernel_size == (5, 5, 5)
            and ct.stride == (2, 2, 2)
            and ct.padding == (2, 2, 2)
            and ct.dilation == (1, 1, 1)
            and ct.output_padding == (0, 0, 0)
            and ct.groups == 1
            and ct.bias is not None
            and ct.out_channels == 64
            and ct.weight.is_cuda
            and ct.weight.dtype == torch.float32
            and x.dim() == 5
            and x.size(2) >= 4
            and x.size(3) >= 4
            and x.size(4) >= 4
        ):
            # Pack weight [IC, OC, 5, 5, 5] -> [IC, 125, OC] (cheap: ~1MB copy).
            wp = ct.weight.permute(0, 2, 3, 4, 1).contiguous()
            return _fused.fused_forward_cuda(
                x, wp, ct.bias, ct.stride[0], ct.padding[0]
            )
        # generic fallback (matches reference semantics)
        x = self.conv_transpose(x)
        x = self.max_pool1(x)
        x = self.max_pool2(x)
        x = torch.sum(x, dim=1, keepdim=True)
        return x
```