```python
import torch
import torch.nn as nn
import torch.nn.functional as F
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>

template <typename T>
__device__ __forceinline__ T activate(T x) {
    if (x <= T(0)) return T(0);
    if (x >= T(3)) return x;
    return x * (x + T(3)) / T(6);
}

__global__ void activate_float4(
    const float* __restrict__ input,
    float* __restrict__ output,
    int64_t count
) {
    const int64_t start =
        static_cast<int64_t>(blockIdx.x) * blockDim.x + threadIdx.x;
    const int64_t stride =
        static_cast<int64_t>(blockDim.x) * gridDim.x;

    const float4* input4 = reinterpret_cast<const float4*>(input);
    float4* output4 = reinterpret_cast<float4*>(output);
    const int64_t vector_count = count / 4;

    for (int64_t i = start; i < vector_count; i += stride) {
        float4 v = input4[i];
        v.x = activate(v.x);
        v.y = activate(v.y);
        v.z = activate(v.z);
        v.w = activate(v.w);
        output4[i] = v;
    }

    for (int64_t i = vector_count * 4 + start; i < count; i += stride) {
        output[i] = activate(input[i]);
    }
}

template <typename scalar_t, typename compute_t>
__global__ void activate_generic(
    const scalar_t* __restrict__ input,
    scalar_t* __restrict__ output,
    int64_t count
) {
    const int64_t start =
        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 = start; i < count; i += stride) {
        output[i] = static_cast<scalar_t>(
            activate(static_cast<compute_t>(input[i]))
        );
    }
}

torch::Tensor fused_hardswish_relu(torch::Tensor input) {
    TORCH_CHECK(input.is_cuda(), "Expected a CUDA tensor");
    TORCH_CHECK(
        input.is_contiguous() ||
        input.is_contiguous(at::MemoryFormat::ChannelsLast),
        "Expected contiguous or channels-last storage"
    );

    const c10::cuda::CUDAGuard guard(input.device());
    auto output = torch::empty_like(input);
    const int64_t count = input.numel();
    if (count == 0) return output;

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

    if (input.scalar_type() == at::kFloat) {
        const int64_t work = (count + 3) / 4;
        const int blocks = static_cast<int>(
            std::min<int64_t>((work + threads - 1) / threads, 4096)
        );
        activate_float4<<<blocks, threads, 0, stream>>>(
            input.data_ptr<float>(), output.data_ptr<float>(), count
        );
    } else {
        const int blocks = static_cast<int>(
            std::min<int64_t>((count + threads - 1) / threads, 4096)
        );
        AT_DISPATCH_FLOATING_TYPES_AND2(
            at::ScalarType::Half,
            at::ScalarType::BFloat16,
            input.scalar_type(),
            "fused_hardswish_relu",
            [&] {
                using compute_t = typename std::conditional<
                    std::is_same<scalar_t, double>::value,
                    double,
                    float
                >::type;
                activate_generic<scalar_t, compute_t>
                    <<<blocks, threads, 0, stream>>>(
                        input.data_ptr<scalar_t>(),
                        output.data_ptr<scalar_t>(),
                        count
                    );
            }
        );
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_fused_ops = load_inline(
    name="conv_hardswish_relu_fused_v1",
    cpp_sources=(
        "torch::Tensor fused_hardswish_relu(torch::Tensor input);"
    ),
    cuda_sources=_cuda_source,
    functions=["fused_hardswish_relu"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    verbose=False,
)


class ModelNew(nn.Module):
    def __init__(self, in_channels, out_channels, kernel_size):
        super().__init__()
        self.conv = nn.Conv2d(in_channels, out_channels, kernel_size)

    def forward(self, x):
        x = self.conv(x)
        if not x.is_cuda:
            return torch.relu(F.hardswish(x))
        return _fused_ops.fused_hardswish_relu(x)
```