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

# Fused CUDA kernel that applies, in a single pass over the conv output:
#   LeakyReLU(0.2) -> + per-channel sum_tensor -> clamp(-1, 1) -> exact GELU
# This eliminates 4 separate kernel launches and their memory round-trips.
fused_epilogue_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <c10/cuda/CUDAGuard.h>
#include <cuda_runtime.h>

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

__global__ void fused_epilogue_vec4_kernel(const float4* __restrict__ x,
                                           const float* __restrict__ bias,
                                           float4* __restrict__ out,
                                           int C,
                                           long long spatial,
                                           long long total4) {
    long long i = (long long)blockIdx.x * blockDim.x + threadIdx.x;
    if (i < total4) {
        // All 4 packed elements belong to the same channel because
        // spatial is a multiple of 4.
        int c = (int)(((i << 2) / spatial) % C);
        float b = __ldg(bias + c);
        float4 v = x[i];
        v.x = epilogue(v.x, b);
        v.y = epilogue(v.y, b);
        v.z = epilogue(v.z, b);
        v.w = epilogue(v.w, b);
        out[i] = v;
    }
}

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

torch::Tensor fused_epilogue_cuda(torch::Tensor x, torch::Tensor bias) {
    TORCH_CHECK(x.is_cuda() && bias.is_cuda(), "inputs must be CUDA tensors");
    TORCH_CHECK(x.scalar_type() == at::kFloat && bias.scalar_type() == at::kFloat,
                "inputs must be float32");

    const c10::cuda::OptionalCUDAGuard device_guard(x.device());

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

    long long total = (long long)xc.numel();
    int C = (int)xc.size(1);
    long long spatial = 1;
    for (int d = 2; d < xc.dim(); ++d) {
        spatial *= (long long)xc.size(d);
    }

    cudaStream_t stream = at::cuda::getCurrentCUDAStream();
    const float* xp = xc.data_ptr<float>();
    const float* bp = bc.data_ptr<float>();
    float* op = out.data_ptr<float>();

    const int block = 256;
    if ((total & 3LL) == 0 && (spatial & 3LL) == 0) {
        long long total4 = total >> 2;
        long long grid = (total4 + block - 1) / block;
        fused_epilogue_vec4_kernel<<<(unsigned int)grid, block, 0, stream>>>(
            reinterpret_cast<const float4*>(xp), bp,
            reinterpret_cast<float4*>(op), C, spatial, total4);
    } else {
        long long grid = (total + block - 1) / block;
        fused_epilogue_kernel<<<(unsigned int)grid, block, 0, stream>>>(
            xp, bp, op, C, spatial, total);
    }
    return out;
}
"""

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

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


class ModelNew(nn.Module):
    """
    Optimized model: the 3D convolution runs through cuDNN (with autotuning
    enabled), and LeakyReLU + channel-wise sum + clamp + GELU are executed in a
    single fused custom CUDA kernel.
    """
    def __init__(self, in_channels, out_channels, kernel_size, sum_tensor_shape):
        super(ModelNew, self).__init__()
        torch.backends.cudnn.benchmark = True
        self.conv = nn.Conv3d(in_channels, out_channels, kernel_size)
        self.sum_tensor = nn.Parameter(torch.randn(sum_tensor_shape))

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