```python
# Copyright (c) Meta Platforms, Inc. and affiliates.
# All rights reserved.
#
# This source code is licensed under the license found in the
# LICENSE file in the root directory of this source tree.

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

# Let cuDNN autotune the best conv algorithm for the fixed input shape
torch.backends.cudnn.benchmark = True

# Fused HardSwish + ReLU activation kernel.
# relu(hardswish(x)) = max(x * clamp(x + 3, 0, 6) / 6, 0)
# Vectorized with float4 loads/stores to saturate memory bandwidth.
fused_act_cuda_source = r"""
#include <torch/extension.h>
#include <cuda_runtime.h>

__device__ __forceinline__ float hs_relu(float x) {
    float t = fminf(fmaxf(x + 3.0f, 0.0f), 6.0f);
    float h = x * t * 0.16666667163372039794921875f;
    return fmaxf(h, 0.0f);
}

__global__ void hardswish_relu_vec_kernel(const float4* __restrict__ in,
                                          float4* __restrict__ out,
                                          long n4) {
    long idx = (long)blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n4) {
        float4 v = in[idx];
        v.x = hs_relu(v.x);
        v.y = hs_relu(v.y);
        v.z = hs_relu(v.z);
        v.w = hs_relu(v.w);
        out[idx] = v;
    }
}

__global__ void hardswish_relu_tail_kernel(const float* __restrict__ in,
                                           float* __restrict__ out,
                                           long n) {
    long idx = (long)blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        out[idx] = hs_relu(in[idx]);
    }
}

torch::Tensor hardswish_relu_cuda(torch::Tensor x) {
    TORCH_CHECK(x.is_cuda(), "input must be a CUDA tensor");
    TORCH_CHECK(x.scalar_type() == torch::kFloat32, "input must be float32");

    auto xc = x.contiguous();
    auto out = torch::empty_like(xc);

    long n = (long)xc.numel();
    if (n == 0) {
        return out;
    }

    const int block = 256;
    const float* in_ptr = xc.data_ptr<float>();
    float* out_ptr = out.data_ptr<float>();

    long n4 = n >> 2;
    if (n4 > 0) {
        long grid = (n4 + block - 1) / block;
        hardswish_relu_vec_kernel<<<grid, block>>>(
            reinterpret_cast<const float4*>(in_ptr),
            reinterpret_cast<float4*>(out_ptr),
            n4);
    }

    long rem = n - (n4 << 2);
    if (rem > 0) {
        long grid = (rem + block - 1) / block;
        hardswish_relu_tail_kernel<<<grid, block>>>(
            in_ptr + (n4 << 2), out_ptr + (n4 << 2), rem);
    }

    return out;
}
"""

fused_act_cpp_source = "torch::Tensor hardswish_relu_cuda(torch::Tensor x);"

fused_act = load_inline(
    name="fused_hardswish_relu",
    cpp_sources=fused_act_cpp_source,
    cuda_sources=fused_act_cuda_source,
    functions=["hardswish_relu_cuda"],
    verbose=False,
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3", "--use_fast_math"],
)


class ModelNew(nn.Module):
    """
    Optimized model: cuDNN Conv2d followed by a single fused
    HardSwish + ReLU custom CUDA kernel (one memory pass instead of two).
    """

    def __init__(self, in_channels, out_channels, kernel_size):
        super(ModelNew, self).__init__()
        self.conv = nn.Conv2d(in_channels, out_channels, kernel_size)
        self.fused_act = fused_act

    def forward(self, x):
        """
        Args:
            x (torch.Tensor): Input tensor of shape (batch_size, in_channels, height, width).

        Returns:
            torch.Tensor: Output tensor of shape (batch_size, out_channels, height, width).
        """
        x = self.conv(x)
        x = self.fused_act.hardswish_relu_cuda(x)
        return x


batch_size = 128
in_channels = 8
out_channels = 64
height, width = 128, 128
kernel_size = 3


def get_inputs():
    return [torch.rand(batch_size, in_channels, height, width)]


def get_init_inputs():
    return [in_channels, out_channels, kernel_size]
```