```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/cuda/CUDAContext.h>
#include <c10/cuda/CUDAGuard.h>
#include <c10/cuda/CUDAException.h>
#include <cuda_runtime.h>

__device__ __forceinline__ float activate(float x) {
    float y = (x * fminf(fmaxf(x + 3.0f, 0.0f), 6.0f)) / 6.0f;
    return y < 0.0f ? 0.0f : y;
}

__global__ void hardswish_relu_kernel(
    const float* __restrict__ input,
    float* __restrict__ output,
    int64_t count
) {
    const int64_t index =
        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 vector_count = count / 4;

    const float4* input4 = reinterpret_cast<const float4*>(input);
    float4* output4 = reinterpret_cast<float4*>(output);

    for (int64_t i = index; 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 + index; i < count; i += stride) {
        output[i] = activate(input[i]);
    }
}

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

    const c10::cuda::CUDAGuard device_guard(x.device());
    auto output = torch::empty_like(x);
    const int64_t count = x.numel();

    if (count != 0) {
        constexpr int threads = 256;
        const int64_t work = (count + 3) / 4;
        const int64_t needed = (work + threads - 1) / threads;
        const int blocks = static_cast<int>(needed < 4096 ? needed : 4096);

        hardswish_relu_kernel<<<
            blocks, threads, 0, at::cuda::getCurrentCUDAStream()
        >>>(
            x.data_ptr<float>(),
            output.data_ptr<float>(),
            count
        );
        C10_CUDA_KERNEL_LAUNCH_CHECK();
    }

    return output;
}
"""

_fused_ops = load_inline(
    name="conv_hardswish_relu_fused_v1",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["hardswish_relu_cuda"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    verbose=False,
)


class _HardSwishReLU(torch.autograd.Function):
    @staticmethod
    def forward(ctx, x):
        if ctx.needs_input_grad[0]:
            ctx.save_for_backward(x)
        return _fused_ops.hardswish_relu_cuda(x)

    @staticmethod
    def backward(ctx, grad_output):
        (x,) = ctx.saved_tensors
        slope = torch.where(
            x <= 0,
            torch.zeros_like(x),
            torch.where(x <= 3, x / 3 + 0.5, torch.ones_like(x)),
        )
        return grad_output * slope


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 x.is_cuda and x.dtype == torch.float32 and x.is_contiguous():
            return _HardSwishReLU.apply(x)
        return F.relu(F.hardswish(x))
```