```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: (x - v1) - v2 followed by Mish activation.
# This replaces three separate memory-bound elementwise kernels
# (sub, sub, mish) with a single vectorized pass over the data.
fused_sub_mish_cuda_source = r"""
#include <torch/extension.h>
#include <cuda_runtime.h>
#include <ATen/cuda/CUDAContext.h>

__device__ __forceinline__ float mish_act(float x) {
    // Matches ATen's CUDA mish: x * tanh(log1p(exp(x))) computed in fp32.
    return x * tanhf(log1pf(expf(x)));
}

__global__ void fused_sub_mish_kernel(const float* __restrict__ in,
                                      float* __restrict__ out,
                                      const long n,
                                      const float v1,
                                      const float v2) {
    const long n_vec = n >> 2;
    const long i = (long)blockIdx.x * blockDim.x + threadIdx.x;

    if (i < n_vec) {
        // Vectorized 128-bit load/store path.
        float4 v = reinterpret_cast<const float4*>(in)[i];
        v.x = mish_act((v.x - v1) - v2);
        v.y = mish_act((v.y - v1) - v2);
        v.z = mish_act((v.z - v1) - v2);
        v.w = mish_act((v.w - v1) - v2);
        reinterpret_cast<float4*>(out)[i] = v;
    } else if (i == n_vec) {
        // Scalar tail for sizes not divisible by 4.
        for (long j = n_vec << 2; j < n; ++j) {
            float x = in[j];
            out[j] = mish_act((x - v1) - v2);
        }
    }
}

torch::Tensor fused_sub_mish_cuda(torch::Tensor x, double v1, double v2) {
    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();
    const long n = xc.numel();
    const long slots = (n >> 2) + ((n & 3) ? 1 : 0);

    const int threads = 256;
    const long blocks = (slots + threads - 1) / threads;
    auto stream = at::cuda::getCurrentCUDAStream();

    // In-place on the freshly produced conv output: saves a full 2.1GB
    // allocation and one read+write round trip compared to eager.
    fused_sub_mish_kernel<<<blocks, threads, 0, stream.stream()>>>(
        xc.data_ptr<float>(),
        xc.data_ptr<float>(),
        n,
        static_cast<float>(v1),
        static_cast<float>(v2));

    return xc;
}
"""

fused_sub_mish_cpp_source = (
    "torch::Tensor fused_sub_mish_cuda(torch::Tensor x, double v1, double v2);"
)

# Compile the inline CUDA extension
fused_sub_mish = load_inline(
    name="fused_sub_mish",
    cpp_sources=fused_sub_mish_cpp_source,
    cuda_sources=fused_sub_mish_cuda_source,
    functions=["fused_sub_mish_cuda"],
    verbose=False,
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Optimized model: keeps the cuDNN Conv2d (which is already highly
    optimized, TF32 tensor-core accelerated) and replaces the three
    memory-bound elementwise ops (subtract, subtract, Mish) with a single
    fused, vectorized custom CUDA kernel.
    """
    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)
        self.subtract_value_1 = float(subtract_value_1)
        self.subtract_value_2 = float(subtract_value_2)
        self.fused_op = fused_sub_mish

    def forward(self, x):
        x = self.conv(x)
        if x.is_cuda and x.dtype == torch.float32:
            return self.fused_op.fused_sub_mish_cuda(
                x, self.subtract_value_1, self.subtract_value_2
            )
        # Fallback for non-CUDA / non-fp32 inputs.
        x = x - self.subtract_value_1
        x = x - self.subtract_value_2
        return F.mish(x)
```