```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_(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>
#include <cstdint>

__device__ __forceinline__ float activate(float x, float b) {
    x += b;
    if (x <= 0.0f) return 0.0f;
    if (x >= 3.0f) return x;
    return (x * (x + 3.0f)) / 6.0f;
}

__device__ __forceinline__ float4 activate4(float4 v, const float4 b) {
    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);
    return v;
}

__global__ void fused_64_inplace(
    float4* __restrict__ data,
    const float4* __restrict__ bias,
    int64_t count
) {
    constexpr int THREADS = 256;
    constexpr int UNROLL = 4;
    const float4 b = bias[threadIdx.x & 15];
    const int64_t stride = (int64_t)gridDim.x * THREADS * UNROLL;

    for (int64_t base =
             (int64_t)blockIdx.x * THREADS * UNROLL + threadIdx.x;
         base < count; base += stride) {
        if (base + 3 * THREADS < count) {
            float4 v0 = data[base];
            float4 v1 = data[base + THREADS];
            float4 v2 = data[base + 2 * THREADS];
            float4 v3 = data[base + 3 * THREADS];

            v0 = activate4(v0, b);
            v1 = activate4(v1, b);
            v2 = activate4(v2, b);
            v3 = activate4(v3, b);

            data[base] = v0;
            data[base + THREADS] = v1;
            data[base + 2 * THREADS] = v2;
            data[base + 3 * THREADS] = v3;
        } else {
            #pragma unroll
            for (int k = 0; k < UNROLL; ++k) {
                const int64_t i = base + k * THREADS;
                if (i < count)
                    data[i] = activate4(data[i], b);
            }
        }
    }
}

__global__ void fused_vector_inplace(
    float4* __restrict__ data,
    const float4* __restrict__ bias,
    int64_t count,
    int 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) {
        data[i] = activate4(data[i], bias[i % channel_vectors]);
    }
}

__global__ void fused_scalar_inplace(
    float* __restrict__ data,
    const float* __restrict__ bias,
    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) {
        data[i] = activate(data[i], bias[i % channels]);
    }
}

torch::Tensor bias_hardswish_relu_(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());
    const int64_t count = x.numel();
    if (count == 0) return x;

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

    const bool aligned =
        reinterpret_cast<uintptr_t>(x.data_ptr<float>()) % 16 == 0 &&
        reinterpret_cast<uintptr_t>(bias.data_ptr<float>()) % 16 == 0;

    if (channels % 4 == 0 && aligned) {
        const int64_t vectors = count / 4;
        auto* data = reinterpret_cast<float4*>(x.data_ptr<float>());
        const auto* b =
            reinterpret_cast<const float4*>(bias.data_ptr<float>());

        if (channels == 64) {
            const int blocks = static_cast<int>(
                std::min<int64_t>((vectors + 1023) / 1024, 4096));
            fused_64_inplace<<<blocks, threads, 0, stream>>>(
                data, b, vectors);
        } else {
            const int blocks = static_cast<int>(
                std::min<int64_t>((vectors + threads - 1) / threads, 4096));
            fused_vector_inplace<<<blocks, threads, 0, stream>>>(
                data, b, vectors, channels / 4);
        }
    } else {
        const int blocks = static_cast<int>(
            std::min<int64_t>((count + threads - 1) / threads, 4096));
        fused_scalar_inplace<<<blocks, threads, 0, stream>>>(
            x.data_ptr<float>(), bias.data_ptr<float>(), count, channels);
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return x;
}
"""

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


class _BiasHardSwishReLUInplace(torch.autograd.Function):
    @staticmethod
    def forward(ctx, x, bias):
        ctx.mark_dirty(x)
        result = _fused_ops.bias_hardswish_relu_(x, bias)
        if ctx.needs_input_grad[0] or ctx.needs_input_grad[1]:
            ctx.save_for_backward(result)
        return result

    @staticmethod
    def backward(ctx, grad_output):
        (output,) = ctx.saved_tensors
        # For 0 < z <= 3, y = z(z + 3)/6, so
        # dy/dz = sqrt(1/4 + 2y/3).
        slope = torch.where(
            output <= 0,
            torch.zeros_like(output),
            torch.where(
                output <= 3,
                torch.sqrt(output * (2.0 / 3.0) + 0.25),
                torch.ones_like(output),
            ),
        )
        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):
        if (
            not x.is_cuda
            or x.ndim != 4
            or x.dtype != torch.float32
            or self.conv.weight.dtype != torch.float32
            or self.conv.bias is None
            or self.conv.bias.dtype != torch.float32
            or torch.is_autocast_enabled()
        ):
            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 not y.is_contiguous(memory_format=torch.channels_last):
            return F.relu(
                F.hardswish(y + self.conv.bias.view(1, -1, 1, 1))
            )

        if torch.is_grad_enabled() and (
            y.requires_grad or self.conv.bias.requires_grad
        ):
            return _BiasHardSwishReLUInplace.apply(y, self.conv.bias)

        return _fused_ops.bias_hardswish_relu_(y, self.conv.bias)
```