```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/cuda/CUDAContext.h>
#include <c10/cuda/CUDAGuard.h>
#include <c10/cuda/CUDAException.h>
#include <cuda_runtime.h>

__device__ __forceinline__ float minimum_subtract(float x, float c) {
    float clipped = isnan(x) ? x : (isnan(c) ? c : (x < c ? x : c));
    return clipped - c;
}

__global__ void minimum_subtract_kernel(
    const float* __restrict__ x,
    const float* __restrict__ constant,
    float* __restrict__ output,
    int64_t size
) {
    const int64_t index =
        static_cast<int64_t>(blockIdx.x) * blockDim.x + threadIdx.x;
    const int64_t vector_count = size / 4;
    const float c = constant[0];

    if (index < vector_count) {
        const float4 value = reinterpret_cast<const float4*>(x)[index];
        float4 result;
        result.x = minimum_subtract(value.x, c);
        result.y = minimum_subtract(value.y, c);
        result.z = minimum_subtract(value.z, c);
        result.w = minimum_subtract(value.w, c);
        reinterpret_cast<float4*>(output)[index] = result;
    }

    const int64_t tail = vector_count * 4 + index;
    if (tail < size) {
        output[tail] = minimum_subtract(x[tail], c);
    }
}

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.scalar_type() == at::kFloat &&
                constant.scalar_type() == at::kFloat,
                "Inputs must be float32");
    TORCH_CHECK(x.is_contiguous() && constant.numel() == 1,
                "Expected contiguous input and scalar constant");

    const c10::cuda::CUDAGuard guard(x.device());
    auto output = torch::empty_like(x);
    const int64_t size = x.numel();

    if (size > 0) {
        constexpr int threads = 256;
        const int64_t work_items = (size + 3) / 4;
        const int blocks = static_cast<int>(
            (work_items + threads - 1) / threads
        );

        minimum_subtract_kernel<<<
            blocks, threads, 0, at::cuda::getCurrentCUDAStream()
        >>>(
            x.data_ptr<float>(),
            constant.data_ptr<float>(),
            output.data_ptr<float>(),
            size
        );
        C10_CUDA_KERNEL_LAUNCH_CHECK();
    }
    return output;
}
"""

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


class _MinimumSubtract(torch.autograd.Function):
    @staticmethod
    def forward(ctx, x, constant):
        ctx.save_for_backward(x, constant)
        return _fused_ops.minimum_subtract_cuda(x, constant)

    @staticmethod
    def backward(ctx, grad_output):
        x, constant = ctx.saved_tensors
        weight = torch.where(
            x == constant,
            0.5,
            (x < constant).to(x.dtype),
        )
        weight = torch.where(
            torch.isnan(x) | torch.isnan(constant),
            torch.ones_like(weight),
            weight,
        )

        grad_x = grad_output * weight if ctx.needs_input_grad[0] else None
        grad_constant = None
        if ctx.needs_input_grad[1]:
            constant_weight = torch.where(
                torch.isnan(x) | torch.isnan(constant),
                torch.zeros_like(weight),
                -weight,
            )
            grad_constant = (
                grad_output * constant_weight
            ).sum().reshape_as(constant)

        return grad_x, grad_constant


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):
        y = self.linear(x)
        if (
            y.is_cuda
            and y.dtype == torch.float32
            and y.is_contiguous()
            and self.constant.device == y.device
            and self.constant.dtype == y.dtype
            and self.constant.numel() == 1
        ):
            if torch.is_grad_enabled() and (
                y.requires_grad or self.constant.requires_grad
            ):
                return _MinimumSubtract.apply(y, self.constant)
            return _fused_ops.minimum_subtract_cuda(y, self.constant)

        return torch.minimum(y, self.constant) - self.constant
```