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


_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>
#include <algorithm>

__device__ __forceinline__ float activate(
    float x, float bias, float subtract_value, float multiply_value
) {
    x = __fadd_rn(x, bias);
    x = __fsub_rn(x, subtract_value);
    x = __fmul_rn(x, multiply_value);
    return x < 0.0f ? 0.0f : x;
}

__global__ void epilogue_vectorized(
    float4* __restrict__ output,
    const float4* __restrict__ bias,
    int64_t count,
    int64_t columns,
    float subtract_value,
    float multiply_value
) {
    for (int64_t i = int64_t(blockIdx.x) * blockDim.x + threadIdx.x;
         i < count;
         i += int64_t(blockDim.x) * gridDim.x) {
        float4 x = output[i];
        const float4 b = bias[i % columns];

        x.x = activate(x.x, b.x, subtract_value, multiply_value);
        x.y = activate(x.y, b.y, subtract_value, multiply_value);
        x.z = activate(x.z, b.z, subtract_value, multiply_value);
        x.w = activate(x.w, b.w, subtract_value, multiply_value);

        output[i] = x;
    }
}

__global__ void epilogue_scalar(
    float* __restrict__ output,
    const float* __restrict__ bias,
    int64_t count,
    int64_t columns,
    float subtract_value,
    float multiply_value
) {
    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],
            subtract_value, multiply_value
        );
    }
}

torch::Tensor fused_epilogue(
    torch::Tensor output,
    torch::Tensor bias,
    double subtract_value,
    double multiply_value
) {
    TORCH_CHECK(output.is_cuda() && bias.is_cuda(),
                "Expected CUDA tensors");
    TORCH_CHECK(output.scalar_type() == at::kFloat &&
                bias.scalar_type() == at::kFloat,
                "Expected float32 tensors");
    TORCH_CHECK(output.device() == bias.device(),
                "Tensors must share a device");
    TORCH_CHECK(output.is_contiguous() && bias.is_contiguous(),
                "Expected contiguous tensors");
    TORCH_CHECK(output.dim() == 2 && bias.dim() == 1 &&
                output.size(1) == bias.numel(),
                "Invalid output or bias shape");

    const c10::cuda::CUDAGuard guard(output.device());
    const int64_t count = output.numel();
    if (count == 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 vector_count = count / 4;
        const int blocks = static_cast<int>(
            std::min<int64_t>((vector_count + threads - 1) / threads, 4096)
        );
        epilogue_vectorized<<<blocks, threads, 0, stream>>>(
            reinterpret_cast<float4*>(output.data_ptr<float>()),
            reinterpret_cast<const float4*>(bias.data_ptr<float>()),
            vector_count,
            columns / 4,
            static_cast<float>(subtract_value),
            static_cast<float>(multiply_value)
        );
    } else {
        const int blocks = static_cast<int>(
            std::min<int64_t>((count + threads - 1) / threads, 4096)
        );
        epilogue_scalar<<<blocks, threads, 0, stream>>>(
            output.data_ptr<float>(),
            bias.data_ptr<float>(),
            count,
            columns,
            static_cast<float>(subtract_value),
            static_cast<float>(multiply_value)
        );
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_extension = load_inline(
    name="linear_subtract_multiply_relu_fused",
    cpp_sources=r"""
        torch::Tensor fused_epilogue(
            torch::Tensor output,
            torch::Tensor bias,
            double subtract_value,
            double multiply_value
        );
    """,
    cuda_sources=_cuda_source,
    functions=["fused_epilogue"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    verbose=False,
)


class ModelNew(nn.Module):
    def __init__(
        self, in_features, out_features, subtract_value, multiply_value
    ):
        super().__init__()
        self.linear = nn.Linear(in_features, out_features)
        self.subtract_value = subtract_value
        self.multiply_value = multiply_value

    def forward(self, x):
        if (
            not x.is_cuda
            or x.dtype != torch.float32
            or self.linear.weight.dtype != torch.float32
            or torch.is_autocast_enabled()
            or (
                torch.is_grad_enabled()
                and (
                    x.requires_grad
                    or self.linear.weight.requires_grad
                    or self.linear.bias.requires_grad
                )
            )
        ):
            y = F.linear(x, self.linear.weight, self.linear.bias)
            return torch.relu(
                (y - self.subtract_value) * self.multiply_value
            )

        output_shape = (*x.shape[:-1], self.linear.out_features)
        flattened = x.reshape(-1, self.linear.in_features)
        y = torch.mm(flattened, self.linear.weight.t())
        y = _extension.fused_epilogue(
            y,
            self.linear.bias,
            float(self.subtract_value),
            float(self.multiply_value),
        )
        return y.view(output_shape)
```