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

# Fused CUDA kernel implementing the entire pointwise chain:
#   LeakyReLU(0.2) -> + per-channel sum_tensor -> clamp(-1, 1) -> exact (erf) GELU
# This replaces 4 separate memory-bound PyTorch kernels with a single pass.
_fused_epilogue_cuda_src = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>
#include <cstdint>

__device__ __forceinline__ float fused_op(float v, float s) {
    // LeakyReLU with negative_slope = 0.2
    v = (v > 0.0f) ? v : (0.2f * v);
    // Add per-channel broadcast tensor
    v = v + s;
    // Clamp to [-1, 1]
    v = fminf(fmaxf(v, -1.0f), 1.0f);
    // Exact GELU (erf formulation, matches torch.nn.functional.gelu default)
    return 0.5f * v * (1.0f + erff(v * 0.70710678118654752440f));
}

__global__ void fused_epilogue_vec4_kernel(const float4* __restrict__ x,
                                           const float*  __restrict__ s,
                                           float4*       __restrict__ out,
                                           const int64_t n4,
                                           const int64_t channel_stride4,
                                           const int C) {
    const int64_t i = (int64_t)blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n4) {
        // channel_stride4 is divisible by 4 so a float4 group never crosses
        // a channel boundary.
        const int c = (int)((i / channel_stride4) % (int64_t)C);
        const float sv = __ldg(s + c);
        float4 v = x[i];
        v.x = fused_op(v.x, sv);
        v.y = fused_op(v.y, sv);
        v.z = fused_op(v.z, sv);
        v.w = fused_op(v.w, sv);
        out[i] = v;
    }
}

__global__ void fused_epilogue_scalar_kernel(const float* __restrict__ x,
                                             const float* __restrict__ s,
                                             float*       __restrict__ out,
                                             const int64_t n,
                                             const int64_t channel_stride,
                                             const int C) {
    const int64_t i = (int64_t)blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        const int c = (int)((i / channel_stride) % (int64_t)C);
        out[i] = fused_op(x[i], s[c]);
    }
}

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

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

    const int64_t n = xc.numel();
    const int C = (int)xc.size(1);
    const int64_t dhw = xc.size(2) * xc.size(3) * xc.size(4);
    TORCH_CHECK(sc.numel() == C, "s must have one element per channel");

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

    if (n > 0) {
        if ((n & 3) == 0 && (dhw & 3) == 0) {
            const int64_t n4 = n >> 2;
            const unsigned int blocks =
                (unsigned int)((n4 + threads - 1) / threads);
            fused_epilogue_vec4_kernel<<<blocks, threads, 0, stream>>>(
                reinterpret_cast<const float4*>(xc.data_ptr<float>()),
                sc.data_ptr<float>(),
                reinterpret_cast<float4*>(out.data_ptr<float>()),
                n4, dhw >> 2, C);
        } else {
            const unsigned int blocks =
                (unsigned int)((n + threads - 1) / threads);
            fused_epilogue_scalar_kernel<<<blocks, threads, 0, stream>>>(
                xc.data_ptr<float>(),
                sc.data_ptr<float>(),
                out.data_ptr<float>(),
                n, dhw, C);
        }
    }
    return out;
}
"""

_fused_epilogue_cpp_src = "torch::Tensor fused_epilogue(torch::Tensor x, torch::Tensor s);"

_fused_epilogue_mod = load_inline(
    name="fused_epilogue_leaky_add_clamp_gelu",
    cpp_sources=_fused_epilogue_cpp_src,
    cuda_sources=_fused_epilogue_cuda_src,
    functions=["fused_epilogue"],
    verbose=False,
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Optimized model: the 3D convolution stays on the highly-optimized cuDNN
    path, while the whole pointwise chain (LeakyReLU + broadcast add + clamp
    + GELU) is executed by a single fused custom CUDA kernel, eliminating 3
    full read/write passes over the activation tensor.
    """
    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_mod

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