```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
from torch.utils.cpp_extension import load_inline

# Fused CUDA kernel: out = min(x, c) - c
# This replaces the two separate element-wise ops (torch.min + sub) with a
# single pass over the linear layer's output, cutting global-memory traffic
# roughly in half for the element-wise stage. The GEMM itself is left to
# cuBLAS (nn.Linear), which is already near-peak for these sizes.
fused_min_sub_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>

__global__ void fused_min_sub_kernel(const float* __restrict__ x,
                                     float* __restrict__ out,
                                     const float* __restrict__ c_ptr,
                                     const long n) {
    const long i = (long)blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        const float c = c_ptr[0];
        out[i] = fminf(x[i], c) - c;
    }
}

__global__ void fused_min_sub_vec4_kernel(const float4* __restrict__ x,
                                          float4* __restrict__ out,
                                          const float* __restrict__ c_ptr,
                                          const long n4) {
    const long i = (long)blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n4) {
        const float c = c_ptr[0];
        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;
    }
}

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(c.is_cuda(), "c must be a CUDA tensor");
    TORCH_CHECK(x.scalar_type() == at::kFloat, "x must be float32");
    TORCH_CHECK(c.scalar_type() == at::kFloat, "c must be float32");
    TORCH_CHECK(x.is_contiguous(), "x must be contiguous");
    TORCH_CHECK(c.numel() == 1, "c must be a scalar");

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

    const float* x_ptr = x.data_ptr<float>();
    float* out_ptr = out.data_ptr<float>();
    const float* c_ptr = c.data_ptr<float>();

    const int block = 256;
    cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    // Fresh torch allocations are 256-byte aligned, so float4 access is safe
    // whenever the element count is a multiple of 4.
    if ((n & 3L) == 0L) {
        const long n4 = n >> 2;
        const int grid = (int)((n4 + block - 1) / block);
        fused_min_sub_vec4_kernel<<<grid, block, 0, stream>>>(
            reinterpret_cast<const float4*>(x_ptr),
            reinterpret_cast<float4*>(out_ptr),
            c_ptr, n4);
    } else {
        const int grid = (int)((n + block - 1) / block);
        fused_min_sub_kernel<<<grid, block, 0, stream>>>(
            x_ptr, out_ptr, c_ptr, 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_cuda_source,
    functions=["fused_min_sub_cuda"],
    verbose=False,
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Optimized model: keeps the cuBLAS GEMM (nn.Linear) and replaces the
    min + subtraction chain with a single fused custom CUDA kernel.
    """
    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 = self.linear(x)
        c = self.constant
        if (x.is_cuda and x.dtype == torch.float32 and
                c.is_cuda and c.dtype == torch.float32):
            return fused_min_sub.fused_min_sub_cuda(x, c)
        # Fallback for CPU / non-float32 execution.
        return torch.min(x, c) - c
```