```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) {
    const float x3 = x * x * x;
    return (0.5f * x) *
           (1.0f + tanhf(0.7978845608028654f * (x + 0.044715f * x3)));
}

__global__ void gelu_kernel(
    const float* __restrict__ x,
    float* __restrict__ y,
    int64_t n
) {
    const int64_t tid =
        static_cast<int64_t>(blockIdx.x) * blockDim.x + threadIdx.x;
    const int64_t stride =
        static_cast<int64_t>(blockDim.x) * gridDim.x;
    const int64_t vectors = n / 4;

    const float4* xv = reinterpret_cast<const float4*>(x);
    float4* yv = reinterpret_cast<float4*>(y);

    for (int64_t i = tid; i < vectors; i += stride) {
        float4 v = xv[i];
        v.x = gelu(v.x);
        v.y = gelu(v.y);
        v.z = gelu(v.z);
        v.w = gelu(v.w);
        yv[i] = v;
    }

    const int64_t tail = vectors * 4 + tid;
    if (tail < n) {
        y[tail] = gelu(x[tail]);
    }
}

__global__ void gelu_scalar_kernel(
    const float* __restrict__ x,
    float* __restrict__ y,
    int64_t n
) {
    const int64_t tid =
        static_cast<int64_t>(blockIdx.x) * blockDim.x + threadIdx.x;
    const int64_t stride =
        static_cast<int64_t>(blockDim.x) * gridDim.x;

    for (int64_t i = tid; i < n; i += stride) {
        y[i] = gelu(x[i]);
    }
}

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;

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

    if ((reinterpret_cast<uintptr_t>(x.data_ptr<float>()) & 15) == 0) {
        const int64_t work = (n + 3) / 4;
        const int blocks = static_cast<int>(
            std::min<int64_t>((work + threads - 1) / threads, 4096));
        gelu_kernel<<<blocks, threads, 0, stream>>>(
            x.data_ptr<float>(), y.data_ptr<float>(), n);
    } else {
        const int blocks = static_cast<int>(
            std::min<int64_t>((n + threads - 1) / threads, 4096));
        gelu_scalar_kernel<<<blocks, threads, 0, stream>>>(
            x.data_ptr<float>(), y.data_ptr<float>(), n);
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return y;
}
"""

_gelu_extension = load_inline(
    name="fused_vectorized_tanh_gelu",
    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)
```