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

torch.backends.cudnn.benchmark = True

# Fused CUDA kernel: LeakyReLU(0.2) + broadcast channel-wise add + clamp(-1, 1) + exact GELU
fused_post_source = r"""
#include <torch/extension.h>
#include <cuda_runtime.h>

__device__ __forceinline__ float gelu_exact(float v) {
    // GELU(x) = x * Phi(x) = 0.5 * x * (1 + erf(x / sqrt(2)))
    return 0.5f * v * (1.0f + erff(v * 0.70710678118654752440f));
}

__device__ __forceinline__ float post_op(float v, float s) {
    v = v > 0.0f ? v : 0.2f * v;              // LeakyReLU(0.2)
    v = v + s;                                // channel-wise bias add
    v = fminf(fmaxf(v, -1.0f), 1.0f);         // clamp(-1, 1)
    return gelu_exact(v);                     // exact GELU
}

__global__ void fused_leaky_add_clamp_gelu_kernel(
    const float* __restrict__ x,
    const float* __restrict__ bias,
    float* __restrict__ out,
    long long total,
    long long inner,
    long long C)
{
    long long idx = (long long)blockIdx.x * blockDim.x + threadIdx.x;
    long long stride = (long long)gridDim.x * blockDim.x;

    if ((inner & 3LL) == 0) {
        // Vectorized path: 4 contiguous elements share the same channel
        long long total4 = total >> 2;
        const float4* x4 = reinterpret_cast<const float4*>(x);
        float4* out4 = reinterpret_cast<float4*>(out);
        for (long long i = idx; i < total4; i += stride) {
            float s = bias[((i << 2) / inner) % C];
            float4 v = x4[i];
            v.x = post_op(v.x, s);
            v.y = post_op(v.y, s);
            v.z = post_op(v.z, s);
            v.w = post_op(v.w, s);
            out4[i] = v;
        }
    } else {
        // Scalar fallback path
        for (long long i = idx; i < total; i += stride) {
            float s = bias[(i / inner) % C];
            out[i] = post_op(x[i], s);
        }
    }
}

torch::Tensor fused_post_cuda(torch::Tensor x, torch::Tensor bias) {
    TORCH_CHECK(x.is_cuda() && bias.is_cuda(), "Inputs must be CUDA tensors");
    auto xc = x.contiguous();
    auto bc = bias.contiguous();
    auto out = torch::empty_like(xc);

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

    const int block = 256;
    long long work = ((inner & 3LL) == 0) ? (total >> 2) : total;
    long long grid = (work + block - 1) / block;

    fused_leaky_add_clamp_gelu_kernel<<<(unsigned int)grid, block>>>(
        xc.data_ptr<float>(),
        bc.data_ptr<float>(),
        out.data_ptr<float>(),
        total, inner, C);

    return out;
}
"""

fused_post_cpp_source = (
    "torch::Tensor fused_post_cuda(torch::Tensor x, torch::Tensor bias);"
)

fused_post = load_inline(
    name="fused_leaky_add_clamp_gelu",
    cpp_sources=fused_post_cpp_source,
    cuda_sources=fused_post_source,
    functions=["fused_post_cuda"],
    verbose=False,
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Optimized model: Conv3d (cuDNN) followed by a single fused custom CUDA kernel
    that performs LeakyReLU + broadcast add + 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))

    def forward(self, x):
        x = self.conv(x)
        x = fused_post.fused_post_cuda(x, self.sum_tensor)
        return x
```