```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: out = min(x, c) - c
# This fuses the two elementwise operators (torch.min and subtraction)
# into a single memory-bound pass over the tensor. The scalar constant
# is read directly from device memory to avoid any host<->device sync.
fused_min_sub_source = """
#include <torch/extension.h>
#include <cuda_runtime.h>

__global__ void fused_min_sub_vec4_kernel(const float4* __restrict__ in,
                                          float4* __restrict__ out,
                                          const float* __restrict__ cp,
                                          long n4) {
    long idx = (long)blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n4) {
        float c = *cp;
        float4 v = in[idx];
        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[idx] = v;
    }
}

__global__ void fused_min_sub_scalar_kernel(const float* __restrict__ in,
                                            float* __restrict__ out,
                                            const float* __restrict__ cp,
                                            long start,
                                            long n) {
    long idx = start + (long)blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        float c = *cp;
        float v = in[idx];
        out[idx] = fminf(v, 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);
    long n = xc.numel();
    if (n == 0) {
        return out;
    }

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

    const int block_size = 256;
    long n4 = n / 4;

    if (n4 > 0) {
        long num_blocks = (n4 + block_size - 1) / block_size;
        fused_min_sub_vec4_kernel<<<num_blocks, block_size>>>(
            reinterpret_cast<const float4*>(in_ptr),
            reinterpret_cast<float4*>(out_ptr),
            cp, n4);
    }
    if (n4 * 4 < n) {
        long rem = n - n4 * 4;
        long num_blocks = (rem + block_size - 1) / block_size;
        fused_min_sub_scalar_kernel<<<num_blocks, block_size>>>(
            in_ptr, out_ptr, cp, n4 * 4, n);
    }

    return out;
}
"""

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

# Compile the inline CUDA extension
fused_min_sub = load_inline(
    name="fused_min_sub",
    cpp_sources=fused_min_sub_cpp_source,
    cuda_sources=fused_min_sub_source,
    functions=["fused_min_sub_cuda"],
    verbose=False,
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Optimized version of Model.

    The matrix multiplication (with fused bias add) is dispatched to cuBLAS
    via F.linear, which is optimal for this large compute-bound GEMM.
    The subsequent elementwise ops (min with the constant, then subtraction
    of the constant) are fused into a single custom vectorized CUDA kernel,
    eliminating one full read/write pass over the activations.
    """
    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):
        # GEMM + bias (cuBLAS), then fused: out = min(x, c) - c
        x = F.linear(x, self.linear.weight, self.linear.bias)
        x = self.fused_min_sub.fused_min_sub_cuda(x, self.constant)
        return x
```