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

# The conv3d is kept as a cuDNN call (it is heavily optimized / uses TF32
# tensor cores). The four pointwise ops that follow it
# (LeakyReLU -> +sum_tensor -> clamp -> GELU) are all memory bound, so they
# are fused into a single custom CUDA kernel, eliminating 3 full
# read+write passes over the (large) activation tensor.

_fused_epilogue_cuda_src = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <c10/cuda/CUDAException.h>
#include <cuda_runtime.h>
#include <cstdint>

__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 sum tensor
    v = v + s;
    // clamp to [-1, 1]
    v = fminf(fmaxf(v, -1.0f), 1.0f);
    // exact (erf) GELU: 0.5 * x * (1 + erf(x / sqrt(2)))
    return 0.5f * v * (1.0f + erff(v * 0.70710678118654752440f));
}

// Vectorized kernel: each thread handles 4 contiguous floats.
// blockIdx.y indexes the (batch * channel) plane so the per-channel
// scalar is fetched without any per-element integer division.
__global__ void fused_epilogue_vec4_kernel(const float* __restrict__ x,
                                           const float* __restrict__ sum_ch,
                                           float* __restrict__ out,
                                           const int plane_vec,
                                           const int channels) {
    const float s = sum_ch[(int)(blockIdx.y % (unsigned)channels)];
    const int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < plane_vec) {
        const int64_t vbase = (int64_t)blockIdx.y * (int64_t)plane_vec;
        const float4* x4 = reinterpret_cast<const float4*>(x) + vbase;
        float4 v = x4[idx];
        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);
        float4* o4 = reinterpret_cast<float4*>(out) + vbase;
        o4[idx] = v;
    }
}

// Fully general fallback kernel (handles any shape).
__global__ void fused_epilogue_scalar_kernel(const float* __restrict__ x,
                                             const float* __restrict__ sum_ch,
                                             float* __restrict__ out,
                                             const int64_t plane_size,
                                             const int64_t channels,
                                             const int64_t total) {
    const int64_t idx = (int64_t)blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < total) {
        const int64_t p = idx / plane_size;
        const float s = sum_ch[p % channels];
        out[idx] = epilogue_op(x[idx], s);
    }
}

torch::Tensor fused_epilogue_cuda(torch::Tensor x, torch::Tensor sum_tensor) {
    TORCH_CHECK(x.is_cuda() && sum_tensor.is_cuda(), "inputs must be CUDA tensors");
    TORCH_CHECK(x.scalar_type() == at::kFloat, "x must be float32");
    TORCH_CHECK(sum_tensor.scalar_type() == at::kFloat, "sum_tensor must be float32");
    TORCH_CHECK(x.dim() >= 2, "x must have at least 2 dims");

    auto x_c = x.contiguous();
    auto s_c = sum_tensor.contiguous();
    auto out = at::empty_like(x_c);

    const int64_t channels = x_c.size(1);
    int64_t plane_size = 1;
    for (int64_t i = 2; i < x_c.dim(); ++i) plane_size *= x_c.size(i);
    const int64_t planes = x_c.size(0) * channels;
    const int64_t total = x_c.numel();
    if (total == 0) return out;

    cudaStream_t stream = at::cuda::getCurrentCUDAStream();
    const float* xp = x_c.data_ptr<float>();
    const float* sp = s_c.data_ptr<float>();
    float* op = out.data_ptr<float>();

    if ((plane_size % 4 == 0) && (planes <= 65535)) {
        const int plane_vec = (int)(plane_size >> 2);
        const int block = 256;
        dim3 grid((unsigned)((plane_vec + block - 1) / block), (unsigned)planes);
        fused_epilogue_vec4_kernel<<<grid, block, 0, stream>>>(
            xp, sp, op, plane_vec, (int)channels);
    } else {
        const int block = 256;
        const int64_t blocks = (total + block - 1) / block;
        fused_epilogue_scalar_kernel<<<(unsigned)blocks, block, 0, stream>>>(
            xp, sp, op, plane_size, channels, total);
    }
    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return out;
}
"""

_fused_epilogue_cpp_src = (
    "torch::Tensor fused_epilogue_cuda(torch::Tensor x, torch::Tensor sum_tensor);"
)

_fused_epilogue_ext = load_inline(
    name="fused_conv3d_epilogue_ext",
    cpp_sources=_fused_epilogue_cpp_src,
    cuda_sources=_fused_epilogue_cuda_src,
    functions=["fused_epilogue_cuda"],
    verbose=False,
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Optimized model: cuDNN Conv3d followed by a single fused CUDA kernel
    that performs LeakyReLU + per-channel sum + clamp + GELU in one pass.
    """

    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_ext.fused_epilogue_cuda

    def forward(self, x):
        x = F.conv3d(x, self.conv.weight, self.conv.bias)
        x = self.fused_epilogue(x, self.sum_tensor)
        return x
```