```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
from torch.utils.cpp_extension import load_inline

# Enable cuDNN autotuning so the (fixed-shape) convolution picks the fastest algorithm.
torch.backends.cudnn.benchmark = True

# Custom CUDA kernel that fuses the two subtractions and the Mish activation
# into a single memory-bound elementwise pass over the convolution output.
# This replaces three full-tensor kernels (sub, sub, mish) with one,
# eliminating two full read+write passes over the activation tensor.
fused_op_source = """
#include <torch/extension.h>
#include <cuda_runtime.h>

// Compute mish((x - s1) - s2), preserving the original order of operations
// so results match PyTorch bit-for-bit on the subtraction part.
__device__ __forceinline__ float sub_sub_mish_scalar(float x, float s1, float s2) {
    float v = (x - s1) - s2;
    // mish(v) = v * tanh(softplus(v)), softplus(v) = log1p(exp(v))
    // For large v, exp(v) -> inf, log1p(inf) -> inf, tanh(inf) -> 1, so result -> v (no NaN).
    return v * tanhf(log1pf(expf(v)));
}

__global__ void sub_sub_mish_kernel(const float* __restrict__ x,
                                    float* __restrict__ out,
                                    float s1, float s2, long n) {
    long idx = (long)blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        out[idx] = sub_sub_mish_scalar(x[idx], s1, s2);
    }
}

// Vectorized (float4) path for better memory throughput when numel % 4 == 0.
__global__ void sub_sub_mish_vec4_kernel(const float4* __restrict__ x,
                                         float4* __restrict__ out,
                                         float s1, float s2, long n4) {
    long idx = (long)blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n4) {
        float4 v = x[idx];
        float4 r;
        r.x = sub_sub_mish_scalar(v.x, s1, s2);
        r.y = sub_sub_mish_scalar(v.y, s1, s2);
        r.z = sub_sub_mish_scalar(v.z, s1, s2);
        r.w = sub_sub_mish_scalar(v.w, s1, s2);
        out[idx] = r;
    }
}

torch::Tensor fused_sub_sub_mish_cuda(torch::Tensor x, double s1, double s2) {
    TORCH_CHECK(x.is_cuda(), "x must be a CUDA tensor");
    TORCH_CHECK(x.scalar_type() == at::kFloat, "x must be float32");

    auto x_c = x.contiguous();
    auto out = torch::empty_like(x_c);
    long n = x_c.numel();
    if (n == 0) {
        return out;
    }

    const int threads = 256;
    const float fs1 = (float)s1;
    const float fs2 = (float)s2;
    cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    if (n % 4 == 0) {
        long n4 = n / 4;
        long blocks = (n4 + threads - 1) / threads;
        sub_sub_mish_vec4_kernel<<<blocks, threads, 0, stream>>>(
            reinterpret_cast<const float4*>(x_c.data_ptr<float>()),
            reinterpret_cast<float4*>(out.data_ptr<float>()),
            fs1, fs2, n4);
    } else {
        long blocks = (n + threads - 1) / threads;
        sub_sub_mish_kernel<<<blocks, threads, 0, stream>>>(
            x_c.data_ptr<float>(), out.data_ptr<float>(), fs1, fs2, n);
    }

    return out;
}
"""

fused_op_cpp_source = (
    "torch::Tensor fused_sub_sub_mish_cuda(torch::Tensor x, double s1, double s2);"
)

fused_sub_sub_mish = load_inline(
    name="fused_sub_sub_mish",
    cpp_sources=fused_op_cpp_source,
    cuda_sources=fused_op_source,
    functions=["fused_sub_sub_mish_cuda"],
    verbose=False,
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Optimized model: keeps the cuDNN convolution (the dominant compute) and
    replaces the two subtractions + Mish activation with a single fused
    custom CUDA kernel, reducing global-memory traffic by ~3x for the
    elementwise portion.
    """
    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 = subtract_value_1
        self.subtract_value_2 = subtract_value_2
        self.fused_op = fused_sub_sub_mish

    def forward(self, x):
        x = self.conv(x)
        x = self.fused_op.fused_sub_sub_mish_cuda(
            x, self.subtract_value_1, self.subtract_value_2
        )
        return x
```