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


_cpp_source = r"""
#include <torch/extension.h>

torch::Tensor fused_bias_min_sub_cuda(
    torch::Tensor output,
    torch::Tensor bias,
    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 activate(
    float value, float bias, float constant) {
    value += bias;
    float minimum = (isnan(value) || isnan(constant))
        ? value + constant
        : fminf(value, constant);
    return minimum - constant;
}

__global__ void fused_vector_kernel(
    float4* __restrict__ output,
    const float4* __restrict__ bias,
    const float* __restrict__ constant,
    int64_t count,
    int64_t columns) {
    const float c = constant[0];
    for (int64_t i = int64_t(blockIdx.x) * blockDim.x + threadIdx.x;
         i < count;
         i += int64_t(blockDim.x) * gridDim.x) {
        float4 v = output[i];
        const float4 b = bias[i % columns];
        v.x = activate(v.x, b.x, c);
        v.y = activate(v.y, b.y, c);
        v.z = activate(v.z, b.z, c);
        v.w = activate(v.w, b.w, c);
        output[i] = v;
    }
}

__global__ void fused_scalar_kernel(
    float* __restrict__ output,
    const float* __restrict__ bias,
    const float* __restrict__ constant,
    int64_t count,
    int64_t columns) {
    const float c = constant[0];
    for (int64_t i = int64_t(blockIdx.x) * blockDim.x + threadIdx.x;
         i < count;
         i += int64_t(blockDim.x) * gridDim.x) {
        output[i] = activate(output[i], bias[i % columns], c);
    }
}

torch::Tensor fused_bias_min_sub_cuda(
    torch::Tensor output,
    torch::Tensor bias,
    torch::Tensor constant) {
    TORCH_CHECK(output.is_cuda() && bias.is_cuda() && constant.is_cuda(),
                "All tensors must be CUDA tensors");
    TORCH_CHECK(output.device() == bias.device() &&
                output.device() == constant.device(),
                "All tensors must be on the same device");
    TORCH_CHECK(output.scalar_type() == at::kFloat &&
                bias.scalar_type() == at::kFloat &&
                constant.scalar_type() == at::kFloat,
                "All tensors must be float32");
    TORCH_CHECK(output.dim() == 2 && bias.dim() == 1 &&
                output.size(1) == bias.numel() && constant.numel() == 1,
                "Invalid tensor shapes");
    TORCH_CHECK(output.is_contiguous() && bias.is_contiguous(),
                "Output and bias must be contiguous");

    const c10::cuda::CUDAGuard guard(output.device());
    const int64_t n = output.numel();
    if (n == 0) return output;

    const int64_t columns = output.size(1);
    constexpr int threads = 256;
    cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    if (columns % 4 == 0) {
        const int64_t count = n / 4;
        const int blocks = static_cast<int>(
            std::min<int64_t>((count + threads - 1) / threads, 4096));
        fused_vector_kernel<<<blocks, threads, 0, stream>>>(
            reinterpret_cast<float4*>(output.data_ptr<float>()),
            reinterpret_cast<const float4*>(bias.data_ptr<float>()),
            constant.data_ptr<float>(), count, columns / 4);
    } else {
        const int blocks = static_cast<int>(
            std::min<int64_t>((n + threads - 1) / threads, 4096));
        fused_scalar_kernel<<<blocks, threads, 0, stream>>>(
            output.data_ptr<float>(), bias.data_ptr<float>(),
            constant.data_ptr<float>(), n, columns);
    }
    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_fused_ops = load_inline(
    name="linear_bias_min_sub_fused",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["fused_bias_min_sub_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):
        if (
            x.is_cuda
            and x.ndim == 2
            and x.dtype == torch.float32
            and self.linear.weight.dtype == torch.float32
            and self.linear.bias.dtype == torch.float32
            and self.constant.dtype == torch.float32
            and not torch.is_autocast_enabled()
            and not (
                torch.is_grad_enabled()
                and any(t.requires_grad for t in (
                    x, self.linear.weight, self.linear.bias, self.constant
                ))
            )
        ):
            output = torch.mm(x, self.linear.weight.t())
            return _fused_ops.fused_bias_min_sub_cuda(
                output, self.linear.bias, self.constant
            )

        output = F.linear(x, self.linear.weight, self.linear.bias)
        return torch.minimum(output, self.constant) - self.constant
```