```python
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline

_cpp_source = """
#include <torch/extension.h>
torch::Tensor minimum_subtract_cuda(torch::Tensor x, torch::Tensor constant);
"""

_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/OpMathType.h>
#include <c10/cuda/CUDAGuard.h>
#include <c10/cuda/CUDAStream.h>
#include <c10/cuda/CUDAException.h>
#include <cuda_runtime.h>

template <typename scalar_t>
__global__ void minimum_subtract_kernel(
    const scalar_t* __restrict__ x,
    const scalar_t* __restrict__ constant,
    scalar_t* __restrict__ output,
    int64_t count)
{
    using acc_t = at::opmath_type<scalar_t>;
    const acc_t c = static_cast<acc_t>(constant[0]);

    for (int64_t i = static_cast<int64_t>(blockIdx.x) * blockDim.x
                     + threadIdx.x;
         i < count;
         i += static_cast<int64_t>(blockDim.x) * gridDim.x) {
        const acc_t value = static_cast<acc_t>(x[i]);
        // Preserve torch.minimum's NaN propagation.
        const acc_t minimum =
            isnan(value) ? value : (value < c ? value : c);
        output[i] = static_cast<scalar_t>(minimum - c);
    }
}

__global__ void minimum_subtract_float4_kernel(
    const float4* __restrict__ x,
    const float* __restrict__ constant,
    float4* __restrict__ output,
    int64_t count)
{
    const float c = constant[0];

    for (int64_t i = static_cast<int64_t>(blockIdx.x) * blockDim.x
                     + threadIdx.x;
         i < count;
         i += static_cast<int64_t>(blockDim.x) * gridDim.x) {
        const float4 value = x[i];
        float4 result;
        result.x = (isnan(value.x) ? value.x :
                    (value.x < c ? value.x : c)) - c;
        result.y = (isnan(value.y) ? value.y :
                    (value.y < c ? value.y : c)) - c;
        result.z = (isnan(value.z) ? value.z :
                    (value.z < c ? value.z : c)) - c;
        result.w = (isnan(value.w) ? value.w :
                    (value.w < c ? value.w : c)) - c;
        output[i] = result;
    }
}

torch::Tensor minimum_subtract_cuda(
    torch::Tensor x, torch::Tensor constant)
{
    TORCH_CHECK(x.is_cuda() && constant.is_cuda(),
                "Inputs must be CUDA tensors");
    TORCH_CHECK(x.device() == constant.device(),
                "Inputs must be on the same device");
    TORCH_CHECK(x.is_contiguous(), "Input must be contiguous");
    TORCH_CHECK(constant.numel() == 1, "Expected a scalar constant");
    TORCH_CHECK(x.scalar_type() == constant.scalar_type(),
                "Input dtypes must match");

    const c10::cuda::CUDAGuard guard(x.device());
    auto output = torch::empty_like(x);
    const int64_t count = x.numel();
    if (count == 0) return output;

    constexpr int threads = 256;
    const auto stream = c10::cuda::getCurrentCUDAStream();

    if (x.scalar_type() == at::ScalarType::Float && count % 4 == 0) {
        const int64_t vectors = count / 4;
        const int blocks = static_cast<int>(
            std::min<int64_t>((vectors + threads - 1) / threads, 4096));
        minimum_subtract_float4_kernel<<<blocks, threads, 0, stream.stream()>>>(
            reinterpret_cast<const float4*>(x.data_ptr<float>()),
            constant.data_ptr<float>(),
            reinterpret_cast<float4*>(output.data_ptr<float>()),
            vectors);
    } else {
        const int blocks = static_cast<int>(
            std::min<int64_t>((count + threads - 1) / threads, 4096));
        AT_DISPATCH_FLOATING_TYPES_AND2(
            at::ScalarType::Half, at::ScalarType::BFloat16,
            x.scalar_type(), "minimum_subtract_cuda", [&] {
                minimum_subtract_kernel<scalar_t>
                    <<<blocks, threads, 0, stream.stream()>>>(
                        x.data_ptr<scalar_t>(),
                        constant.data_ptr<scalar_t>(),
                        output.data_ptr<scalar_t>(),
                        count);
            });
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_fused_ops = load_inline(
    name="fused_minimum_subtract_cuda",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["minimum_subtract_cuda"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    verbose=False,
)


class ModelNew(nn.Module):
    def __init__(self, in_features, out_features, constant):
        super().__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 not x.is_cuda:
            return torch.minimum(x, self.constant) - self.constant
        constant = self.constant.to(dtype=x.dtype)
        return _fused_ops.minimum_subtract_cuda(x, constant)
```