```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 autotune the best algorithm for the fixed conv shape.
torch.backends.cudnn.benchmark = True

# Fused CUDA kernel implementing:
#   LeakyReLU(0.2) -> broadcast per-channel add -> clamp(-1, 1) -> exact GELU
# in a single memory pass over the conv output.
fused_epilogue_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <c10/cuda/CUDAException.h>
#include <cuda_runtime.h>

namespace {

constexpr float kNegSlope = 0.2f;
constexpr float kInvSqrt2 = 0.70710678118654752440f;

__device__ __forceinline__ float fused_op(float v, float s) {
    // LeakyReLU(negative_slope=0.2)
    v = v >= 0.0f ? v : kNegSlope * v;
    // Broadcast add of the per-channel sum tensor
    v += s;
    // clamp(x, -1, 1)
    v = fminf(fmaxf(v, -1.0f), 1.0f);
    // Exact GELU: x * 0.5 * (1 + erf(x / sqrt(2)))
    return 0.5f * v * (1.0f + erff(v * kInvSqrt2));
}

__global__ void fused_epilogue_vec4_kernel(
        const float4* __restrict__ x,
        const float* __restrict__ sum_t,
        float4* __restrict__ out,
        long long n4,
        long long spatial4,
        int channels) {
    long long i = (long long)blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n4) {
        // NCDHW contiguous layout: channel = (linear_idx / spatial) % C
        int c = (int)((i / spatial4) % (long long)channels);
        float s = __ldg(sum_t + c);
        float4 v = x[i];
        v.x = fused_op(v.x, s);
        v.y = fused_op(v.y, s);
        v.z = fused_op(v.z, s);
        v.w = fused_op(v.w, s);
        out[i] = v;
    }
}

__global__ void fused_epilogue_kernel(
        const float* __restrict__ x,
        const float* __restrict__ sum_t,
        float* __restrict__ out,
        long long total,
        long long spatial,
        int channels) {
    long long i = (long long)blockIdx.x * blockDim.x + threadIdx.x;
    if (i < total) {
        int c = (int)((i / spatial) % (long long)channels);
        out[i] = fused_op(x[i], __ldg(sum_t + c));
    }
}

} // namespace

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.scalar_type() == at::kFloat, "x must be float32");
    TORCH_CHECK(sum_t.is_cuda() && sum_t.scalar_type() == at::kFloat,
                "sum_t must be a CUDA float32 tensor");

    auto xc = x.contiguous();
    auto st = sum_t.contiguous();
    auto out = at::empty_like(xc);

    const int64_t total = xc.numel();
    const int channels = (int)xc.size(1);
    const int64_t spatial = total / ((int64_t)xc.size(0) * (int64_t)channels);

    const int block = 256;
    cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    if ((total & 3LL) == 0 && (spatial & 3LL) == 0) {
        const int64_t n4 = total >> 2;
        const int64_t s4 = spatial >> 2;
        const int64_t blocks = (n4 + block - 1) / block;
        fused_epilogue_vec4_kernel<<<blocks, block, 0, stream>>>(
            reinterpret_cast<const float4*>(xc.data_ptr<float>()),
            st.data_ptr<float>(),
            reinterpret_cast<float4*>(out.data_ptr<float>()),
            (long long)n4, (long long)s4, channels);
    } else {
        const int64_t blocks = (total + block - 1) / block;
        fused_epilogue_kernel<<<blocks, block, 0, stream>>>(
            xc.data_ptr<float>(),
            st.data_ptr<float>(),
            out.data_ptr<float>(),
            (long long)total, (long long)spatial, channels);
    }
    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return out;
}
"""

fused_epilogue_cpp_source = (
    "torch::Tensor fused_epilogue_cuda(torch::Tensor x, torch::Tensor sum_t);"
)

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


class ModelNew(nn.Module):
    """
    Optimized model: Conv3d runs through cuDNN (benchmark mode), and the
    LeakyReLU + per-channel add + clamp + GELU chain is fused into a single
    custom CUDA kernel, eliminating 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
```