```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 fused CUDA kernel:
#     out = min(x, c) - c        ( == min(x - c, 0) )
#
# This fuses the two element-wise operators (torch.min followed by subtract)
# into a single pass over the data, removing one full read+write of the
# activation tensor and one kernel launch. The heavy matmul (nn.Linear) is
# kept on the highly-optimized cuBLAS path.
# ---------------------------------------------------------------------------
fused_min_sub_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>

// Vectorized (float4) path: used when the number of elements is divisible by 4.
__global__ void __launch_bounds__(256) fused_min_sub_vec4_kernel(
        const float4* __restrict__ x,
        const float*  __restrict__ c,
        float4* __restrict__ out,
        const int n4) {
    const int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n4) {
        const float cv = c[0];
        float4 v = x[idx];
        v.x = fminf(v.x, cv) - cv;
        v.y = fminf(v.y, cv) - cv;
        v.z = fminf(v.z, cv) - cv;
        v.w = fminf(v.w, cv) - cv;
        out[idx] = v;
    }
}

// Scalar fallback path for arbitrary element counts.
__global__ void __launch_bounds__(256) fused_min_sub_kernel(
        const float* __restrict__ x,
        const float* __restrict__ c,
        float* __restrict__ out,
        const int n) {
    const int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        const float cv = c[0];
        out[idx] = fminf(x[idx], cv) - cv;
    }
}

torch::Tensor fused_min_sub(torch::Tensor x, torch::Tensor c) {
    TORCH_CHECK(x.is_cuda() && c.is_cuda(), "Inputs must be CUDA tensors");
    TORCH_CHECK(x.scalar_type() == at::kFloat, "x must be float32");
    TORCH_CHECK(c.scalar_type() == at::kFloat, "c must be float32");

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

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

    const int block_size = 256;
    cudaStream_t stream = at::cuda::getCurrentCUDAStream();
    const float* c_ptr = c.data_ptr<float>();

    if ((n & 3) == 0) {
        const int n4 = (int)(n >> 2);
        const int num_blocks = (n4 + block_size - 1) / block_size;
        fused_min_sub_vec4_kernel<<<num_blocks, block_size, 0, stream>>>(
            reinterpret_cast<const float4*>(x_c.data_ptr<float>()),
            c_ptr,
            reinterpret_cast<float4*>(out.data_ptr<float>()),
            n4);
    } else {
        const int num_blocks = (int)((n + block_size - 1) / block_size);
        fused_min_sub_kernel<<<num_blocks, block_size, 0, stream>>>(
            x_c.data_ptr<float>(), c_ptr, out.data_ptr<float>(), (int)n);
    }
    return out;
}
"""

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

# Compile the inline CUDA extension.
fused_min_sub_ext = 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"],
    verbose=False,
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Optimized model: the matrix multiplication stays on cuBLAS (nn.Linear),
    while `torch.min(x, constant)` and `x - constant` are fused into a single
    custom CUDA kernel (one pass over the data instead of two).
    """
    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))

    def forward(self, x):
        x = self.linear(x)
        if (x.is_cuda and x.dtype == torch.float32
                and self.constant.is_cuda
                and self.constant.dtype == torch.float32):
            # Fused: out = min(x, c) - c
            return fused_min_sub_ext.fused_min_sub(x, self.constant)
        # Fallback (exact original semantics) for non-CUDA tensors.
        x = torch.min(x, self.constant)
        x = x - self.constant
        return x


batch_size = 128
in_features = 16384
out_features = 16384
constant = 2.0


def get_inputs():
    return [torch.rand(batch_size, in_features)]


def get_init_inputs():
    return [in_features, out_features, constant]
```