```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 bias_hardswish_relu_nhwc(torch::Tensor x, torch::Tensor bias);
"""

_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>

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

template<int FixedChannelVectors>
__global__ void fused_nhwc_vector(
    const float4* __restrict__ input,
    const float4* __restrict__ bias,
    float4* __restrict__ output,
    int64_t count,
    int channel_vectors
) {
    const int cv = FixedChannelVectors > 0
        ? FixedChannelVectors : channel_vectors;
    const int64_t stride = (int64_t)blockDim.x * gridDim.x;

    for (int64_t i = (int64_t)blockIdx.x * blockDim.x + threadIdx.x;
         i < count; i += stride) {
        float4 v = input[i];
        const float4 b = bias[i % cv];
        v.x = activate(v.x, b.x);
        v.y = activate(v.y, b.y);
        v.z = activate(v.z, b.z);
        v.w = activate(v.w, b.w);
        output[i] = v;
    }
}

__global__ void fused_nhwc_scalar(
    const float* __restrict__ input,
    const float* __restrict__ bias,
    float* __restrict__ output,
    int64_t count,
    int channels
) {
    const int64_t stride = (int64_t)blockDim.x * gridDim.x;
    for (int64_t i = (int64_t)blockIdx.x * blockDim.x + threadIdx.x;
         i < count; i += stride) {
        output[i] = activate(input[i], bias[i % channels]);
    }
}

torch::Tensor bias_hardswish_relu_nhwc(
    torch::Tensor x, torch::Tensor bias
) {
    TORCH_CHECK(x.is_cuda() && bias.is_cuda(), "Expected CUDA tensors");
    TORCH_CHECK(x.device() == bias.device(), "Device mismatch");
    TORCH_CHECK(x.scalar_type() == at::kFloat &&
                bias.scalar_type() == at::kFloat, "Expected float32");
    TORCH_CHECK(x.dim() == 4 &&
                x.is_contiguous(at::MemoryFormat::ChannelsLast),
                "Expected channels-last tensor");
    TORCH_CHECK(bias.dim() == 1 && bias.is_contiguous() &&
                bias.numel() == x.size(1), "Invalid bias");

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

    const int channels = static_cast<int>(x.size(1));
    constexpr int threads = 256;
    auto stream = at::cuda::getCurrentCUDAStream();

    if (channels % 4 == 0) {
        const int64_t vectors = count / 4;
        const int blocks = static_cast<int>(
            std::min<int64_t>((vectors + threads - 1) / threads, 4096));
        const float4* input_ptr =
            reinterpret_cast<const float4*>(x.data_ptr<float>());
        const float4* bias_ptr =
            reinterpret_cast<const float4*>(bias.data_ptr<float>());
        float4* output_ptr =
            reinterpret_cast<float4*>(output.data_ptr<float>());

        if (channels == 64) {
            fused_nhwc_vector<16><<<blocks, threads, 0, stream>>>(
                input_ptr, bias_ptr, output_ptr, vectors, 16);
        } else {
            fused_nhwc_vector<0><<<blocks, threads, 0, stream>>>(
                input_ptr, bias_ptr, output_ptr, vectors, channels / 4);
        }
    } else {
        const int blocks = static_cast<int>(
            std::min<int64_t>((count + threads - 1) / threads, 4096));
        fused_nhwc_scalar<<<blocks, threads, 0, stream>>>(
            x.data_ptr<float>(), bias.data_ptr<float>(),
            output.data_ptr<float>(), count, channels);
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

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


class _BiasHardSwishReLU(torch.autograd.Function):
    @staticmethod
    def forward(ctx, x, bias):
        if ctx.needs_input_grad[0] or ctx.needs_input_grad[1]:
            ctx.save_for_backward(x, bias)
        return _fused_ops.bias_hardswish_relu_nhwc(x, bias)

    @staticmethod
    def backward(ctx, grad_output):
        x, bias = ctx.saved_tensors
        z = x + bias.view(1, -1, 1, 1)
        slope = torch.where(
            z <= 0,
            torch.zeros_like(z),
            torch.where(z <= 3, z / 3 + 0.5, torch.ones_like(z)),
        )
        grad = grad_output * slope
        return (
            grad if ctx.needs_input_grad[0] else None,
            grad.sum((0, 2, 3)) if ctx.needs_input_grad[1] else None,
        )


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)
        self.conv.to(memory_format=torch.channels_last)

    def forward(self, x):
        use_custom = (
            x.is_cuda
            and x.ndim == 4
            and x.dtype == torch.float32
            and self.conv.weight.dtype == torch.float32
            and self.conv.bias.dtype == torch.float32
            and not torch.is_autocast_enabled()
        )
        if not use_custom:
            return F.relu(F.hardswish(self.conv(x)))

        x = x.contiguous(memory_format=torch.channels_last)
        y = F.conv2d(
            x,
            self.conv.weight,
            None,
            self.conv.stride,
            self.conv.padding,
            self.conv.dilation,
            self.conv.groups,
        )
        if y.is_contiguous(memory_format=torch.channels_last):
            return _BiasHardSwishReLU.apply(y, self.conv.bias)

        return F.relu(
            F.hardswish(y + self.conv.bias.view(1, -1, 1, 1))
        )
```