```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 for: out = min(x, c) - c
# Combines the two elementwise passes (torch.min followed by subtraction)
# into a single vectorized (float4) memory pass over the linear output.
fused_min_sub_source = """
#include <torch/extension.h>
#include <cuda_runtime.h>

__global__ void fused_min_sub_vec4_kernel(const float4* __restrict__ x,
                                          float4* __restrict__ out,
                                          const float* __restrict__ cptr,
                                          int n4) {
    const float c = *cptr;
    const int stride = blockDim.x * gridDim.x;
    for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < n4; i += stride) {
        float4 v = x[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__ x,
                                          float* __restrict__ out,
                                          const float* __restrict__ cptr,
                                          int start,
                                          int n) {
    const float c = *cptr;
    int i = start + blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        out[i] = fminf(x[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() == torch::kFloat32, "x must be float32");

    auto xc = x.contiguous();
    auto cc = c.to(torch::kFloat32).contiguous();
    auto out = torch::empty_like(xc);

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

    const float* xp = xc.data_ptr<float>();
    float* op = out.data_ptr<float>();
    const float* cp = cc.data_ptr<float>();

    const int block = 256;
    if (n4 > 0) {
        int grid = (n4 + block - 1) / block;
        if (grid > 8192) grid = 8192;
        fused_min_sub_vec4_kernel<<<grid, block>>>(
            reinterpret_cast<const float4*>(xp),
            reinterpret_cast<float4*>(op),
            cp, n4);
    }

    const int rem = (int)(n & 3);
    if (rem > 0) {
        fused_min_sub_tail_kernel<<<1, rem>>>(xp, op, cp, n4 * 4, (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",
    cpp_sources=fused_min_sub_cpp_source,
    cuda_sources=fused_min_sub_source,
    functions=["fused_min_sub_cuda"],
    verbose=False,
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3", "--use_fast_math"],
    extra_ldflags=[""],
)


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

    The GEMM (nn.Linear) is kept on cuBLAS via F.linear, while the two
    elementwise operators (torch.min with the constant and the subtraction
    of the constant) are fused into a single custom CUDA kernel, halving
    the elementwise memory traffic and removing one kernel launch.
    """
    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 = F.linear(x, self.linear.weight, self.linear.bias)
        if x.is_cuda and x.dtype == torch.float32:
            x = fused_min_sub.fused_min_sub_cuda(x, self.constant)
        else:
            x = torch.min(x, self.constant)
            x = x - self.constant
        return x
```