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

_cpp_source = """
#include <torch/extension.h>
torch::Tensor gelu_forward(torch::Tensor x);
"""

_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/OpMathType.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 *
              (x + 0.044715f * ((x * x) * x));
    return (0.5f * x) * (1.0f + tanhf(u));
}

__device__ __forceinline__ double gelu(double x) {
    double u = 0.7978845608028654 *
               (x + 0.044715 * ((x * x) * x));
    return (0.5 * x) * (1.0 + tanh(u));
}

__global__ void gelu_float4_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>(gridDim.x) * blockDim.x;
    const int64_t n4 = n / 4;

    const float4* x4 = reinterpret_cast<const float4*>(x);
    float4* y4 = reinterpret_cast<float4*>(y);

    for (int64_t i = tid; i < n4; i += stride) {
        float4 v = x4[i];
        y4[i] = make_float4(
            gelu(v.x), gelu(v.y), gelu(v.z), gelu(v.w));
    }

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

template <typename scalar_t>
__global__ void gelu_scalar_kernel(
    const scalar_t* __restrict__ x,
    scalar_t* __restrict__ y,
    int64_t n
) {
    using acc_t = at::opmath_type<scalar_t>;
    const int64_t stride =
        static_cast<int64_t>(gridDim.x) * blockDim.x;

    for (int64_t i =
             static_cast<int64_t>(blockIdx.x) * blockDim.x + threadIdx.x;
         i < n;
         i += stride) {
        y[i] = static_cast<scalar_t>(gelu(static_cast<acc_t>(x[i])));
    }
}

torch::Tensor gelu_forward(torch::Tensor x) {
    TORCH_CHECK(x.is_cuda(), "Input must be a CUDA tensor");
    TORCH_CHECK(x.is_floating_point(), "Input must be floating point");

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

    if (n == 0) {
        return output;
    }

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

    if (input.scalar_type() == at::kFloat &&
        reinterpret_cast<uintptr_t>(input.data_ptr<float>()) % 16 == 0 &&
        reinterpret_cast<uintptr_t>(output.data_ptr<float>()) % 16 == 0) {
        const int blocks = static_cast<int>(
            std::min<int64_t>((n + threads * 4 - 1) / (threads * 4), 65535));

        gelu_float4_kernel<<<blocks, threads, 0, stream>>>(
            input.data_ptr<float>(), output.data_ptr<float>(), n);
    } else {
        const int blocks = static_cast<int>(
            std::min<int64_t>((n + threads - 1) / threads, 65535));

        AT_DISPATCH_FLOATING_TYPES_AND2(
            at::ScalarType::Half,
            at::ScalarType::BFloat16,
            input.scalar_type(),
            "fused_tanh_gelu",
            [&] {
                gelu_scalar_kernel<scalar_t><<<blocks, threads, 0, stream>>>(
                    input.data_ptr<scalar_t>(),
                    output.data_ptr<scalar_t>(),
                    n);
            });
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_gelu_extension = load_inline(
    name="fused_vectorized_tanh_gelu",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["gelu_forward"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    verbose=False,
)


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

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