```python
import math
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 <cuda_fp16.h>
#include <cublas_v2.h>
#include <algorithm>
#include <climits>
#include <cstdint>

template<int Scale>
__global__ void pack_half(
    const float* __restrict__ src,
    half* __restrict__ dst,
    int64_t count
) {
    const int64_t vectors = (count + 3) / 4;
    for (int64_t t = int64_t(blockIdx.x) * blockDim.x + threadIdx.x;
         t < vectors;
         t += int64_t(blockDim.x) * gridDim.x) {
        const int64_t i = t * 4;
        if (i + 3 < count &&
            (reinterpret_cast<uintptr_t>(src + i) & 15) == 0) {
            float4 v = *reinterpret_cast<const float4*>(src + i);
            half2* d = reinterpret_cast<half2*>(dst + i);
            d[0] = __floats2half2_rn(v.x * Scale, v.y * Scale);
            d[1] = __floats2half2_rn(v.z * Scale, v.w * Scale);
        } else {
            #pragma unroll
            for (int j = 0; j < 4; ++j) {
                if (i + j < count)
                    dst[i + j] = __float2half_rn(src[i + j] * Scale);
            }
        }
    }
}

torch::Tensor pack_weight_cuda(torch::Tensor weight) {
    TORCH_CHECK(weight.is_cuda() &&
                weight.scalar_type() == at::kFloat &&
                weight.dim() == 2, "Expected a float32 CUDA matrix");
    const c10::cuda::CUDAGuard guard(weight.device());
    weight = weight.contiguous();
    auto packed = torch::empty(
        weight.sizes(), weight.options().dtype(at::kHalf)
    );
    if (weight.numel() != 0) {
        const int64_t vectors = (weight.numel() + 3) / 4;
        const int blocks = static_cast<int>(
            std::min<int64_t>((vectors + 255) / 256, 8192)
        );
        pack_half<256><<<blocks, 256, 0,
                        at::cuda::getCurrentCUDAStream()>>>(
            weight.data_ptr<float>(),
            reinterpret_cast<half*>(packed.data_ptr<at::Half>()),
            weight.numel()
        );
        C10_CUDA_KERNEL_LAUNCH_CHECK();
    }
    return packed;
}

__device__ __forceinline__ float activate(
    float value, float bias, float sub, float mul
) {
    value = __fadd_rn(value, bias);
    value = __fsub_rn(value, sub);
    value = __fmul_rn(value, mul);
    return value < 0.0f ? 0.0f : value;
}

__global__ void epilogue4(
    float4* __restrict__ output,
    const float4* __restrict__ bias,
    int64_t count,
    int columns,
    float sub,
    float mul
) {
    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, sub, mul);
        v.y = activate(v.y, b.y, sub, mul);
        v.z = activate(v.z, b.z, sub, mul);
        v.w = activate(v.w, b.w, sub, mul);
        output[i] = v;
    }
}

__global__ void epilogue1(
    float* __restrict__ output,
    const float* __restrict__ bias,
    int64_t count,
    int columns,
    float sub,
    float mul
) {
    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], sub, mul);
    }
}

torch::Tensor linear_activate_cuda(
    torch::Tensor x,
    torch::Tensor packed_weight,
    torch::Tensor bias,
    double subtract_value,
    double multiply_value
) {
    TORCH_CHECK(x.is_cuda() && packed_weight.is_cuda() && bias.is_cuda(),
                "Expected CUDA tensors");
    TORCH_CHECK(x.scalar_type() == at::kFloat &&
                packed_weight.scalar_type() == at::kHalf &&
                bias.scalar_type() == at::kFloat, "Invalid tensor dtypes");
    TORCH_CHECK(x.device() == packed_weight.device() &&
                x.device() == bias.device(), "Device mismatch");
    TORCH_CHECK(x.dim() == 2 && packed_weight.dim() == 2 &&
                bias.dim() == 1 &&
                x.size(1) == packed_weight.size(1) &&
                bias.numel() == packed_weight.size(0), "Invalid shapes");
    TORCH_CHECK(packed_weight.is_contiguous(),
                "Packed weight must be contiguous");
    TORCH_CHECK(x.size(0) <= INT_MAX &&
                packed_weight.size(0) <= INT_MAX &&
                x.size(1) <= INT_MAX, "Matrix dimensions are too large");

    const c10::cuda::CUDAGuard guard(x.device());
    x = x.contiguous();
    bias = bias.contiguous();

    const int m = static_cast<int>(x.size(0));
    const int n = static_cast<int>(packed_weight.size(0));
    const int k = static_cast<int>(x.size(1));
    auto output = torch::empty({m, n}, x.options());
    if (m == 0 || n == 0)
        return output;

    cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    if (k == 0) {
        C10_CUDA_CHECK(cudaMemsetAsync(
            output.data_ptr<float>(), 0,
            output.numel() * sizeof(float), stream
        ));
    } else {
        auto xh = torch::empty(x.sizes(), x.options().dtype(at::kHalf));
        const int64_t vectors = (x.numel() + 3) / 4;
        const int blocks = static_cast<int>(
            std::min<int64_t>((vectors + 255) / 256, 4096)
        );
        pack_half<1><<<blocks, 256, 0, stream>>>(
            x.data_ptr<float>(),
            reinterpret_cast<half*>(xh.data_ptr<at::Half>()),
            x.numel()
        );
        C10_CUDA_KERNEL_LAUNCH_CHECK();

        cublasHandle_t handle = at::cuda::getCurrentCUDABlasHandle();
        cublasMath_t old_math;
        cublasPointerMode_t old_pointer_mode;
        TORCH_CHECK(cublasGetMathMode(handle, &old_math) ==
                    CUBLAS_STATUS_SUCCESS, "Cannot read cuBLAS math mode");
        TORCH_CHECK(cublasGetPointerMode(handle, &old_pointer_mode) ==
                    CUBLAS_STATUS_SUCCESS, "Cannot read cuBLAS pointer mode");

        const cublasMath_t math_mode = static_cast<cublasMath_t>(
            CUBLAS_DEFAULT_MATH |
            CUBLAS_MATH_DISALLOW_REDUCED_PRECISION_REDUCTION
        );
        TORCH_CHECK(cublasSetMathMode(handle, math_mode) ==
                    CUBLAS_STATUS_SUCCESS, "Cannot set cuBLAS math mode");

        const auto pointer_status =
            cublasSetPointerMode(handle, CUBLAS_POINTER_MODE_HOST);
        if (pointer_status != CUBLAS_STATUS_SUCCESS) {
            cublasSetMathMode(handle, old_math);
            TORCH_CHECK(false, "Cannot set cuBLAS pointer mode");
        }

        const float alpha = 1.0f / 256.0f;
        const float beta = 0.0f;
        const auto status = cublasGemmEx(
            handle, CUBLAS_OP_T, CUBLAS_OP_N,
            n, m, k,
            &alpha,
            packed_weight.data_ptr<at::Half>(), CUDA_R_16F, k,
            xh.data_ptr<at::Half>(), CUDA_R_16F, k,
            &beta,
            output.data_ptr<float>(), CUDA_R_32F, n,
            CUBLAS_COMPUTE_32F,
            CUBLAS_GEMM_DEFAULT_TENSOR_OP
        );

        const auto restore_pointer =
            cublasSetPointerMode(handle, old_pointer_mode);
        const auto restore_math = cublasSetMathMode(handle, old_math);
        TORCH_CHECK(status == CUBLAS_STATUS_SUCCESS,
                    "cuBLAS GEMM failed: ", static_cast<int>(status));
        TORCH_CHECK(restore_pointer == CUBLAS_STATUS_SUCCESS &&
                    restore_math == CUBLAS_STATUS_SUCCESS,
                    "Cannot restore cuBLAS state");
    }

    const int64_t count = output.numel();
    if (n % 4 == 0 &&
        (reinterpret_cast<uintptr_t>(bias.data_ptr<float>()) & 15) == 0) {
        const int64_t vectors = count / 4;
        const int blocks = static_cast<int>(
            std::min<int64_t>((vectors + 255) / 256, 4096)
        );
        epilogue4<<<blocks, 256, 0, stream>>>(
            reinterpret_cast<float4*>(output.data_ptr<float>()),
            reinterpret_cast<const float4*>(bias.data_ptr<float>()),
            vectors, n / 4,
            static_cast<float>(subtract_value),
            static_cast<float>(multiply_value)
        );
    } else {
        const int blocks = static_cast<int>(
            std::min<int64_t>((count + 255) / 256, 4096)
        );
        epilogue1<<<blocks, 256, 0, stream>>>(
            output.data_ptr<float>(), bias.data_ptr<float>(),
            count, n,
            static_cast<float>(subtract_value),
            static_cast<float>(multiply_value)
        );
    }
    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_extension = load_inline(
    name="linear_versioned_packed_tensorcore_v4",
    cpp_sources=r"""
torch::Tensor pack_weight_cuda(torch::Tensor weight);
torch::Tensor linear_activate_cuda(
    torch::Tensor x,
    torch::Tensor packed_weight,
    torch::Tensor bias,
    double subtract_value,
    double multiply_value
);
""",
    cuda_sources=_cuda_source,
    functions=["pack_weight_cuda", "linear_activate_cuda"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    extra_ldflags=["-lcublas"],
    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
        self.register_buffer("_packed_weight", None, persistent=False)
        self._packed_weight_key = None

    def _apply(self, fn, recurse=True):
        self._packed_weight = None
        self._packed_weight_key = None
        return super()._apply(fn, recurse=recurse)

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

        if x.ndim == 0 or x.shape[-1] != self.linear.in_features:
            raise ValueError("Input's last dimension must match in_features")
        if x.device != weight.device or x.device != bias.device:
            raise ValueError("Input and parameters must share a device")

        # Reuse only the parameter representation, never input-dependent data.
        # Normal parameter updates and state_dict loads increment _version.
        try:
            version = weight._version
        except RuntimeError:
            version = None

        stream = torch.cuda.current_stream(weight.device)
        key = (
            id(weight),
            weight.data_ptr(),
            version,
            tuple(weight.shape),
            tuple(weight.stride()),
            weight.device,
            stream.cuda_stream,
        )

        if (
            version is None
            or self._packed_weight is None
            or self._packed_weight_key != key
        ):
            packed_weight = _extension.pack_weight_cuda(weight)
            if version is not None:
                self._packed_weight = packed_weight
                self._packed_weight_key = key
        else:
            packed_weight = self._packed_weight

        shape = (*x.shape[:-1], self.linear.out_features)
        rows = math.prod(x.shape[:-1])
        return _extension.linear_activate_cuda(
            x.reshape(rows, self.linear.in_features),
            packed_weight,
            bias,
            float(self.subtract_value),
            float(self.multiply_value),
        ).view(shape)
```