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

__device__ __forceinline__ float mish_fast(float x) {
    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__ float activate(
    float x, float bias, float s1, float s2
) {
    x = x + bias;
    x = x - s1;
    x = x - s2;
    return mish_fast(x);
}

template<int FixedChannels>
__global__ void nhwc_vector_kernel(
    float4* __restrict__ x,
    const float4* __restrict__ bias,
    int64_t vectors,
    int channel_vectors,
    float s1,
    float s2
) {
    const int cv = FixedChannels ? FixedChannels / 4 : channel_vectors;
    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 % cv];
        float4 v = x[i];
        v.x = activate(v.x, b.x, s1, s2);
        v.y = activate(v.y, b.y, s1, s2);
        v.z = activate(v.z, b.z, s1, s2);
        v.w = activate(v.w, b.w, s1, s2);
        x[i] = v;
    }
}

__global__ void nhwc_scalar_kernel(
    float* __restrict__ x,
    const float* __restrict__ bias,
    int64_t n,
    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 < n; i += stride) {
        x[i] = activate(x[i], bias[i % channels], s1, s2);
    }
}

torch::Tensor nhwc_bias_subtract_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");
    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 n = x.numel();
    if (n == 0) return x;

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

    if (channels % 4 == 0) {
        const int64_t vectors = n / 4;
        const int blocks = static_cast<int>(
            std::min<int64_t>((vectors + threads - 1) / threads, 4096));
        auto* data = reinterpret_cast<float4*>(x.data_ptr<float>());
        const auto* b =
            reinterpret_cast<const float4*>(bias.data_ptr<float>());

        if (channels == 64) {
            nhwc_vector_kernel<64><<<blocks, threads, 0, stream>>>(
                data, b, vectors, channels / 4, float(s1), float(s2));
        } else {
            nhwc_vector_kernel<0><<<blocks, threads, 0, stream>>>(
                data, b, vectors, channels / 4, float(s1), float(s2));
        }
    } else {
        const int blocks = static_cast<int>(
            std::min<int64_t>((n + threads - 1) / threads, 4096));
        nhwc_scalar_kernel<<<blocks, threads, 0, stream>>>(
            x.data_ptr<float>(), bias.data_ptr<float>(), n, channels,
            float(s1), float(s2));
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return x;
}
"""

_fused_ops = load_inline(
    name="nhwc_conv_bias_subtract_mish_v3",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["nhwc_bias_subtract_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:
            y = self.conv(x)
            return F.mish(
                (y - self.subtract_value_1) - self.subtract_value_2
            )

        x = x.contiguous(memory_format=torch.channels_last)
        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 _fused_ops.nhwc_bias_subtract_mish(
            y,
            self.conv.bias,
            self.subtract_value_1,
            self.subtract_value_2,
        )
```