```python
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline

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

__device__ __forceinline__ float gelu(float x) {
    float u = 0.7978845608028654f *
              fmaf(0.044715f * x, x * x, x);
    float t;
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 750
    asm("tanh.approx.f32 %0, %1;" : "=f"(t) : "f"(u));
#else
    t = tanhf(u);
#endif
    return (0.5f * x) * (1.0f + t);
}

__device__ __forceinline__ float4 load_vector(const float4* p) {
    float4 v;
    asm volatile(
        "ld.global.cg.v4.f32 {%0, %1, %2, %3}, [%4];"
        : "=f"(v.x), "=f"(v.y), "=f"(v.z), "=f"(v.w)
        : "l"(p)
        : "memory"
    );
    return v;
}

__device__ __forceinline__ void store_vector(float4* p, float4 v) {
    asm volatile(
        "st.global.cs.v4.f32 [%0], {%1, %2, %3, %4};"
        :
        : "l"(p), "f"(v.x), "f"(v.y), "f"(v.z), "f"(v.w)
        : "memory"
    );
}

__device__ __forceinline__ float4 activate(float4 v) {
    v.x = gelu(v.x);
    v.y = gelu(v.y);
    v.z = gelu(v.z);
    v.w = gelu(v.w);
    return v;
}

template <typename Index, bool FullTiles>
__global__ __launch_bounds__(256, 4)
void gelu_vector_kernel(
    const float* __restrict__ input,
    float* __restrict__ output,
    Index n
) {
    constexpr int threads = 256;
    constexpr int tile_vectors = threads * 4;
    const Index nv = n / 4;
    const auto* x = reinterpret_cast<const float4*>(input);
    auto* y = reinterpret_cast<float4*>(output);
    const Index stride = static_cast<Index>(gridDim.x) * tile_vectors;

    for (Index tile = static_cast<Index>(blockIdx.x) * tile_vectors;
         tile < nv;
         tile += stride) {
        const Index base = tile + threadIdx.x;
        float4 v[4];

        #pragma unroll
        for (int j = 0; j < 4; ++j) {
            const Index i = base + j * threads;
            if (FullTiles || i < nv)
                v[j] = load_vector(x + i);
        }

        #pragma unroll
        for (int j = 0; j < 4; ++j) {
            const Index i = base + j * threads;
            if (FullTiles || i < nv)
                store_vector(y + i, activate(v[j]));
        }
    }

    if (!FullTiles && blockIdx.x == 0) {
        const Index i = nv * 4 + threadIdx.x;
        if (i < n)
            output[i] = gelu(input[i]);
    }
}

template <typename Index>
__global__ void gelu_scalar_kernel(
    const float* __restrict__ x,
    float* __restrict__ y,
    Index n
) {
    const Index stride = static_cast<Index>(gridDim.x) * blockDim.x;
    for (Index i = static_cast<Index>(blockIdx.x) * blockDim.x
                   + threadIdx.x;
         i < n;
         i += stride) {
        y[i] = gelu(x[i]);
    }
}

template <typename Index>
void launch_gelu(
    const float* x,
    float* y,
    int64_t n,
    cudaStream_t stream
) {
    constexpr int threads = 256;

    if ((reinterpret_cast<uintptr_t>(x) & 15) == 0) {
        constexpr int tile_elements = threads * 16;
        const int blocks = static_cast<int>(
            std::min<int64_t>(
                (n + tile_elements - 1) / tile_elements, 4096));

        if (n % tile_elements == 0) {
            gelu_vector_kernel<Index, true>
                <<<blocks, threads, 0, stream>>>(
                    x, y, static_cast<Index>(n));
        } else {
            gelu_vector_kernel<Index, false>
                <<<blocks, threads, 0, stream>>>(
                    x, y, static_cast<Index>(n));
        }
    } else {
        const int blocks = static_cast<int>(
            std::min<int64_t>((n + threads - 1) / threads, 4096));
        gelu_scalar_kernel<Index><<<blocks, threads, 0, stream>>>(
            x, y, static_cast<Index>(n));
    }
}

torch::Tensor gelu_cuda(torch::Tensor input) {
    TORCH_CHECK(input.is_cuda(), "Expected a CUDA tensor");
    TORCH_CHECK(input.scalar_type() == at::kFloat, "Expected float32");

    const c10::cuda::CUDAGuard guard(input.device());
    auto x = input.contiguous();
    auto y = torch::empty_like(x);
    const int64_t n = x.numel();

    if (n == 0)
        return y;

    const cudaStream_t stream = at::cuda::getCurrentCUDAStream();
    if (n < (int64_t(1) << 30)) {
        launch_gelu<int>(
            x.data_ptr<float>(), y.data_ptr<float>(), n, stream);
    } else {
        launch_gelu<int64_t>(
            x.data_ptr<float>(), y.data_ptr<float>(), n, stream);
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return y;
}
"""

_gelu_extension = load_inline(
    name="fused_gelu_streaming_v4",
    cpp_sources="torch::Tensor gelu_cuda(torch::Tensor input);",
    cuda_sources=_cuda_source,
    functions=["gelu_cuda"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3", "--use_fast_math"],
    verbose=False,
)


class ModelNew(nn.Module):
    def __init__(self):
        super().__init__()

    def forward(self, x):
        return _gelu_extension.gelu_cuda(x)
```