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

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

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

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

__global__ void activate_float4(float* __restrict__ x, 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;
    float4* values = reinterpret_cast<float4*>(x);

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

    for (int64_t i = vectors * 4 + tid; i < n; i += stride) {
        x[i] = activate(x[i]);
    }
}

template <typename scalar_t>
__global__ void activate_generic(scalar_t* __restrict__ x, int64_t n) {
    using acc_t = at::acc_type<scalar_t, true>;
    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) {
        // Preserve the intermediate rounding of the original two operators.
        const acc_t value = static_cast<acc_t>(x[i]);
        acc_t gate = value + acc_t(3);
        gate = gate < acc_t(0) ? acc_t(0) : gate;
        gate = gate > acc_t(6) ? acc_t(6) : gate;
        scalar_t hswish = static_cast<scalar_t>(
            (value * gate) / acc_t(6));
        x[i] = static_cast<acc_t>(hswish) < acc_t(0)
            ? scalar_t(0) : hswish;
    }
}

torch::Tensor hardswish_relu_cuda(torch::Tensor x) {
    TORCH_CHECK(x.is_cuda(), "Expected a CUDA tensor");
    TORCH_CHECK(
        x.is_contiguous() ||
        x.is_contiguous(at::MemoryFormat::ChannelsLast),
        "Expected a dense contiguous tensor"
    );

    const c10::cuda::CUDAGuard device_guard(x.device());
    const int64_t n = x.numel();
    if (n == 0) return x;

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

    if (x.scalar_type() == at::kFloat) {
        const int64_t work = (n + 3) / 4;
        const int blocks = static_cast<int>(
            std::min<int64_t>((work + threads - 1) / threads, 4096));
        activate_float4<<<blocks, threads, 0, stream>>>(
            x.data_ptr<float>(), n);
    } else {
        const int blocks = static_cast<int>(
            std::min<int64_t>((n + threads - 1) / threads, 4096));
        AT_DISPATCH_FLOATING_TYPES_AND2(
            at::ScalarType::Half,
            at::ScalarType::BFloat16,
            x.scalar_type(),
            "hardswish_relu_cuda",
            [&] {
                activate_generic<scalar_t>
                    <<<blocks, threads, 0, stream>>>(
                        x.data_ptr<scalar_t>(), n);
            });
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return x;
}
"""

_fused_activation = load_inline(
    name="fused_hardswish_relu_inplace_v1",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["hardswish_relu_cuda"],
    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 or (torch.is_grad_enabled() and x.requires_grad):
            return torch.relu(F.hardswish(x))
        return _fused_activation.hardswish_relu_cuda(x)
```