You write custom CUDA kernels to replace the pytorch operators in the given architecture to get speedups. 

    You have complete freedom to choose the set of operators you want to replace. You may make the decision to replace some operators with custom CUDA kernels and leave others unchanged. You may replace multiple operators with custom implementations, consider operator fusion opportunities (combining multiple operators into a single kernel, for example, combining matmul+relu), or algorithmic changes (such as online softmax). You are only limited by your imagination.


        Here's an example to show you the syntax of inline embedding custom CUDA operators in torch: The example given architecture is: 

        ``` 

        # 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


class Model(nn.Module):
    def __init__(self) -> None:
        super().__init__()

    def forward(self, a, b):
        return a + b


def get_inputs():
    # randomly generate input tensors based on the model architecture
    a = torch.randn(1, 128).cuda()
    b = torch.randn(1, 128).cuda()
    return [a, b]


def get_init_inputs():
    # randomly generate tensors required for initialization based on the model architecture
    return []

        ``` 

        The example new arch with custom CUDA kernels looks like this:
        ```
        # 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

# Define the custom CUDA kernel for element-wise addition
elementwise_add_source = """
#include <torch/extension.h>
#include <cuda_runtime.h>

__global__ void elementwise_add_kernel(const float* a, const float* b, float* out, int size) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < size) {
        out[idx] = a[idx] + b[idx];
    }
}

torch::Tensor elementwise_add_cuda(torch::Tensor a, torch::Tensor b) {
    auto size = a.numel();
    auto out = torch::zeros_like(a);

    const int block_size = 256;
    const int num_blocks = (size + block_size - 1) / block_size;

    elementwise_add_kernel<<<num_blocks, block_size>>>(a.data_ptr<float>(), b.data_ptr<float>(), out.data_ptr<float>(), size);

    return out;
}
"""

elementwise_add_cpp_source = (
    "torch::Tensor elementwise_add_cuda(torch::Tensor a, torch::Tensor b);"
)

# Compile the inline CUDA code for element-wise addition
elementwise_add = load_inline(
    name="elementwise_add",
    cpp_sources=elementwise_add_cpp_source,
    cuda_sources=elementwise_add_source,
    functions=["elementwise_add_cuda"],
    verbose=True,
    extra_cflags=[""],
    extra_ldflags=[""],
)


class ModelNew(nn.Module):
    def __init__(self) -> None:
        super().__init__()
        self.elementwise_add = elementwise_add

    def forward(self, a, b):
        return self.elementwise_add.elementwise_add_cuda(a, b)

        ``` 

        
    You are given the following architecture: 

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

class Model(nn.Module):
    """
    Model that performs a convolution, applies HardSwish, and then ReLU.
    """
    def __init__(self, in_channels, out_channels, kernel_size):
        super(Model, self).__init__()
        self.conv = nn.Conv2d(in_channels, out_channels, kernel_size)

    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 = torch.nn.functional.hardswish(x)
        x = torch.relu(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]
    ```
    
Optimize the architecture named Model with custom CUDA operators! Name your optimized output architecture ModelNew. Output the new code in codeblocks. Please generate real code, NOT pseudocode, make sure the code compiles and is fully functional. Just output the new model code, no other text, and NO testing code! 



This is turn 4 of 6 for this task. Your previous solution (turn 3) was:
```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

# Fused Conv2d + HardSwish + ReLU.
# Fast path: implicit-GEMM tensor-core kernel specialized for C=8, K=3x3 using
# "3xTF32" (split each fp32 into hi+lo tf32 and issue 3 mma's). This gives
# tensor-core throughput with ~fp32 accuracy (error ~1e-6, far below the ~2e-4
# error of plain TF32), so it passes tight correctness checks.
# Each block computes a 128(ow) x 64(oc) output tile with 4 warps via 16x16x8
# tf32 wmma fragments; the epilogue applies bias + HardSwish + ReLU.
# Fallbacks: direct FP32 fused SIMT kernel (any shape/arch), then ATen.
# The TC path is only enabled after a one-time runtime self-check against a
# strict-fp32 reference; on any mismatch it silently stays on the FP32 kernel.
_fused_cpp_src = (
    "torch::Tensor fused_conv_act(torch::Tensor input, torch::Tensor weight, "
    "torch::Tensor bias, int64_t stride, int64_t pad, bool use_tc);"
)

_fused_cuda_src = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <c10/cuda/CUDAException.h>
#include <cuda_runtime.h>
#include <mma.h>

#define OC_TILE 8
#define BM 128
#define BN 64
#define BK 24
#define CLD (BN + 4)

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

// Split fp32 into (hi, lo) tf32 bit patterns with x ~= hi + lo.
__device__ __forceinline__ void split_tf32(float x, float& hi, float& lo) {
#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 800)
    unsigned int h, l;
    asm("cvt.rna.tf32.f32 %0, %1;\n" : "=r"(h) : "f"(x));
    float xh = __uint_as_float(h);
    asm("cvt.rna.tf32.f32 %0, %1;\n" : "=r"(l) : "f"(x - xh));
    hi = xh;
    lo = __uint_as_float(l);
#else
    hi = x;
    lo = 0.0f;
#endif
}

// ---------------- Tensor-core (3xTF32) fused conv + HardSwish + ReLU -----------
template <int C, int KH, int KW>
__global__ void __launch_bounds__(BM) conv_act_tc_kernel(
    const float* __restrict__ input,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ output,
    int N, int H, int W, int OC,
    int OH, int OW, int stride, int pad, int octiles)
{
#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 800)
    constexpr int K = C * KH * KW;
    constexpr int KK = KH * KW;
    extern __shared__ __align__(32) float smem[];
    float* As_hi = smem;                 // [BK][BM] col-major A tile (hi)
    float* As_lo = As_hi + BK * BM;      // [BK][BM] (lo)
    float* Bs_hi = As_lo + BK * BM;      // [BN][BK] col-major B tile (hi)
    float* Bs_lo = Bs_hi + BN * BK;      // [BN][BK] (lo)
    float* Cs    = smem;                 // [BM][CLD] reused after the k-loop

    const int tid  = threadIdx.x;        // 128 threads = 4 warps
    const int warp = tid >> 5;

    const int n   = blockIdx.z / octiles;
    const int oct = blockIdx.z - n * octiles;
    const int oc0 = oct * BN;
    const int oh  = blockIdx.y;
    const int ow0 = blockIdx.x * BM;
    const int ow  = ow0 + tid;

    namespace wmma = nvcuda::wmma;
    wmma::fragment<wmma::accumulator, 16, 16, 8, float> acc[2][4];
    #pragma unroll
    for (int i = 0; i < 2; ++i)
        #pragma unroll
        for (int j = 0; j < 4; ++j)
            wmma::fill_fragment(acc[i][j], 0.0f);

    const float* in_n = input + (long)n * (C * H * W);

    #pragma unroll
    for (int k0 = 0; k0 < K; k0 += BK) {
        __syncthreads();
        // Stage A tile (col-major): fully unrolled -> k is compile-time.
        #pragma unroll
        for (int kk = 0; kk < BK; ++kk) {
            const int k  = k0 + kk;
            const int c  = k / KK;
            const int r  = k - c * KK;
            const int kh = r / KW;
            const int kw = r - kh * KW;
            const int ih = oh * stride + kh - pad;
            const int iw = ow * stride + kw - pad;
            float v = 0.0f;
            if (ih >= 0 && ih < H && iw >= 0 && iw < W)
                v = in_n[(c * H + ih) * W + iw];
            float vh, vl;
            split_tf32(v, vh, vl);
            As_hi[kk * BM + tid] = vh;
            As_lo[kk * BM + tid] = vl;
        }
        // Stage B tile (col-major): [BN][BK].
        #pragma unroll
        for (int i = 0; i < (BN * BK) / BM; ++i) {
            const int idx = tid + i * BM;
            const int ocl = idx / BK;
            const int kk  = idx - ocl * BK;
            const int k   = k0 + kk;
            const int oc  = oc0 + ocl;
            float v = 0.0f;
            if (oc < OC) v = weight[(long)oc * K + k];
            float vh, vl;
            split_tf32(v, vh, vl);
            Bs_hi[idx] = vh;
            Bs_lo[idx] = vl;
        }
        __syncthreads();
        #pragma unroll
        for (int kk = 0; kk < BK; kk += 8) {
            wmma::fragment<wmma::matrix_a, 16, 16, 8, wmma::precision::tf32, wmma::col_major> a0h, a1h, a0l, a1l;
            wmma::load_matrix_sync(a0h, As_hi + kk * BM + warp * 32, BM);
            wmma::load_matrix_sync(a1h, As_hi + kk * BM + warp * 32 + 16, BM);
            wmma::load_matrix_sync(a0l, As_lo + kk * BM + warp * 32, BM);
            wmma::load_matrix_sync(a1l, As_lo + kk * BM + warp * 32 + 16, BM);
            #pragma unroll
            for (int j = 0; j < 4; ++j) {
                wmma::fragment<wmma::matrix_b, 16, 16, 8, wmma::precision::tf32, wmma::col_major> bh, bl;
                wmma::load_matrix_sync(bh, Bs_hi + j * 16 * BK + kk, BK);
                wmma::load_matrix_sync(bl, Bs_lo + j * 16 * BK + kk, BK);
                // 3xTF32: hi*hi + hi*lo + lo*hi (drop lo*lo, ~2^-22 relative).
                wmma::mma_sync(acc[0][j], a0h, bh, acc[0][j]);
                wmma::mma_sync(acc[0][j], a0h, bl, acc[0][j]);
                wmma::mma_sync(acc[0][j], a0l, bh, acc[0][j]);
                wmma::mma_sync(acc[1][j], a1h, bh, acc[1][j]);
                wmma::mma_sync(acc[1][j], a1h, bl, acc[1][j]);
                wmma::mma_sync(acc[1][j], a1l, bh, acc[1][j]);
            }
        }
    }
    __syncthreads();
    // Stage C tile in shared (padded leading dim avoids bank conflicts).
    #pragma unroll
    for (int i = 0; i < 2; ++i)
        #pragma unroll
        for (int j = 0; j < 4; ++j)
            wmma::store_matrix_sync(Cs + (warp * 32 + i * 16) * CLD + j * 16,
                                    acc[i][j], CLD, wmma::mem_row_major);
    __syncthreads();
    // Fused epilogue: bias + HardSwish + ReLU, coalesced global writes.
    if (ow < OW) {
        float* out_p = output + (long)n * OC * OH * OW + (long)oh * OW + ow;
        const float* cs_p = Cs + tid * CLD;
        #pragma unroll 4
        for (int j = 0; j < BN; ++j) {
            const int oc = oc0 + j;
            if (oc < OC) {
                out_p[(long)oc * OH * OW] = relu_hardswish(cs_p[j] + bias[oc]);
            }
        }
    }
#endif
}

// ---------------- Generic direct FP32 fallback (any shape/arch) -----------------
template <int C, int K>
__global__ void conv_hswish_relu_kernel(
    const float* __restrict__ input,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ output,
    int N, int H, int W, int OC,
    int OH, int OW, int stride, int pad)
{
    constexpr int KSIZE = C * K * K;
    __shared__ float s_weight[OC_TILE * KSIZE];

    const int tid = threadIdx.y * blockDim.x + threadIdx.x;
    const int nthreads = blockDim.x * blockDim.y;
    const int oc0 = blockIdx.y * OC_TILE;

    for (int i = tid; i < OC_TILE * KSIZE; i += nthreads) {
        int oc = oc0 + i / KSIZE;
        int w = i % KSIZE;
        s_weight[i] = (oc < OC) ? weight[oc * KSIZE + w] : 0.0f;
    }
    __syncthreads();

    const int tilesX = (OW + (int)blockDim.x - 1) / (int)blockDim.x;
    const int tileY = blockIdx.x / tilesX;
    const int tileX = blockIdx.x % tilesX;
    const int oh = tileY * (int)blockDim.y + (int)threadIdx.y;
    const int ow = tileX * (int)blockDim.x + (int)threadIdx.x;
    if (oh >= OH || ow >= OW) return;

    const int n = blockIdx.z;
    const float* in_n = input + (long)n * (C * H * W);
    const int ih0 = oh * stride - pad;
    const int iw0 = ow * stride - pad;

    float acc[OC_TILE];
    #pragma unroll
    for (int i = 0; i < OC_TILE; ++i) acc[i] = 0.0f;

    #pragma unroll
    for (int c = 0; c < C; ++c) {
        const float* in_c = in_n + c * (H * W);
        #pragma unroll
        for (int kh = 0; kh < K; ++kh) {
            const int ih = ih0 + kh;
            const bool hok = (ih >= 0) && (ih < H);
            #pragma unroll
            for (int kw = 0; kw < K; ++kw) {
                const int iw = iw0 + kw;
                float v = (hok && iw >= 0 && iw < W) ? in_c[ih * W + iw] : 0.0f;
                const int woff = (c * K + kh) * K + kw;
                #pragma unroll
                for (int i = 0; i < OC_TILE; ++i) {
                    acc[i] = fmaf(v, s_weight[i * KSIZE + woff], acc[i]);
                }
            }
        }
    }

    float* out_n = output + (long)n * OC * OH * OW + oh * OW + ow;
    #pragma unroll
    for (int i = 0; i < OC_TILE; ++i) {
        const int oc = oc0 + i;
        if (oc < OC) {
            out_n[(long)oc * OH * OW] = relu_hardswish(acc[i] + bias[oc]);
        }
    }
}

__global__ void conv_hswish_relu_generic_kernel(
    const float* __restrict__ input,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ output,
    int N, int C, int H, int W, int OC,
    int KH, int KW, int OH, int OW, int stride, int pad)
{
    extern __shared__ float s_weight[];
    const int ksize = C * KH * KW;
    const int tid = threadIdx.y * blockDim.x + threadIdx.x;
    const int nthreads = blockDim.x * blockDim.y;
    const int oc0 = blockIdx.y * OC_TILE;

    for (int i = tid; i < OC_TILE * ksize; i += nthreads) {
        int oc = oc0 + i / ksize;
        int w = i % ksize;
        s_weight[i] = (oc < OC) ? weight[oc * ksize + w] : 0.0f;
    }
    __syncthreads();

    const int tilesX = (OW + (int)blockDim.x - 1) / (int)blockDim.x;
    const int tileY = blockIdx.x / tilesX;
    const int tileX = blockIdx.x % tilesX;
    const int oh = tileY * (int)blockDim.y + (int)threadIdx.y;
    const int ow = tileX * (int)blockDim.x + (int)threadIdx.x;
    if (oh >= OH || ow >= OW) return;

    const int n = blockIdx.z;
    const float* in_n = input + (long)n * C * H * W;
    const int ih0 = oh * stride - pad;
    const int iw0 = ow * stride - pad;

    float acc[OC_TILE];
    #pragma unroll
    for (int i = 0; i < OC_TILE; ++i) acc[i] = 0.0f;

    for (int c = 0; c < C; ++c) {
        const float* in_c = in_n + (long)c * H * W;
        for (int kh = 0; kh < KH; ++kh) {
            const int ih = ih0 + kh;
            const bool hok = (ih >= 0) && (ih < H);
            for (int kw = 0; kw < KW; ++kw) {
                const int iw = iw0 + kw;
                float v = (hok && iw >= 0 && iw < W) ? in_c[ih * W + iw] : 0.0f;
                const int woff = (c * KH + kh) * KW + kw;
                #pragma unroll
                for (int i = 0; i < OC_TILE; ++i) {
                    acc[i] = fmaf(v, s_weight[i * ksize + woff], acc[i]);
                }
            }
        }
    }

    float* out_n = output + (long)n * OC * OH * OW + oh * OW + ow;
    #pragma unroll
    for (int i = 0; i < OC_TILE; ++i) {
        const int oc = oc0 + i;
        if (oc < OC) {
            out_n[(long)oc * OH * OW] = relu_hardswish(acc[i] + bias[oc]);
        }
    }
}

torch::Tensor fused_conv_act(torch::Tensor input, torch::Tensor weight, torch::Tensor bias,
                             int64_t stride, int64_t pad, bool use_tc)
{
    TORCH_CHECK(input.is_cuda() && weight.is_cuda() && bias.is_cuda(), "tensors must be CUDA");
    TORCH_CHECK(input.scalar_type() == at::kFloat, "only float32 supported");
    TORCH_CHECK(input.dim() == 4 && weight.dim() == 4, "bad ranks");

    const long N  = input.size(0);
    const long C  = input.size(1);
    const long H  = input.size(2);
    const long W  = input.size(3);
    const long OC = weight.size(0);
    const long KH = weight.size(2);
    const long KW = weight.size(3);
    const long OH = (H + 2 * pad - KH) / stride + 1;
    const long OW = (W + 2 * pad - KW) / stride + 1;

    auto output = at::empty({N, OC, OH, OW}, input.options());
    if (output.numel() == 0 || N == 0) return output;

    cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    // Tensor-core 3xTF32 fast path.
    const long octiles = (OC + BN - 1) / BN;
    const long tilesX  = (OW + BM - 1) / BM;
    if (use_tc && C == 8 && KH == 3 && KW == 3 &&
        at::cuda::getCurrentDeviceProperties()->major >= 8 &&
        OH <= 65535 && N * octiles <= 65535) {
        const dim3 grid((unsigned)tilesX, (unsigned)OH, (unsigned)(N * octiles));
        constexpr size_t shmem = (size_t)(2 * BK * BM + 2 * BN * BK) * sizeof(float);
        conv_act_tc_kernel<8, 3, 3><<<grid, BM, shmem, stream>>>(
            input.data_ptr<float>(), weight.data_ptr<float>(), bias.data_ptr<float>(),
            output.data_ptr<float>(), (int)N, (int)H, (int)W, (int)OC,
            (int)OH, (int)OW, (int)stride, (int)pad, (int)octiles);
        C10_CUDA_KERNEL_LAUNCH_CHECK();
        return output;
    }

    // Direct FP32 fallback.
    const dim3 block(32, 8);
    const long tx = (OW + 31) / 32;
    const long ty = (OH + 7) / 8;
    if (tx * ty <= 65535 && N <= 65535 &&
        (size_t)OC_TILE * C * KH * KW * sizeof(float) <= 48 * 1024) {
        const dim3 grid((unsigned)(tx * ty), (unsigned)((OC + OC_TILE - 1) / OC_TILE), (unsigned)N);
        if (KH == 3 && KW == 3 && C == 8) {
            conv_hswish_relu_kernel<8, 3><<<grid, block, 0, stream>>>(
                input.data_ptr<float>(), weight.data_ptr<float>(), bias.data_ptr<float>(),
                output.data_ptr<float>(), (int)N, (int)H, (int)W, (int)OC,
                (int)OH, (int)OW, (int)stride, (int)pad);
        } else {
            const size_t shmem = (size_t)OC_TILE * C * KH * KW * sizeof(float);
            conv_hswish_relu_generic_kernel<<<grid, block, shmem, stream>>>(
                input.data_ptr<float>(), weight.data_ptr<float>(), bias.data_ptr<float>(),
                output.data_ptr<float>(), (int)N, (int)C, (int)H, (int)W, (int)OC,
                (int)KH, (int)KW, (int)OH, (int)OW, (int)stride, (int)pad);
        }
        C10_CUDA_KERNEL_LAUNCH_CHECK();
        return output;
    }

    // Last-resort ATen fallback (identical semantics).
    auto y = at::conv2d(input, weight, bias, {stride, stride}, {pad, pad});
    y = at::hardswish(y);
    y = at::relu(y);
    output.copy_(y);
    return output;
}
"""

fused_conv_act_mod = load_inline(
    name="fused_conv_hardswish_relu_v3",
    cpp_sources=_fused_cpp_src,
    cuda_sources=_fused_cuda_src,
    functions=["fused_conv_act"],
    extra_cuda_cflags=["-O3"],
    verbose=False,
)


class ModelNew(nn.Module):
    """
    Optimized model: Conv2d + HardSwish + ReLU fused into a single custom CUDA kernel.
    Fast path uses tensor cores with 3xTF32 (near-fp32 accuracy, enabled only after a
    one-time runtime self-check against a strict fp32 reference). Falls back to a
    direct FP32 fused kernel (or the original PyTorch ops) for other shapes/devices.
    Keeps the same nn.Conv2d submodule so state_dict / initialization is identical
    to the reference Model.
    """
    def __init__(self, in_channels, out_channels, kernel_size):
        super(ModelNew, self).__init__()
        self.conv = nn.Conv2d(in_channels, out_channels, kernel_size)
        kh, kw = self.conv.kernel_size
        self._ok = (self.conv.groups == 1 and
                    self.conv.dilation == (1, 1) and
                    self.conv.stride[0] == self.conv.stride[1] and
                    self.conv.padding[0] == self.conv.padding[1] and
                    self.conv.bias is not None and
                    (in_channels * kh * kw * 8 * 4) <= 48 * 1024)
        self._use_tc = False
        self._checked = False
        if torch.cuda.is_available():
            self._try_check()

    @torch.no_grad()
    def _try_check(self):
        self._checked = True
        self._use_tc = False
        try:
            ic = self.conv.in_channels
            oc = self.conv.out_channels
            kh, kw = self.conv.kernel_size
            if not (ic == 8 and kh == 3 and kw == 3):
                return
            if not torch.cuda.is_available():
                return
            if torch.cuda.get_device_capability(0)[0] < 8:
                return
            dev = torch.device('cuda')
            tmp = nn.Conv2d(ic, oc, kh,
                            stride=self.conv.stride,
                            padding=self.conv.padding).to(dev)
            x = torch.rand(2, ic, 34, 34, device=dev)
            prev = torch.backends.cudnn.allow_tf32
            try:
                torch.backends.cudnn.allow_tf32 = False
                ref = F.relu(F.hardswish(F.conv2d(x, tmp.weight, tmp.bias,
                                                  tmp.stride, tmp.padding)))
            finally:
                torch.backends.cudnn.allow_tf32 = prev
            tc = fused_conv_act_mod.fused_conv_act(
                x, tmp.weight, tmp.bias, tmp.stride[0], tmp.padding[0], True)
            # Accept the TC path only if it is numerically ~fp32-exact.
            self._use_tc = bool(torch.allclose(tc, ref, rtol=1e-4, atol=1e-5))
        except Exception:
            self._use_tc = False

    def forward(self, x):
        if self._ok and x.is_cuda and x.dtype == torch.float32 and x.dim() == 4:
            if not self._checked:
                self._try_check()
            return fused_conv_act_mod.fused_conv_act(
                x.contiguous(),
                self.conv.weight.contiguous(),
                self.conv.bias.contiguous(),
                self.conv.stride[0],
                self.conv.padding[0],
                self._use_tc,
            )
        x = self.conv(x)
        x = F.hardswish(x)
        x = torch.relu(x)
        return x
```
Result of evaluating it:
The solution is correct on the test inputs. Mean runtime 2.4 ms; the PyTorch reference (with TF32 enabled) takes 4.84 ms; speedup 2.017x.
Profiler, top device kernels of one forward pass of your solution:
| device kernel | total device time (us) | calls |
|---|---|---|
| void conv_act_tc_kernel<8, 3, 3>(float const*, float const*, float const*, float*, int, in | 2382.1 | 1 |

The fastest correct solution so far (2.495x speedup) is:
```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

# Fused CUDA kernel: Conv2d (stride=1, direct) + HardSwish + ReLU in a single pass.
# Key ideas:
#  - Each thread computes OC_TILE output channels for one spatial location, so each
#    loaded input value is reused OC_TILE times (higher arithmetic intensity).
#  - The (tiny) filter weights are staged in shared memory and broadcast to all threads.
#  - HardSwish (x * clamp(x+3,0,6) / 6) and the trailing ReLU are fused into the
#    epilogue, eliminating two full element-wise passes over the output tensor.
_fused_cpp_src = "torch::Tensor fused_conv_act(torch::Tensor input, torch::Tensor weight, torch::Tensor bias, int64_t stride, int64_t pad);"

_fused_cuda_src = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <c10/cuda/CUDAException.h>
#include <cuda_runtime.h>

#define OC_TILE 8

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

// Specialized kernel: compile-time channel count / kernel size (fully unrolled).
template <int C, int K>
__global__ void conv_hswish_relu_kernel(
    const float* __restrict__ input,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ output,
    int N, int H, int W, int OC,
    int OH, int OW, int stride, int pad)
{
    constexpr int KSIZE = C * K * K;
    __shared__ float s_weight[OC_TILE * KSIZE];

    const int tid = threadIdx.y * blockDim.x + threadIdx.x;
    const int nthreads = blockDim.x * blockDim.y;
    const int oc0 = blockIdx.y * OC_TILE;

    for (int i = tid; i < OC_TILE * KSIZE; i += nthreads) {
        int oc = oc0 + i / KSIZE;
        int w = i % KSIZE;
        s_weight[i] = (oc < OC) ? weight[oc * KSIZE + w] : 0.0f;
    }
    __syncthreads();

    const int tilesX = (OW + (int)blockDim.x - 1) / (int)blockDim.x;
    const int tileY = blockIdx.x / tilesX;
    const int tileX = blockIdx.x % tilesX;
    const int oh = tileY * (int)blockDim.y + (int)threadIdx.y;
    const int ow = tileX * (int)blockDim.x + (int)threadIdx.x;
    if (oh >= OH || ow >= OW) return;

    const int n = blockIdx.z;
    const float* in_n = input + (long)n * (C * H * W);
    const int ih0 = oh * stride - pad;
    const int iw0 = ow * stride - pad;

    float acc[OC_TILE];
    #pragma unroll
    for (int i = 0; i < OC_TILE; ++i) acc[i] = 0.0f;

    #pragma unroll
    for (int c = 0; c < C; ++c) {
        const float* in_c = in_n + c * (H * W);
        #pragma unroll
        for (int kh = 0; kh < K; ++kh) {
            const int ih = ih0 + kh;
            const bool hok = (ih >= 0) && (ih < H);
            #pragma unroll
            for (int kw = 0; kw < K; ++kw) {
                const int iw = iw0 + kw;
                float v = (hok && iw >= 0 && iw < W) ? in_c[ih * W + iw] : 0.0f;
                const int woff = (c * K + kh) * K + kw;
                #pragma unroll
                for (int i = 0; i < OC_TILE; ++i) {
                    acc[i] = fmaf(v, s_weight[i * KSIZE + woff], acc[i]);
                }
            }
        }
    }

    float* out_n = output + (long)n * OC * OH * OW + oh * OW + ow;
    #pragma unroll
    for (int i = 0; i < OC_TILE; ++i) {
        const int oc = oc0 + i;
        if (oc < OC) {
            out_n[(long)oc * OH * OW] = relu_hardswish(acc[i] + bias[oc]);
        }
    }
}

// Generic fallback kernel: runtime channel count / kernel size.
__global__ void conv_hswish_relu_generic_kernel(
    const float* __restrict__ input,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ output,
    int N, int C, int H, int W, int OC,
    int KH, int KW, int OH, int OW, int stride, int pad)
{
    extern __shared__ float s_weight[];
    const int ksize = C * KH * KW;
    const int tid = threadIdx.y * blockDim.x + threadIdx.x;
    const int nthreads = blockDim.x * blockDim.y;
    const int oc0 = blockIdx.y * OC_TILE;

    for (int i = tid; i < OC_TILE * ksize; i += nthreads) {
        int oc = oc0 + i / ksize;
        int w = i % ksize;
        s_weight[i] = (oc < OC) ? weight[oc * ksize + w] : 0.0f;
    }
    __syncthreads();

    const int tilesX = (OW + (int)blockDim.x - 1) / (int)blockDim.x;
    const int tileY = blockIdx.x / tilesX;
    const int tileX = blockIdx.x % tilesX;
    const int oh = tileY * (int)blockDim.y + (int)threadIdx.y;
    const int ow = tileX * (int)blockDim.x + (int)threadIdx.x;
    if (oh >= OH || ow >= OW) return;

    const int n = blockIdx.z;
    const float* in_n = input + (long)n * C * H * W;
    const int ih0 = oh * stride - pad;
    const int iw0 = ow * stride - pad;

    float acc[OC_TILE];
    #pragma unroll
    for (int i = 0; i < OC_TILE; ++i) acc[i] = 0.0f;

    for (int c = 0; c < C; ++c) {
        const float* in_c = in_n + (long)c * H * W;
        for (int kh = 0; kh < KH; ++kh) {
            const int ih = ih0 + kh;
            const bool hok = (ih >= 0) && (ih < H);
            for (int kw = 0; kw < KW; ++kw) {
                const int iw = iw0 + kw;
                float v = (hok && iw >= 0 && iw < W) ? in_c[ih * W + iw] : 0.0f;
                const int woff = (c * KH + kh) * KW + kw;
                #pragma unroll
                for (int i = 0; i < OC_TILE; ++i) {
                    acc[i] = fmaf(v, s_weight[i * ksize + woff], acc[i]);
                }
            }
        }
    }

    float* out_n = output + (long)n * OC * OH * OW + oh * OW + ow;
    #pragma unroll
    for (int i = 0; i < OC_TILE; ++i) {
        const int oc = oc0 + i;
        if (oc < OC) {
            out_n[(long)oc * OH * OW] = relu_hardswish(acc[i] + bias[oc]);
        }
    }
}

torch::Tensor fused_conv_act(torch::Tensor input, torch::Tensor weight, torch::Tensor bias, int64_t stride, int64_t pad)
{
    TORCH_CHECK(input.is_cuda() && weight.is_cuda() && bias.is_cuda(), "tensors must be CUDA");
    TORCH_CHECK(input.scalar_type() == at::kFloat, "only float32 supported");

    const int N  = (int)input.size(0);
    const int C  = (int)input.size(1);
    const int H  = (int)input.size(2);
    const int W  = (int)input.size(3);
    const int OC = (int)weight.size(0);
    const int KH = (int)weight.size(2);
    const int KW = (int)weight.size(3);
    const int OH = (int)((H + 2 * pad - KH) / stride + 1);
    const int OW = (int)((W + 2 * pad - KW) / stride + 1);

    auto output = at::empty({N, OC, OH, OW}, input.options());
    if (output.numel() == 0) return output;

    const dim3 block(32, 8);
    const int tilesX = (OW + 31) / 32;
    const int tilesY = (OH + 7) / 8;
    const dim3 grid(tilesX * tilesY, (OC + OC_TILE - 1) / OC_TILE, N);

    cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    if (KH == 3 && KW == 3 && C == 8) {
        conv_hswish_relu_kernel<8, 3><<<grid, block, 0, stream>>>(
            input.data_ptr<float>(), weight.data_ptr<float>(), bias.data_ptr<float>(),
            output.data_ptr<float>(), N, H, W, OC, OH, OW, (int)stride, (int)pad);
    } else {
        const size_t shmem = (size_t)OC_TILE * C * KH * KW * sizeof(float);
        TORCH_CHECK(shmem <= 48 * 1024, "kernel too large for fused conv");
        conv_hswish_relu_generic_kernel<<<grid, block, shmem, stream>>>(
            input.data_ptr<float>(), weight.data_ptr<float>(), bias.data_ptr<float>(),
            output.data_ptr<float>(), N, C, H, W, OC, KH, KW, OH, OW, (int)stride, (int)pad);
    }
    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

fused_conv_act_mod = load_inline(
    name="fused_conv_hardswish_relu_v1",
    cpp_sources=_fused_cpp_src,
    cuda_sources=_fused_cuda_src,
    functions=["fused_conv_act"],
    extra_cuda_cflags=["-O3"],
    verbose=False,
)


class ModelNew(nn.Module):
    """
    Optimized model: Conv2d + HardSwish + ReLU fused into a single custom CUDA kernel.
    Keeps the same nn.Conv2d submodule so state_dict / initialization is identical
    to the reference Model.
    """
    def __init__(self, in_channels, out_channels, kernel_size):
        super(ModelNew, self).__init__()
        self.conv = nn.Conv2d(in_channels, out_channels, kernel_size)
        kh, kw = self.conv.kernel_size
        # Fused kernel stages OC_TILE(=8) filters in shared memory.
        self._fused_ok = (in_channels * kh * kw * 8 * 4) <= 48 * 1024

    def forward(self, x):
        if self._fused_ok and x.is_cuda and x.dtype == torch.float32 and x.dim() == 4:
            return fused_conv_act_mod.fused_conv_act(
                x.contiguous(),
                self.conv.weight.contiguous(),
                self.conv.bias.contiguous(),
                self.conv.stride[0],
                self.conv.padding[0],
            )
        x = self.conv(x)
        x = F.hardswish(x)
        x = torch.relu(x)
        return x
```

Write an improved complete ModelNew: fix any error, and if the solution is already correct, make it faster. Output the new code in one ```python codeblock. Just the model code, no other text, and NO testing code.
