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

# Key improvement over turn 2: the profiler showed that converting the input
# to channels_last with PyTorch's generic strided-copy elementwise kernel cost
# ~3.1 ms (only ~170 GB/s effective). This replaces it with a custom tiled
# NCHW->NHWC transpose kernel that is fully coalesced on both loads and
# stores (float4 on the NHWC side), running at close to memory bandwidth
# (~0.3 ms for the 268 MB input). Everything else stays: cuDNN's native NHWC
# TF32 conv engine plus the layout-agnostic fused (subtract + Mish) epilogue.
fused_pipe_cuda_source = r"""
#include <torch/extension.h>
#include <cuda_runtime.h>
#include <cstdint>

// ------------------------- fused subtract + Mish -------------------------
__device__ __forceinline__ float mish_stable(float x) {
    // mish(x) = x * tanh(softplus(x)); threshold matches PyTorch softplus.
    float sp = (x > 20.0f) ? x : log1pf(expf(x));
    return x * tanhf(sp);
}

// Two float4 (32 bytes) per thread for extra memory-level parallelism.
__global__ void sub_mish_vec8_kernel(const float4* __restrict__ in,
                                     float4* __restrict__ out,
                                     const float sub,
                                     const long n8) {
    const long i = (long)blockIdx.x * (long)blockDim.x + threadIdx.x;
    if (i < n8) {
        float4 a = in[2 * i];
        float4 b = in[2 * i + 1];
        a.x = mish_stable(a.x - sub);
        a.y = mish_stable(a.y - sub);
        a.z = mish_stable(a.z - sub);
        a.w = mish_stable(a.w - sub);
        b.x = mish_stable(b.x - sub);
        b.y = mish_stable(b.y - sub);
        b.z = mish_stable(b.z - sub);
        b.w = mish_stable(b.w - sub);
        out[2 * i]     = a;
        out[2 * i + 1] = b;
    }
}

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

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

    // Any dense layout works for an elementwise op.
    torch::Tensor xc = x;
    if (!x.is_contiguous(at::MemoryFormat::Contiguous) &&
        !x.is_contiguous(at::MemoryFormat::ChannelsLast)) {
        xc = x.contiguous();
    }

    auto out = torch::empty_like(xc);  // preserves xc's memory format
    const long n = xc.numel();
    if (n == 0) {
        return out;
    }

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

    const bool aligned =
        ((reinterpret_cast<uintptr_t>(in_ptr) |
          reinterpret_cast<uintptr_t>(out_ptr)) & 0xF) == 0;

    const long n8 = aligned ? (n / 8) : 0;
    if (n8 > 0) {
        const long blocks = (n8 + block - 1) / block;
        sub_mish_vec8_kernel<<<(unsigned int)blocks, block>>>(
            reinterpret_cast<const float4*>(in_ptr),
            reinterpret_cast<float4*>(out_ptr),
            sub_f, n8);
    }

    const long start = n8 * 8;
    if (start < n) {
        const long rem = n - start;
        const long blocks = (rem + block - 1) / block;
        sub_mish_scalar_kernel<<<(unsigned int)blocks, block>>>(
            in_ptr, out_ptr, sub_f, start, n);
    }

    return out;
}

// -------------------- fast NCHW -> NHWC input transpose --------------------
// Fast path: C % 4 == 0. Each thread produces one float4 of NHWC output
// (4 channels at one spatial location). Loads are coalesced per channel
// plane, stores are perfectly coalesced (contiguous NHWC run).
__global__ void nchw_to_nhwc_c4_kernel(const float* __restrict__ in,
                                       float* __restrict__ out,
                                       const long N,
                                       const long C,
                                       const long HW) {
    const long idx = (long)blockIdx.x * (long)blockDim.x + threadIdx.x;
    const long C4 = C >> 2;
    const long total = N * HW * C4;
    if (idx >= total) return;

    const long c4 = idx % C4;
    const long p  = idx / C4;
    const long hw = p % HW;
    const long n  = p / HW;

    const float* base = in + (n * C + (c4 << 2)) * HW + hw;
    float4 v;
    v.x = base[0];
    v.y = base[HW];
    v.z = base[2 * HW];
    v.w = base[3 * HW];
    *reinterpret_cast<float4*>(out + p * C + (c4 << 2)) = v;
}

// Generic fallback for arbitrary C. Stores fully coalesced.
__global__ void nchw_to_nhwc_generic_kernel(const float* __restrict__ in,
                                            float* __restrict__ out,
                                            const long N,
                                            const long C,
                                            const long HW) {
    const long idx = (long)blockIdx.x * (long)blockDim.x + threadIdx.x;
    const long total = N * HW * C;
    if (idx >= total) return;

    const long c  = idx % C;
    const long p  = idx / C;
    const long hw = p % HW;
    const long n  = p / HW;

    out[idx] = in[(n * C + c) * HW + hw];
}

torch::Tensor nchw_to_nhwc_cuda(torch::Tensor x) {
    TORCH_CHECK(x.is_cuda(), "x must be a CUDA tensor");
    TORCH_CHECK(x.dim() == 4, "x must be 4D NCHW");
    TORCH_CHECK(x.scalar_type() == torch::kFloat32, "x must be float32");

    torch::Tensor xc = x.is_contiguous() ? x : x.contiguous();
    auto out = torch::empty_like(xc, xc.options(),
                                 torch::MemoryFormat::ChannelsLast);

    const long N  = xc.size(0);
    const long C  = xc.size(1);
    const long HW = (long)xc.size(2) * (long)xc.size(3);
    if (N * C * HW == 0) {
        return out;
    }

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

    if ((C & 3) == 0) {
        const long total = N * HW * (C >> 2);
        const long blocks = (total + block - 1) / block;
        nchw_to_nhwc_c4_kernel<<<(unsigned int)blocks, block>>>(
            in_ptr, out_ptr, N, C, HW);
    } else {
        const long total = N * HW * C;
        const long blocks = (total + block - 1) / block;
        nchw_to_nhwc_generic_kernel<<<(unsigned int)blocks, block>>>(
            in_ptr, out_ptr, N, C, HW);
    }

    return out;
}
"""

fused_pipe_cpp_source = (
    "torch::Tensor fused_sub_mish_cuda(torch::Tensor x, double sub);\n"
    "torch::Tensor nchw_to_nhwc_cuda(torch::Tensor x);"
)

fused_pipe = load_inline(
    name="fused_conv_pipe_v3",
    cpp_sources=fused_pipe_cpp_source,
    cuda_sources=fused_pipe_cuda_source,
    functions=["fused_sub_mish_cuda", "nchw_to_nhwc_cuda"],
    verbose=False,
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Optimized model.

    1. Converts the NCHW input to channels_last with a custom, fully
       coalesced CUDA transpose kernel (~10x faster than PyTorch's generic
       strided copy for this permutation).
    2. Runs the convolution natively in channels_last (NHWC): cuDNN's
       fastest TF32 tensor-core kernels are NHWC-based, so NHWC input with
       NHWC weights avoids the internal transpose engines.
    3. Fuses (x - subtract_value_1 - subtract_value_2) + Mish into a single
       custom CUDA kernel that operates directly on whatever dense layout
       the convolution produced (no extra contiguous() copy).
    """
    def __init__(self, in_channels, out_channels, kernel_size, subtract_value_1, subtract_value_2):
        super(ModelNew, self).__init__()
        self.conv = nn.Conv2d(in_channels, out_channels, kernel_size)
        # Fold the two scalar subtractions into one (plain attribute, so it
        # stays out of the state_dict and weight loading is unaffected).
        self.sub_total = float(subtract_value_1) + float(subtract_value_2)
        self._cl_ready = False
        # Let cuDNN autotune the fastest (NHWC) convolution algorithm.
        torch.backends.cudnn.benchmark = True
        torch.backends.cudnn.allow_tf32 = True
        try:
            torch.backends.cudnn.benchmark_limit = 0  # exhaustive search
        except Exception:
            pass

    def forward(self, x):
        if not x.is_cuda:
            # Safe fallback for CPU tensors.
            return F.mish(self.conv(x) - self.sub_total)

        if not self._cl_ready:
            # Convert weights to channels_last once (done lazily here so it
            # happens after any state_dict loading, not at __init__ time).
            with torch.no_grad():
                self.conv.weight.data = self.conv.weight.data.to(
                    memory_format=torch.channels_last)
            self._cl_ready = True

        # Fast custom NCHW -> NHWC transpose (bandwidth-bound, coalesced).
        if x.dim() == 4 and x.dtype == torch.float32:
            x = fused_pipe.nchw_to_nhwc_cuda(x)
        else:
            x = x.to(memory_format=torch.channels_last)

        # NHWC input -> cuDNN runs its native NHWC engine, NHWC output.
        x = self.conv(x)
        # Fused subtract + Mish directly on the dense NHWC buffer.
        return fused_pipe.fused_sub_mish_cuda(x, self.sub_total)
```