Looking at the profiler breakdown: the conv (4.0ms) and my fused epilogue (2.6ms, ~3.5GB of traffic at ~1.35TB/s — already near roofline since the output tensor is 441M elements/1.76GB) are fine. The big waste is the `x.contiguous(channels_last_3d)` call taking 2587µs via PyTorch's generic permute-copy (uncoalesced strided writes). I'll replace it with a custom NCDHW→NDHWC transpose kernel that is fully coalesced on both sides: each thread handles one spatial pixel, reads C channel values (each read is warp-contiguous), and writes C/4 consecutive float4s (warp-contiguous) — no shared memory needed. Expected ~350µs instead of ~2.6ms.

```python
# Copyright (c) Meta Platforms, Inc. and affiliates.
# All rights reserved.
#
# This source code is licensed under the license found in the
# LICENSE file in the root directory of this source tree.

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

# Autotune cuDNN (exhaustive search) for the fixed conv shapes used here.
torch.backends.cudnn.benchmark = True
torch.backends.cudnn.benchmark_limit = 0
torch.backends.cudnn.allow_tf32 = True
torch.backends.cuda.matmul.allow_tf32 = True

# Strategy (unchanged in structure, improved in the layout conversion):
#  * Run the 3D conv natively in channels-last (NDHWC) so cuDNN uses its
#    tensor-core NHWC kernels with no internal transpose kernels.
#  * NEW: replace PyTorch's slow generic permute-copy for the NCDHW->NDHWC
#    input conversion (was ~2.6ms) with a custom transpose kernel that is
#    fully coalesced on BOTH sides: each thread handles one spatial pixel,
#    issues C warp-contiguous strided loads, and writes C/4 consecutive
#    float4 stores (warp-contiguous). No shared memory, near-BW (~0.35ms).
#  * Keep the fused LeakyReLU -> +sum_tensor -> clamp -> GELU epilogue as a
#    single vectorized kernel over the NHWC data (already near roofline).
fused_ops_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>

// ------------------- NCDHW -> NDHWC transpose -------------------
// y[n,d,h,w,c] = x[n,c,d,h,w].  Both sides fully coalesced.
__global__ void __launch_bounds__(256) ncdhw_to_ndhwc_vec_kernel(
        const float* __restrict__ x,
        float* __restrict__ y,
        long P,      // N * D*H*W  (number of spatial pixels)
        long S,      // D*H*W      (per-channel spatial size)
        int C) {     // channels, multiple of 4
    const long p = (long)blockIdx.x * blockDim.x + threadIdx.x;
    if (p >= P) return;
    const long n = p / S;
    const long s = p - n * S;
    const float* src = x + n * (long)C * S + s;   // &x[n,0,s]
    float4* dst = reinterpret_cast<float4*>(y + p * (long)C);
    const int cvec = C >> 2;
    for (int cv = 0; cv < cvec; ++cv) {
        float4 v;
        v.x = src[(long)(4 * cv + 0) * S];
        v.y = src[(long)(4 * cv + 1) * S];
        v.z = src[(long)(4 * cv + 2) * S];
        v.w = src[(long)(4 * cv + 3) * S];
        dst[cv] = v;
    }
}

__global__ void __launch_bounds__(256) ncdhw_to_ndhwc_scalar_kernel(
        const float* __restrict__ x,
        float* __restrict__ y,
        long P, long S, int C) {
    const long p = (long)blockIdx.x * blockDim.x + threadIdx.x;
    if (p >= P) return;
    const long n = p / S;
    const long s = p - n * S;
    const float* src = x + n * (long)C * S + s;
    float* dst = y + p * (long)C;
    for (int c = 0; c < C; ++c) dst[c] = src[(long)c * S];
}

torch::Tensor ncdhw_to_ndhwc_cuda(torch::Tensor x) {
    TORCH_CHECK(x.is_cuda(), "x must be a CUDA tensor");
    TORCH_CHECK(x.scalar_type() == at::kFloat, "x must be float32");
    TORCH_CHECK(x.dim() == 5, "x must be a 5D NCDHW tensor");
    if (x.is_contiguous(at::MemoryFormat::ChannelsLast3d)) return x;

    auto xc = x.contiguous();
    auto y = torch::empty(xc.sizes(),
                          xc.options().memory_format(at::MemoryFormat::ChannelsLast3d));
    const int C = (int)xc.size(1);
    long S = 1;
    for (int i = 2; i < 5; ++i) S *= xc.size(i);
    const long P = xc.size(0) * S;
    const int block = 256;
    const long grid = (P + block - 1) / block;
    cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    if ((C & 3) == 0) {
        ncdhw_to_ndhwc_vec_kernel<<<(unsigned)grid, block, 0, stream>>>(
            xc.data_ptr<float>(), y.data_ptr<float>(), P, S, C);
    } else {
        ncdhw_to_ndhwc_scalar_kernel<<<(unsigned)grid, block, 0, stream>>>(
            xc.data_ptr<float>(), y.data_ptr<float>(), P, S, C);
    }
    return y;
}

// ------------- fused LeakyReLU -> add -> clamp -> GELU epilogue -------------
__device__ __forceinline__ float fused_op(float v, float s) {
    v = (v > 0.0f) ? v : v * 0.2f;                                   // LeakyReLU(0.2)
    v += s;                                                          // per-channel add
    v = fminf(fmaxf(v, -1.0f), 1.0f);                                // clamp(-1, 1)
    return 0.5f * v * (1.0f + erff(v * 0.70710678118654752440f));    // exact GELU
}

// ---- channels-last (NDHWC) path: float4 spans 4 consecutive channels ----
__global__ void fused_cl_vec4_kernel(const float4* __restrict__ x,
                                     const float4* __restrict__ s4,
                                     float4* __restrict__ out,
                                     long total_vec,
                                     int c_vec,
                                     long c_mask) {
    const long stride = (long)gridDim.x * blockDim.x;
    for (long i = (long)blockIdx.x * blockDim.x + threadIdx.x; i < total_vec; i += stride) {
        const int cv = (c_mask >= 0) ? (int)(i & c_mask) : (int)(i % c_vec);
        const float4 sv = s4[cv];
        const float4 v = x[i];
        float4 r;
        r.x = fused_op(v.x, sv.x);
        r.y = fused_op(v.y, sv.y);
        r.z = fused_op(v.z, sv.z);
        r.w = fused_op(v.w, sv.w);
        out[i] = r;
    }
}

__global__ void fused_cl_scalar_kernel(const float* __restrict__ x,
                                       const float* __restrict__ s,
                                       float* __restrict__ out,
                                       long total,
                                       int channels) {
    const long stride = (long)gridDim.x * blockDim.x;
    for (long i = (long)blockIdx.x * blockDim.x + threadIdx.x; i < total; i += stride) {
        out[i] = fused_op(x[i], s[i % channels]);
    }
}

// ---- NCDHW fallback path: float4 spans 4 consecutive spatial locations ----
__global__ void fused_vec4_kernel(const float4* __restrict__ x,
                                  const float* __restrict__ sum_tensor,
                                  float4* __restrict__ out,
                                  int spatial_vec,
                                  int channels) {
    const int nc = blockIdx.y;
    const float sv = sum_tensor[nc % channels];
    const long base = (long)nc * spatial_vec;
    const float4* xs = x + base;
    float4* os = out + base;
    const int stride = gridDim.x * blockDim.x;
    for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < spatial_vec; i += stride) {
        const float4 v = xs[i];
        float4 r;
        r.x = fused_op(v.x, sv);
        r.y = fused_op(v.y, sv);
        r.z = fused_op(v.z, sv);
        r.w = fused_op(v.w, sv);
        os[i] = r;
    }
}

__global__ void fused_scalar_kernel(const float* __restrict__ x,
                                    const float* __restrict__ sum_tensor,
                                    float* __restrict__ out,
                                    long spatial_size,
                                    int channels) {
    const int nc = blockIdx.y;
    const float sv = sum_tensor[nc % channels];
    const long base = (long)nc * spatial_size;
    const float* xs = x + base;
    float* os = out + base;
    const long stride = (long)gridDim.x * blockDim.x;
    for (long i = (long)blockIdx.x * blockDim.x + threadIdx.x; i < spatial_size; i += stride) {
        os[i] = fused_op(xs[i], sv);
    }
}

torch::Tensor fused_postprocess_cuda(torch::Tensor x, torch::Tensor sum_tensor) {
    TORCH_CHECK(x.is_cuda(), "x must be a CUDA tensor");
    TORCH_CHECK(x.scalar_type() == at::kFloat, "x must be float32");
    TORCH_CHECK(sum_tensor.is_cuda(), "sum_tensor must be a CUDA tensor");

    auto sc = sum_tensor.contiguous();
    const int channels = (int)x.size(1);
    const int block = 256;
    cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    if (x.is_contiguous(at::MemoryFormat::ChannelsLast3d)) {
        const long total = x.numel();
        if ((channels & 3) == 0) {
            const long total_vec = total >> 2;
            const int c_vec = channels >> 2;
            const long c_mask = ((c_vec & (c_vec - 1)) == 0) ? (long)(c_vec - 1) : -1L;
            long want = (total_vec + block - 1) / block;
            int gx = (int)(want > 65535 ? 65535 : want);
            fused_cl_vec4_kernel<<<gx, block, 0, stream>>>(
                reinterpret_cast<const float4*>(x.data_ptr<float>()),
                reinterpret_cast<const float4*>(sc.data_ptr<float>()),
                reinterpret_cast<float4*>(x.data_ptr<float>()),
                total_vec, c_vec, c_mask);
        } else {
            long want = (total + block - 1) / block;
            int gx = (int)(want > 65535 ? 65535 : want);
            fused_cl_scalar_kernel<<<gx, block, 0, stream>>>(
                x.data_ptr<float>(), sc.data_ptr<float>(), x.data_ptr<float>(),
                total, channels);
        }
        return x;
    }

    auto xc = x.contiguous();
    long spatial = 1;
    for (int i = 2; i < xc.dim(); ++i) spatial *= xc.size(i);
    const long nc_total = xc.numel() / spatial;

    if (spatial % 4 == 0) {
        const int spatial_vec = (int)(spatial >> 2);
        long want = (spatial_vec + block - 1) / block;
        int gx = (int)(want > 65535 ? 65535 : want);
        dim3 grid(gx, (unsigned)nc_total);
        fused_vec4_kernel<<<grid, block, 0, stream>>>(
            reinterpret_cast<const float4*>(xc.data_ptr<float>()),
            sc.data_ptr<float>(),
            reinterpret_cast<float4*>(xc.data_ptr<float>()),
            spatial_vec, channels);
    } else {
        long want = (spatial + block - 1) / block;
        int gx = (int)(want > 65535 ? 65535 : want);
        dim3 grid(gx, (unsigned)nc_total);
        fused_scalar_kernel<<<grid, block, 0, stream>>>(
            xc.data_ptr<float>(), sc.data_ptr<float>(), xc.data_ptr<float>(),
            spatial, channels);
    }
    return xc;
}
"""

fused_ops_cpp_source = (
    "torch::Tensor ncdhw_to_ndhwc_cuda(torch::Tensor x);\n"
    "torch::Tensor fused_postprocess_cuda(torch::Tensor x, torch::Tensor sum_tensor);"
)

fused_ops = load_inline(
    name="fused_model_ops_v3",
    cpp_sources=fused_ops_cpp_source,
    cuda_sources=fused_ops_source,
    functions=["ncdhw_to_ndhwc_cuda", "fused_postprocess_cuda"],
    verbose=False,
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Optimized model: custom fully-coalesced NCDHW->NDHWC transpose kernel,
    then the 3D convolution runs natively in channels-last layout so cuDNN
    uses its tensor-core NHWC kernels with no layout-transpose kernels, and
    finally LeakyReLU + per-channel sum + clamp + GELU are applied in a
    single fused, vectorized CUDA kernel over the NHWC data.
    """
    def __init__(self, in_channels, out_channels, kernel_size, sum_tensor_shape):
        super(ModelNew, self).__init__()
        self.conv = nn.Conv3d(in_channels, out_channels, kernel_size)
        self.sum_tensor = nn.Parameter(torch.randn(sum_tensor_shape))
        # Keep conv weights in channels-last layout (one-time cost) so cuDNN
        # never inserts NCHW<->NHWC transpose kernels.
        with torch.no_grad():
            self.conv.weight.data = self.conv.weight.data.contiguous(
                memory_format=torch.channels_last_3d)

    def forward(self, x):
        x = fused_ops.ncdhw_to_ndhwc_cuda(x)
        x = self.conv(x)
        x = fused_ops.fused_postprocess_cuda(
            x, self.sum_tensor.view(-1))
        return x
```