```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
import torch.nn.functional as F
from torch.utils.cpp_extension import load_inline

# Let cuDNN pick the fastest convolution algorithm and allow TF32 tensor cores.
torch.backends.cudnn.benchmark = True
torch.backends.cudnn.allow_tf32 = True

# Fused CUDA kernel: LeakyReLU(0.2) -> + per-channel sum_tensor -> clamp(-1,1) -> GELU(erf)
fused_epilogue_cpp_source = (
    "torch::Tensor fused_epilogue_cuda(torch::Tensor x, torch::Tensor sum_t);"
)

fused_epilogue_cuda_source = r"""
#include <torch/extension.h>
#include <cuda_runtime.h>

__device__ __forceinline__ float epilogue_op(float v, float s) {
    // LeakyReLU with negative_slope = 0.2
    v = v >= 0.0f ? v : 0.2f * v;
    // Add per-channel broadcast sum tensor
    v = v + s;
    // Clamp to [-1, 1]
    v = fminf(fmaxf(v, -1.0f), 1.0f);
    // Exact (erf-based) GELU: 0.5 * x * (1 + erf(x / sqrt(2)))
    return 0.5f * v * (1.0f + erff(v * 0.70710678118654752440f));
}

__global__ void fused_epilogue_scalar_kernel(
    const float* __restrict__ in,
    const float* __restrict__ sum_t,
    float* __restrict__ out,
    long total,
    long vol,
    int C)
{
    long i = (long)blockIdx.x * blockDim.x + threadIdx.x;
    if (i < total) {
        int c = (int)((i / vol) % (long)C);
        out[i] = epilogue_op(in[i], sum_t[c]);
    }
}

// Vectorized (float4) path. Requires vol % 4 == 0, which guarantees all 4
// elements of a vector belong to the same channel.
__global__ void fused_epilogue_vec4_kernel(
    const float4* __restrict__ in,
    const float* __restrict__ sum_t,
    float4* __restrict__ out,
    long total4,
    long vol,
    int C)
{
    long i = (long)blockIdx.x * blockDim.x + threadIdx.x;
    if (i < total4) {
        int c = (int)(((i * 4) / vol) % (long)C);
        float s = sum_t[c];
        float4 v = in[i];
        v.x = epilogue_op(v.x, s);
        v.y = epilogue_op(v.y, s);
        v.z = epilogue_op(v.z, s);
        v.w = epilogue_op(v.w, s);
        out[i] = v;
    }
}

torch::Tensor fused_epilogue_cuda(torch::Tensor x, torch::Tensor sum_t) {
    TORCH_CHECK(x.is_cuda(), "x must be a CUDA tensor");
    TORCH_CHECK(x.dtype() == torch::kFloat32, "x must be float32");
    TORCH_CHECK(sum_t.is_cuda() && sum_t.dtype() == torch::kFloat32,
                "sum_t must be a float32 CUDA tensor");

    auto xc = x.contiguous();
    auto out = torch::empty_like(xc);

    long total = (long)xc.numel();
    int C = (int)xc.size(1);
    long vol = (long)xc.size(2) * (long)xc.size(3) * (long)xc.size(4);

    const int block = 256;

    if (total == 0) {
        return out;
    }

    if (vol % 4 == 0 && total % 4 == 0) {
        long total4 = total / 4;
        int grid = (int)((total4 + block - 1) / block);
        fused_epilogue_vec4_kernel<<<grid, block>>>(
            reinterpret_cast<const float4*>(xc.data_ptr<float>()),
            sum_t.data_ptr<float>(),
            reinterpret_cast<float4*>(out.data_ptr<float>()),
            total4, vol, C);
    } else {
        int grid = (int)((total + block - 1) / block);
        fused_epilogue_scalar_kernel<<<grid, block>>>(
            xc.data_ptr<float>(),
            sum_t.data_ptr<float>(),
            out.data_ptr<float>(),
            total, vol, C);
    }

    return out;
}
"""

fused_epilogue = load_inline(
    name="fused_epilogue_conv3d",
    cpp_sources=fused_epilogue_cpp_source,
    cuda_sources=fused_epilogue_cuda_source,
    functions=["fused_epilogue_cuda"],
    verbose=False,
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Optimized model: the 3D convolution is kept on the highly-tuned cuDNN
    backend, while the entire elementwise chain
    (LeakyReLU -> +sum_tensor -> clamp -> GELU) is fused into a single
    custom CUDA kernel to eliminate multiple full-tensor memory round-trips.
    """
    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))
        self.fused_epilogue = fused_epilogue

    def forward(self, x):
        x = self.conv(x)
        x = self.fused_epilogue.fused_epilogue_cuda(x, self.sum_tensor)
        return x
```