```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

# Allow TF32 tensor-core math for FP32 matmuls (explicitly permitted by the
# precision rule). This keeps FP32 storage/semantics while using tensor cores.
torch.backends.cuda.matmul.allow_tf32 = True
try:
    torch.set_float32_matmul_precision("high")
except Exception:
    pass

# Custom CUDA kernel that fuses the two elementwise epilogue ops:
#     out = min(x, c) - c
# into a single pass over the GEMM output (halves the epilogue memory traffic
# and kernel launches compared to two separate PyTorch ops).
min_sub_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>
#include <algorithm>

__global__ void min_sub_const_vec4_kernel(const float4* __restrict__ in,
                                          float4* __restrict__ out,
                                          const float c,
                                          const long n4) {
    long i = blockIdx.x * (long)blockDim.x + threadIdx.x;
    const long stride = (long)gridDim.x * blockDim.x;
    for (; i < n4; i += stride) {
        float4 v = in[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 min_sub_const_kernel(const float* __restrict__ in,
                                     float* __restrict__ out,
                                     const float c,
                                     const long n) {
    long i = blockIdx.x * (long)blockDim.x + threadIdx.x;
    const long stride = (long)gridDim.x * blockDim.x;
    for (; i < n; i += stride) {
        out[i] = fminf(in[i], c) - c;
    }
}

torch::Tensor min_sub_const_cuda(torch::Tensor x, double c) {
    TORCH_CHECK(x.is_cuda(), "x must be a CUDA tensor");
    TORCH_CHECK(x.scalar_type() == at::kFloat, "x must be float32");

    auto xc = x.contiguous();
    auto out = at::empty_like(xc);
    const long n = xc.numel();
    if (n == 0) {
        return out;
    }

    const float cf = static_cast<float>(c);
    const int block = 256;
    cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    if ((n & 3) == 0) {
        const long n4 = n >> 2;
        const long blocks_needed = (n4 + block - 1) / block;
        const int grid = (int)std::min<long>(blocks_needed, 524288L);
        min_sub_const_vec4_kernel<<<grid, block, 0, stream>>>(
            reinterpret_cast<const float4*>(xc.data_ptr<float>()),
            reinterpret_cast<float4*>(out.data_ptr<float>()),
            cf, n4);
    } else {
        const long blocks_needed = (n + block - 1) / block;
        const int grid = (int)std::min<long>(blocks_needed, 524288L);
        min_sub_const_kernel<<<grid, block, 0, stream>>>(
            xc.data_ptr<float>(), out.data_ptr<float>(), cf, n);
    }
    return out;
}
"""

min_sub_cpp_source = "torch::Tensor min_sub_const_cuda(torch::Tensor x, double c);"

fused_min_sub = load_inline(
    name="fused_min_sub_const",
    cpp_sources=min_sub_cpp_source,
    cuda_sources=min_sub_source,
    functions=["min_sub_const_cuda"],
    verbose=False,
    extra_cuda_cflags=["-O3", "--use_fast_math"],
)


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

    - The linear layer (GEMM + fused bias epilogue) runs through cuBLAS with
      TF32 tensor cores enabled (allowed by the precision rule; tensors stay
      in FP32).
    - The `torch.min(x, constant)` and `x - constant` ops are fused into a
      single custom CUDA kernel, cutting the epilogue memory passes in half.
    """
    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._const_val = float(constant)

    def forward(self, x):
        x = F.linear(x, self.linear.weight, self.linear.bias)
        if x.is_cuda and x.dtype == torch.float32:
            return fused_min_sub.min_sub_const_cuda(x, self._const_val)
        # Fallback for non-CUDA execution.
        return torch.min(x, self.constant) - self.constant
```