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

# Custom CUDA kernel that fuses the two element-wise ops:
#   x = torch.min(x, constant)
#   x = x - constant
# into a single memory pass (out[i] = min(x[i], c) - c).
# The constant is read directly from GPU memory to avoid a host<->device sync.
fused_min_sub_source = """
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>

__global__ void fused_min_sub_vec4_kernel(const float4* __restrict__ x,
                                          float4* __restrict__ out,
                                          const float* __restrict__ c_ptr,
                                          int size4) {
    const float c = *c_ptr;
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = gridDim.x * blockDim.x;
    for (int i = idx; i < size4; i += stride) {
        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;
    }
}

__global__ void fused_min_sub_kernel(const float* __restrict__ x,
                                     float* __restrict__ out,
                                     const float* __restrict__ c_ptr,
                                     int size) {
    const float c = *c_ptr;
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = gridDim.x * blockDim.x;
    for (int i = idx; i < size; i += stride) {
        out[i] = fminf(x[i], c) - c;
    }
}

torch::Tensor fused_min_sub_cuda(torch::Tensor x, torch::Tensor c) {
    TORCH_CHECK(x.is_cuda() && c.is_cuda(), "Inputs must be CUDA tensors");
    TORCH_CHECK(x.scalar_type() == torch::kFloat32, "x must be float32");
    TORCH_CHECK(c.scalar_type() == torch::kFloat32, "c must be float32");

    auto x_contig = x.contiguous();
    auto out = torch::empty_like(x_contig);
    const int size = static_cast<int>(x_contig.numel());
    if (size == 0) {
        return out;
    }

    const int block_size = 256;
    const float* x_ptr = x_contig.data_ptr<float>();
    float* out_ptr = out.data_ptr<float>();
    const float* c_ptr = c.data_ptr<float>();
    cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    if (size % 4 == 0) {
        const int size4 = size / 4;
        int num_blocks = (size4 + block_size - 1) / block_size;
        fused_min_sub_vec4_kernel<<<num_blocks, block_size, 0, stream>>>(
            reinterpret_cast<const float4*>(x_ptr),
            reinterpret_cast<float4*>(out_ptr),
            c_ptr, size4);
    } else {
        int num_blocks = (size + block_size - 1) / block_size;
        fused_min_sub_kernel<<<num_blocks, block_size, 0, stream>>>(
            x_ptr, out_ptr, c_ptr, size);
    }
    return out;
}
"""

fused_min_sub_cpp_source = (
    "torch::Tensor fused_min_sub_cuda(torch::Tensor x, torch::Tensor c);"
)

# Compile the inline CUDA code for the fused min + subtract epilogue
fused_min_sub = load_inline(
    name="fused_min_sub",
    cpp_sources=fused_min_sub_cpp_source,
    cuda_sources=fused_min_sub_source,
    functions=["fused_min_sub_cuda"],
    verbose=False,
)


class ModelNew(nn.Module):
    """
    Optimized model: the matmul+bias stays on the highly-tuned cuBLAS path
    (nn.Linear / addmm), while the element-wise minimum and subtraction are
    fused into a single custom CUDA kernel, eliminating one full
    read-modify-write pass over the 128x16384 activation tensor and avoiding
    a host<->device synchronization for the scalar constant.
    """
    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.fused_min_sub = fused_min_sub

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