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

# Custom CUDA kernel that fuses the two element-wise ops:
#   y = min(x, c) - c
# into a single kernel. This removes one full read+write pass over the
# (batch_size x out_features) activation tensor and one kernel launch
# compared to the eager PyTorch implementation. The matmul (the dominant
# cost) is left to the highly optimized cuBLAS path via F.linear, keeping
# bit-exact fp32 numerics for the GEMM while the fused epilogue performs
# exactly the same fminf/subtract floating-point operations as eager.

fused_min_sub_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>

__global__ void fused_min_sub_vec4_kernel(const float4* __restrict__ in,
                                          float4* __restrict__ out,
                                          const float* __restrict__ c_ptr,
                                          const int n4) {
    const float c = *c_ptr;
    const int stride = blockDim.x * gridDim.x;
    for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < n4; i += stride) {
        float4 v = in[i];
        v.x = fminf(v.x, c) - c;
        v.y = fminf(v.y, c) - c;
        v.z = fminf(v.z, c) - c;
        v.w = fminf(v.w, c) - c;
        out[i] = v;
    }
}

__global__ void fused_min_sub_tail_kernel(const float* __restrict__ in,
                                          float* __restrict__ out,
                                          const float* __restrict__ c_ptr,
                                          const int start,
                                          const int n) {
    const float c = *c_ptr;
    const int i = start + threadIdx.x;
    if (i < n) {
        out[i] = fminf(in[i], c) - c;
    }
}

torch::Tensor fused_min_sub_cuda(torch::Tensor x, torch::Tensor c) {
    TORCH_CHECK(x.is_cuda(), "x must be a CUDA tensor");
    TORCH_CHECK(x.scalar_type() == at::kFloat, "x must be float32");
    TORCH_CHECK(c.is_cuda(), "c must be a CUDA tensor");
    TORCH_CHECK(c.scalar_type() == at::kFloat, "c must be float32");

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

    const long n = xc.numel();
    const int n4 = static_cast<int>(n >> 2);

    const float* in_ptr = xc.data_ptr<float>();
    float* out_ptr = out.data_ptr<float>();
    const float* c_ptr = c.data_ptr<float>();

    cudaStream_t stream = at::cuda::getCurrentCUDAStream();
    const int block = 256;

    if (n4 > 0) {
        int grid = (n4 + block - 1) / block;
        if (grid > 65535) grid = 65535;
        fused_min_sub_vec4_kernel<<<grid, block, 0, stream>>>(
            reinterpret_cast<const float4*>(in_ptr),
            reinterpret_cast<float4*>(out_ptr),
            c_ptr,
            n4);
    }

    const int rem = static_cast<int>(n & 3);
    if (rem > 0) {
        fused_min_sub_tail_kernel<<<1, rem, 0, stream>>>(
            in_ptr, out_ptr, c_ptr, n4 << 2, static_cast<int>(n));
    }

    return out;
}
"""

fused_min_sub_cpp_source = (
    "torch::Tensor fused_min_sub_cuda(torch::Tensor x, torch::Tensor c);"
)

fused_min_sub = load_inline(
    name="fused_min_sub_ext",
    cpp_sources=fused_min_sub_cpp_source,
    cuda_sources=fused_min_sub_cuda_source,
    functions=["fused_min_sub_cuda"],
    verbose=False,
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Optimized model: cuBLAS linear followed by a single fused CUDA kernel
    that computes min(x, constant) - constant in one pass.
    """

    def __init__(self, in_features, out_features, constant):
        super(ModelNew, self).__init__()
        self.linear = nn.Linear(in_features, out_features)
        self.constant = nn.Parameter(torch.tensor(constant))
        self.fused_min_sub = fused_min_sub

    def forward(self, x):
        x = self.linear(x)
        if x.is_cuda and x.dtype == torch.float32 and self.constant.is_cuda:
            return self.fused_min_sub.fused_min_sub_cuda(x, self.constant)
        # Fallback for non-CUDA execution.
        return torch.min(x, self.constant) - self.constant
```