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


_cpp_source = r"""
#include <torch/extension.h>

torch::Tensor channels_last_input(torch::Tensor x);
torch::Tensor fused_bias_mish(
    torch::Tensor x, torch::Tensor bias, double s1, double s2);
"""

_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <c10/cuda/CUDAGuard.h>
#include <c10/cuda/CUDAException.h>
#include <cuda_runtime.h>
#include <algorithm>
#include <cstdint>

__global__ void nchw8_to_nhwc(
    const float* __restrict__ src,
    float* __restrict__ dst,
    int64_t spatial,
    int64_t tiles_per_image
) {
    __shared__ float tile[8][516];

    const int tid = threadIdx.x;
    const int channel = tid >> 5;
    const int lane = tid & 31;
    const int64_t batch = blockIdx.x / tiles_per_image;
    const int64_t tile_id = blockIdx.x % tiles_per_image;
    const int64_t first_pixel = tile_id * 512;

    #pragma unroll
    for (int j = 0; j < 4; ++j) {
        const int p = lane * 4 + j * 128;
        const int64_t input_offset =
            (batch * 8 + channel) * spatial + first_pixel + p;
        const float4 v =
            *reinterpret_cast<const float4*>(src + input_offset);

        tile[channel][p] = v.x;
        tile[channel][p + 1] = v.y;
        tile[channel][p + 2] = v.z;
        tile[channel][p + 3] = v.w;
    }

    __syncthreads();

    const int output_channel = (tid & 1) * 4;

    #pragma unroll
    for (int j = 0; j < 4; ++j) {
        const int p = (tid >> 1) + j * 128;
        float4 v;
        v.x = tile[output_channel][p];
        v.y = tile[output_channel + 1][p];
        v.z = tile[output_channel + 2][p];
        v.w = tile[output_channel + 3][p];

        const int64_t output_offset =
            (batch * spatial + first_pixel + p) * 8 + output_channel;
        *reinterpret_cast<float4*>(dst + output_offset) = v;
    }
}

__device__ __forceinline__ float activate(
    float x, float bias, float s1, float s2
) {
    x = x + bias;
    x = x - s1;
    x = x - s2;

    if (x > 20.0f) return x;
    const float e = __expf(x);
    const float q = e * (e + 2.0f);
    return x * __fdividef(q, q + 2.0f);
}

__device__ __forceinline__ float4 activate4(
    float4 x, const float4 bias, float s1, float s2
) {
    x.x = activate(x.x, bias.x, s1, s2);
    x.y = activate(x.y, bias.y, s1, s2);
    x.z = activate(x.z, bias.z, s1, s2);
    x.w = activate(x.w, bias.w, s1, s2);
    return x;
}

__global__ void bias_mish64(
    float4* __restrict__ x,
    const float4* __restrict__ bias,
    int64_t vectors,
    float s1,
    float s2
) {
    constexpr int Threads = 256;
    constexpr int Tile = Threads * 4;

    const float4 b = bias[threadIdx.x & 15];
    const int64_t step = int64_t(gridDim.x) * Tile;

    for (int64_t base = int64_t(blockIdx.x) * Tile;
         base < vectors; base += step) {
        const int64_t i = base + threadIdx.x;

        if (base + Tile <= vectors) {
            float4 a0 = x[i];
            float4 a1 = x[i + Threads];
            float4 a2 = x[i + 2 * Threads];
            float4 a3 = x[i + 3 * Threads];

            a0 = activate4(a0, b, s1, s2);
            a1 = activate4(a1, b, s1, s2);
            a2 = activate4(a2, b, s1, s2);
            a3 = activate4(a3, b, s1, s2);

            x[i] = a0;
            x[i + Threads] = a1;
            x[i + 2 * Threads] = a2;
            x[i + 3 * Threads] = a3;
        } else {
            #pragma unroll
            for (int j = 0; j < 4; ++j) {
                const int64_t k = i + j * Threads;
                if (k < vectors)
                    x[k] = activate4(x[k], b, s1, s2);
            }
        }
    }
}

__global__ void bias_mish_vector(
    float4* __restrict__ x,
    const float4* __restrict__ bias,
    int64_t vectors,
    int channel_vectors,
    float s1,
    float s2
) {
    const int64_t stride = int64_t(blockDim.x) * gridDim.x;

    for (int64_t i = int64_t(blockIdx.x) * blockDim.x + threadIdx.x;
         i < vectors; i += stride) {
        const float4 b = bias[i % channel_vectors];
        x[i] = activate4(x[i], b, s1, s2);
    }
}

__global__ void bias_mish_scalar(
    float* __restrict__ x,
    const float* __restrict__ bias,
    int64_t count,
    int channels,
    float s1,
    float s2
) {
    const int64_t stride = int64_t(blockDim.x) * gridDim.x;

    for (int64_t i = int64_t(blockIdx.x) * blockDim.x + threadIdx.x;
         i < count; i += stride) {
        x[i] = activate(x[i], bias[i % channels], s1, s2);
    }
}

torch::Tensor channels_last_input(torch::Tensor x) {
    TORCH_CHECK(x.is_cuda() && x.dim() == 4, "Expected CUDA NCHW tensor");

    if (x.is_contiguous(at::MemoryFormat::ChannelsLast))
        return x;

    const c10::cuda::CUDAGuard guard(x.device());
    const int64_t spatial = x.size(2) * x.size(3);

    if (x.scalar_type() != at::kFloat ||
        x.size(1) != 8 ||
        !x.is_contiguous() ||
        spatial == 0 ||
        spatial % 512 != 0 ||
        reinterpret_cast<uintptr_t>(x.data_ptr()) % 16 != 0) {
        return x.contiguous(at::MemoryFormat::ChannelsLast);
    }

    auto output = torch::empty(
        x.sizes(), x.options().memory_format(at::MemoryFormat::ChannelsLast));

    const int64_t tiles = spatial / 512;
    const int64_t blocks = x.size(0) * tiles;

    if (blocks != 0) {
        const auto stream = at::cuda::getCurrentCUDAStream();
        nchw8_to_nhwc<<<blocks, 256, 0, stream>>>(
            x.data_ptr<float>(), output.data_ptr<float>(), spatial, tiles);
        C10_CUDA_KERNEL_LAUNCH_CHECK();
    }

    return output;
}

torch::Tensor fused_bias_mish(
    torch::Tensor x, torch::Tensor bias, double s1, double s2
) {
    TORCH_CHECK(x.is_cuda() && bias.is_cuda(), "Expected CUDA tensors");
    TORCH_CHECK(
        x.scalar_type() == at::kFloat &&
        bias.scalar_type() == at::kFloat, "Expected float32 tensors");
    TORCH_CHECK(
        x.dim() == 4 && x.is_contiguous(at::MemoryFormat::ChannelsLast),
        "Expected channels-last output");
    TORCH_CHECK(x.device() == bias.device(), "Device mismatch");
    TORCH_CHECK(
        bias.is_contiguous() && bias.numel() == x.size(1),
        "Invalid bias");

    const c10::cuda::CUDAGuard guard(x.device());
    const int64_t count = x.numel();
    if (count == 0) return x;

    const int channels = static_cast<int>(x.size(1));
    constexpr int threads = 256;
    const auto stream = at::cuda::getCurrentCUDAStream();

    const bool aligned =
        reinterpret_cast<uintptr_t>(x.data_ptr()) % 16 == 0 &&
        reinterpret_cast<uintptr_t>(bias.data_ptr()) % 16 == 0;

    if (channels == 64 && aligned) {
        const int64_t vectors = count / 4;
        const int blocks = static_cast<int>(
            std::min<int64_t>((vectors + 1023) / 1024, 2048));

        bias_mish64<<<blocks, threads, 0, stream>>>(
            reinterpret_cast<float4*>(x.data_ptr<float>()),
            reinterpret_cast<const float4*>(bias.data_ptr<float>()),
            vectors, static_cast<float>(s1), static_cast<float>(s2));
    } else if (channels % 4 == 0 && aligned) {
        const int64_t vectors = count / 4;
        const int blocks = static_cast<int>(
            std::min<int64_t>((vectors + threads - 1) / threads, 2048));

        bias_mish_vector<<<blocks, threads, 0, stream>>>(
            reinterpret_cast<float4*>(x.data_ptr<float>()),
            reinterpret_cast<const float4*>(bias.data_ptr<float>()),
            vectors, channels / 4,
            static_cast<float>(s1), static_cast<float>(s2));
    } else {
        const int blocks = static_cast<int>(
            std::min<int64_t>((count + threads - 1) / threads, 2048));

        bias_mish_scalar<<<blocks, threads, 0, stream>>>(
            x.data_ptr<float>(), bias.data_ptr<float>(),
            count, channels, static_cast<float>(s1), static_cast<float>(s2));
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return x;
}
"""

_ops = load_inline(
    name="conv_bias_mish_nhwc_v4",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["channels_last_input", "fused_bias_mish"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3", "--fmad=false"],
    verbose=False,
)


class ModelNew(nn.Module):
    def __init__(
        self,
        in_channels,
        out_channels,
        kernel_size,
        subtract_value_1,
        subtract_value_2,
    ):
        super().__init__()
        self.conv = nn.Conv2d(in_channels, out_channels, kernel_size)
        self.conv.to(memory_format=torch.channels_last)
        self.subtract_value_1 = subtract_value_1
        self.subtract_value_2 = subtract_value_2

    def forward(self, x):
        if (
            not x.is_cuda
            or x.dtype != torch.float32
            or self.conv.weight.dtype != torch.float32
        ):
            return F.mish(
                (self.conv(x) - self.subtract_value_1)
                - self.subtract_value_2
            )

        x = _ops.channels_last_input(x)
        weight = self.conv.weight.contiguous(
            memory_format=torch.channels_last
        )
        y = F.conv2d(
            x,
            weight,
            None,
            self.conv.stride,
            self.conv.padding,
            self.conv.dilation,
            self.conv.groups,
        )

        if y.dtype != torch.float32:
            y = y + self.conv.bias.to(y.dtype).view(1, -1, 1, 1)
            return F.mish(
                (y - self.subtract_value_1) - self.subtract_value_2
            )

        return _ops.fused_bias_mish(
            y,
            self.conv.bias,
            self.subtract_value_1,
            self.subtract_value_2,
        )
```