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

# Enable cuDNN autotuning so the (fixed-shape) convolution picks the fastest algorithm.
torch.backends.cudnn.benchmark = True

# ---------------------------------------------------------------------------
# Fused HardSwish + ReLU CUDA kernel.
#
# Mathematical simplification used:
#   HardSwish(x) = x * clamp(x + 3, 0, 6) / 6
#   ReLU(HardSwish(x)):
#       - for x <= 0  -> HardSwish(x) <= 0  -> result is 0
#       - for x >  0  -> result is x * min(x + 3, 6) / 6
# So the two elementwise passes collapse into a single cheap fused pass with
# no branches other than the sign test, saving one full global-memory
# read/write round trip and one kernel launch.
# ---------------------------------------------------------------------------
fused_hardswish_relu_source = """
#include <torch/extension.h>
#include <cuda_runtime.h>

__device__ __forceinline__ float hardswish_relu(float x) {
    // ReLU(HardSwish(x)) == (x > 0) ? x * min(x + 3, 6) / 6 : 0
    return (x > 0.0f) ? (x * fminf(x + 3.0f, 6.0f) / 6.0f) : 0.0f;
}

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

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

torch::Tensor fused_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 = xc.numel();
    if (n == 0) {
        return out;
    }

    const int block_size = 256;
    if ((n & 3) == 0) {
        // Vectorized path: 128-bit loads/stores for maximum bandwidth.
        long n4 = n >> 2;
        long blocks = (n4 + block_size - 1) / block_size;
        fused_hardswish_relu_vec4_kernel<<<blocks, block_size>>>(
            reinterpret_cast<const float4*>(xc.data_ptr<float>()),
            reinterpret_cast<float4*>(out.data_ptr<float>()),
            n4);
    } else {
        long blocks = (n + block_size - 1) / block_size;
        fused_hardswish_relu_kernel<<<blocks, block_size>>>(
            xc.data_ptr<float>(), out.data_ptr<float>(), n);
    }
    return out;
}
"""

fused_hardswish_relu_cpp_source = (
    "torch::Tensor fused_hardswish_relu_cuda(torch::Tensor x);"
)

# Compile the inline CUDA fused-activation kernel.
fused_hardswish_relu = load_inline(
    name="fused_hardswish_relu",
    cpp_sources=fused_hardswish_relu_cpp_source,
    cuda_sources=fused_hardswish_relu_source,
    functions=["fused_hardswish_relu_cuda"],
    verbose=True,
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    extra_ldflags=[""],
)


class ModelNew(nn.Module):
    """
    Optimized model: Conv2d is executed by cuDNN (with autotuning), while the
    HardSwish and ReLU activations are fused into a single custom CUDA kernel
    using the algebraic identity ReLU(HardSwish(x)) = (x > 0) ? x*min(x+3,6)/6 : 0.
    """

    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_op = fused_hardswish_relu

    def forward(self, x):
        x = self.conv(x)
        if x.is_cuda and x.dtype == torch.float32:
            return self.fused_op.fused_hardswish_relu_cuda(x)
        # Fallback (e.g. CPU tensors) to reference implementation.
        return torch.relu(F.hardswish(x))
```