```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: y = mish((x - v1) - v2)
# where mish(t) = t * tanh(log1p(exp(t)))
# This fuses the two scalar subtractions and the Mish activation into a
# single memory-bound pass over the convolution output (instead of three
# separate elementwise kernels), and matches PyTorch's CUDA math exactly.
fused_sub_mish_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <c10/cuda/CUDAGuard.h>
#include <cuda_runtime.h>

__device__ __forceinline__ float sub_sub_mish_scalar(float x, float v1, float v2) {
    float t = (x - v1) - v2;
    return t * tanhf(log1pf(expf(t)));
}

__global__ void sub_sub_mish_vec4_kernel(const float4* __restrict__ in,
                                         float4* __restrict__ out,
                                         const float v1,
                                         const float v2,
                                         const long n4) {
    const long stride = (long)gridDim.x * (long)blockDim.x;
    for (long i = (long)blockIdx.x * (long)blockDim.x + (long)threadIdx.x; i < n4; i += stride) {
        float4 v = in[i];
        v.x = sub_sub_mish_scalar(v.x, v1, v2);
        v.y = sub_sub_mish_scalar(v.y, v1, v2);
        v.z = sub_sub_mish_scalar(v.z, v1, v2);
        v.w = sub_sub_mish_scalar(v.w, v1, v2);
        out[i] = v;
    }
}

__global__ void sub_sub_mish_tail_kernel(const float* __restrict__ in,
                                         float* __restrict__ out,
                                         const float v1,
                                         const float v2,
                                         const long start,
                                         const long n) {
    long i = start + (long)threadIdx.x;
    if (i < n) {
        out[i] = sub_sub_mish_scalar(in[i], v1, v2);
    }
}

torch::Tensor fused_sub_mish(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");

    const c10::cuda::OptionalCUDAGuard device_guard(x.device());

    auto xc = x.is_contiguous() ? x : x.contiguous();
    auto out = torch::empty_like(xc);

    const long n = (long)xc.numel();
    if (n == 0) {
        return out;
    }

    const float f1 = static_cast<float>(v1);
    const float f2 = static_cast<float>(v2);
    cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    const long n4 = n >> 2;
    if (n4 > 0) {
        const int threads = 256;
        const int sm_count = at::cuda::getCurrentDeviceProperties()->multiProcessorCount;
        long want = (n4 + threads - 1) / threads;
        long cap = (long)sm_count * 32L;
        int blocks = (int)(want < cap ? want : cap);
        sub_sub_mish_vec4_kernel<<<blocks, threads, 0, stream>>>(
            reinterpret_cast<const float4*>(xc.data_ptr<float>()),
            reinterpret_cast<float4*>(out.data_ptr<float>()),
            f1, f2, n4);
    }

    const long rem = n & 3L;
    if (rem != 0) {
        sub_sub_mish_tail_kernel<<<1, 32, 0, stream>>>(
            xc.data_ptr<float>(), out.data_ptr<float>(), f1, f2, n4 << 2, n);
    }

    return out;
}
"""

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

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


class ModelNew(nn.Module):
    """
    Optimized model:
      * keeps the cuDNN convolution (identical numerics to the baseline)
      * replaces (subtract, subtract, Mish) with a single fused custom
        CUDA kernel, eliminating two full elementwise passes over the
        (large) convolution output tensor.
    """
    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)

    def forward(self, x):
        x = self.conv(x)
        if x.is_cuda and x.dtype == torch.float32:
            return fused_sub_mish_ext.fused_sub_mish(
                x, self.subtract_value_1, self.subtract_value_2
            )
        # Fallback for non-CUDA / non-float32 tensors
        x = x - self.subtract_value_1
        x = x - self.subtract_value_2
        x = torch.nn.functional.mish(x)
        return x
```