```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 convolution algorithm for the fixed input shape.
torch.backends.cudnn.enabled = True
torch.backends.cudnn.benchmark = True

# Fused activation kernel.
#
# Math note: HardSwish(x) = x * clamp(x + 3, 0, 6) / 6
#            ReLU(HardSwish(x)) == ReLU(x) * clamp(x + 3, 0, 6) / 6
# because HardSwish(x) <= 0 for x <= 0 and >= 0 for x >= 0.
# This lets us fuse the two full elementwise passes into a single pass,
# cutting global memory traffic for the activations roughly in half.
fused_activation_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>

__device__ __forceinline__ float hardswish_relu(float x) {
    float c = fminf(fmaxf(x + 3.0f, 0.0f), 6.0f);
    return fmaxf(x, 0.0f) * c / 6.0f;
}

// Vectorized (float4) kernel: 128-bit loads/stores to saturate bandwidth.
__global__ void hardswish_relu_vec4_kernel(const float4* __restrict__ in,
                                           float4* __restrict__ out,
                                           const int n4) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n4) {
        float4 v = in[idx];
        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[idx] = v;
    }
}

// Scalar kernel to handle the tail when numel is not a multiple of 4.
__global__ void hardswish_relu_tail_kernel(const float* __restrict__ in,
                                           float* __restrict__ out,
                                           const int start,
                                           const int n) {
    int idx = start + threadIdx.x;
    if (idx < n) {
        out[idx] = hardswish_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() == at::kFloat, "input must be float32");

    auto xc = x.contiguous();
    auto out = at::empty_like(xc);
    const int64_t n = xc.numel();
    if (n == 0) {
        return out;
    }

    cudaStream_t stream = at::cuda::getCurrentCUDAStream();
    const int threads = 256;
    const int64_t n4 = n >> 2;

    if (n4 > 0) {
        const int blocks = (int)((n4 + threads - 1) / threads);
        hardswish_relu_vec4_kernel<<<blocks, threads, 0, stream>>>(
            reinterpret_cast<const float4*>(xc.data_ptr<float>()),
            reinterpret_cast<float4*>(out.data_ptr<float>()),
            (int)n4);
    }

    const int rem = (int)(n & 3);
    if (rem != 0) {
        hardswish_relu_tail_kernel<<<1, rem, 0, stream>>>(
            xc.data_ptr<float>(),
            out.data_ptr<float>(),
            (int)(n4 << 2),
            (int)n);
    }

    return out;
}
"""

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

# Compile the fused HardSwish+ReLU CUDA kernel.
fused_activation = load_inline(
    name="fused_hardswish_relu",
    cpp_sources=fused_activation_cpp_source,
    cuda_sources=fused_activation_cuda_source,
    functions=["hardswish_relu_cuda"],
    verbose=False,
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3", "--use_fast_math"],
)


class ModelNew(nn.Module):
    """
    Optimized model:
      - Conv2d stays on the highly tuned cuDNN path (with benchmark autotuning).
      - HardSwish + ReLU are replaced by a single fused custom CUDA kernel,
        using the identity ReLU(HardSwish(x)) = ReLU(x) * clamp(x+3, 0, 6) / 6,
        which removes one full elementwise global-memory pass.
    """
    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_activation = fused_activation

    def forward(self, x):
        x = self.conv(x)
        x = self.fused_activation.hardswish_relu_cuda(x)
        return x
```