```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 pack_channels_last_8(torch::Tensor x);
torch::Tensor fused_activation(
    torch::Tensor x, torch::Tensor bias, bool inplace);
"""

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

__global__ void pack8_kernel(
    const float* __restrict__ input,
    float4* __restrict__ output,
    int64_t spatial,
    int64_t tiles_per_image,
    int64_t total_tiles
) {
    __shared__ float tile[8][257];

    const int t = threadIdx.x;
    const int read_channel = t >> 6;
    const int read_pixel = (t & 63) * 4;
    const int write_pixel = t >> 1;
    const int write_channel = (t & 1) * 4;

    for (int64_t task = blockIdx.x;
         task < total_tiles;
         task += gridDim.x) {
        const int64_t n = task / tiles_per_image;
        const int64_t start = (task % tiles_per_image) * 256;

        #pragma unroll
        for (int half = 0; half < 2; ++half) {
            const int c = read_channel + half * 4;
            if (start + read_pixel < spatial) {
                const float4 v = *reinterpret_cast<const float4*>(
                    input + (n * 8 + c) * spatial + start + read_pixel);
                tile[c][read_pixel] = v.x;
                tile[c][read_pixel + 1] = v.y;
                tile[c][read_pixel + 2] = v.z;
                tile[c][read_pixel + 3] = v.w;
            }
        }
        __syncthreads();

        #pragma unroll
        for (int half = 0; half < 2; ++half) {
            const int p = write_pixel + half * 128;
            if (start + p < spatial) {
                output[(n * spatial + start + p) * 2 + (t & 1)] =
                    make_float4(
                        tile[write_channel][p],
                        tile[write_channel + 1][p],
                        tile[write_channel + 2][p],
                        tile[write_channel + 3][p]);
            }
        }
        __syncthreads();
    }
}

torch::Tensor pack_channels_last_8(torch::Tensor x) {
    TORCH_CHECK(x.is_cuda() && x.scalar_type() == at::kFloat,
                "Expected CUDA float32 input");
    TORCH_CHECK(x.dim() == 4 && x.size(1) == 8 && x.is_contiguous(),
                "Expected contiguous NCHW input with eight channels");

    const int64_t spatial = x.size(2) * x.size(3);
    TORCH_CHECK(spatial % 4 == 0, "Spatial size must be divisible by four");
    TORCH_CHECK(
        reinterpret_cast<uintptr_t>(x.data_ptr<float>()) % 16 == 0,
        "Input must be aligned");

    const c10::cuda::CUDAGuard guard(x.device());
    auto output = torch::empty_strided(
        x.sizes(), {8 * spatial, 1, 8 * x.size(3), 8}, x.options());
    if (x.numel() == 0) return output;

    const int64_t tiles = (spatial + 255) / 256;
    const int64_t total_tiles = tiles * x.size(0);
    const int blocks = static_cast<int>(
        std::min<int64_t>(total_tiles, 2048));

    pack8_kernel<<<blocks, 256, 0, at::cuda::getCurrentCUDAStream()>>>(
        x.data_ptr<float>(),
        reinterpret_cast<float4*>(output.data_ptr<float>()),
        spatial, tiles, total_tiles);
    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}

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

__device__ __forceinline__ float4 activate4(float4 v, float4 b) {
    return make_float4(
        activate(v.x, b.x),
        activate(v.y, b.y),
        activate(v.z, b.z),
        activate(v.w, b.w));
}

__device__ __forceinline__ void store_streaming(
    float4* address, float4 value
) {
    asm volatile(
        "st.global.cs.v4.f32 [%0], {%1, %2, %3, %4};"
        :
        : "l"(address), "f"(value.x), "f"(value.y),
          "f"(value.z), "f"(value.w)
        : "memory");
}

__global__ void fused64_kernel(
    const float4* input,
    const float4* __restrict__ bias,
    float4* output,
    int64_t count
) {
    constexpr int threads = 256;
    constexpr int unroll = 4;
    constexpr int tile_size = threads * unroll;

    const float4 b = __ldg(bias + (threadIdx.x & 15));
    const int64_t step = static_cast<int64_t>(gridDim.x) * tile_size;

    for (int64_t base = static_cast<int64_t>(blockIdx.x) * tile_size;
         base < count;
         base += step) {
        float4 values[unroll];

        if (base + tile_size <= count) {
            #pragma unroll
            for (int j = 0; j < unroll; ++j) {
                values[j] = input[base + threadIdx.x + j * threads];
            }

            #pragma unroll
            for (int j = 0; j < unroll; ++j) {
                values[j] = activate4(values[j], b);
            }

            #pragma unroll
            for (int j = 0; j < unroll; ++j) {
                store_streaming(
                    output + base + threadIdx.x + j * threads,
                    values[j]);
            }
        } else {
            #pragma unroll
            for (int j = 0; j < unroll; ++j) {
                const int64_t i = base + threadIdx.x + j * threads;
                if (i < count) {
                    const float4 v = input[i];
                    store_streaming(output + i, activate4(v, b));
                }
            }
        }
    }
}

__global__ void fused_vector_kernel(
    const float4* input,
    const float4* __restrict__ bias,
    float4* output,
    int64_t count,
    int channel_vectors
) {
    const int64_t step = static_cast<int64_t>(blockDim.x) * gridDim.x;
    for (int64_t i = static_cast<int64_t>(blockIdx.x) * blockDim.x
                     + threadIdx.x;
         i < count;
         i += step) {
        const float4 v = input[i];
        const float4 b = __ldg(bias + i % channel_vectors);
        store_streaming(output + i, activate4(v, b));
    }
}

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

torch::Tensor fused_activation(
    torch::Tensor x, torch::Tensor bias, bool inplace
) {
    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 tensors");
    TORCH_CHECK(
        x.dim() == 4 &&
        x.is_contiguous(at::MemoryFormat::ChannelsLast),
        "Expected channels-last input");
    TORCH_CHECK(
        bias.dim() == 1 && bias.is_contiguous() &&
        bias.numel() == x.size(1),
        "Invalid bias");

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

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

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

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

        if (channels == 64) {
            const int blocks = static_cast<int>(
                std::min<int64_t>((vectors + 1023) / 1024, 2048));
            fused64_kernel<<<blocks, threads, 0, stream>>>(
                input_ptr, bias_ptr, output_ptr, vectors);
        } else {
            const int blocks = static_cast<int>(
                std::min<int64_t>((vectors + 255) / 256, 2048));
            fused_vector_kernel<<<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 + 255) / 256, 2048));
        fused_scalar_kernel<<<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;
}
"""

_ops = load_inline(
    name="conv_hardswish_relu_streaming_v6",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["pack_channels_last_8", "fused_activation"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    verbose=False,
)


class _PackChannelsLast8(torch.autograd.Function):
    @staticmethod
    def forward(ctx, x):
        return _ops.pack_channels_last_8(x)

    @staticmethod
    def backward(ctx, grad_output):
        return grad_output


class _BiasHardSwishReLU(torch.autograd.Function):
    @staticmethod
    def forward(ctx, x, bias):
        ctx.save_for_backward(x, bias)
        return _ops.fused_activation(x, bias, False)

    @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):
        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)))

        if not x.is_contiguous(memory_format=torch.channels_last):
            if (
                x.size(1) == 8
                and x.is_contiguous()
                and (x.size(2) * x.size(3)) % 4 == 0
                and x.data_ptr() % 16 == 0
            ):
                x = _PackChannelsLast8.apply(x)
            else:
                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 _BiasHardSwishReLU.apply(y, self.conv.bias)

        return _ops.fused_activation(y, self.conv.bias, True)
```