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


_cpp_source = r"""
#include <torch/extension.h>

torch::Tensor scale_leaky_relu_cuda(
    torch::Tensor input,
    double multiplier,
    double negative_slope);
"""

_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/Dispatch.h>
#include <ATen/OpMathType.h>
#include <ATen/cuda/CUDAContext.h>
#include <c10/cuda/CUDAGuard.h>
#include <c10/cuda/CUDAException.h>
#include <cuda_runtime.h>
#include <algorithm>
#include <cstdint>

__global__ void scale_leaky_float4_kernel(
    const float* __restrict__ input,
    float* __restrict__ output,
    int64_t count,
    float multiplier,
    float negative_slope)
{
    const int64_t index =
        static_cast<int64_t>(blockIdx.x) * blockDim.x + threadIdx.x;
    const int64_t stride =
        static_cast<int64_t>(gridDim.x) * blockDim.x;
    const int64_t vectors = count / 4;

    const float4* input4 = reinterpret_cast<const float4*>(input);
    float4* output4 = reinterpret_cast<float4*>(output);

    for (int64_t i = index; i < vectors; i += stride) {
        float4 value = input4[i];

        value.x *= multiplier;
        value.y *= multiplier;
        value.z *= multiplier;
        value.w *= multiplier;

        value.x = value.x > 0.0f ? value.x : value.x * negative_slope;
        value.y = value.y > 0.0f ? value.y : value.y * negative_slope;
        value.z = value.z > 0.0f ? value.z : value.z * negative_slope;
        value.w = value.w > 0.0f ? value.w : value.w * negative_slope;

        output4[i] = value;
    }

    for (int64_t i = vectors * 4 + index; i < count; i += stride) {
        float value = input[i] * multiplier;
        output[i] = value > 0.0f ? value : value * negative_slope;
    }
}

template <typename scalar_t>
__global__ void scale_leaky_kernel(
    const scalar_t* __restrict__ input,
    scalar_t* __restrict__ output,
    int64_t count,
    at::opmath_type<scalar_t> multiplier,
    at::opmath_type<scalar_t> negative_slope)
{
    using acc_t = at::opmath_type<scalar_t>;

    const int64_t index =
        static_cast<int64_t>(blockIdx.x) * blockDim.x + threadIdx.x;
    const int64_t stride =
        static_cast<int64_t>(gridDim.x) * blockDim.x;

    for (int64_t i = index; i < count; i += stride) {
        // Preserve the intermediate rounding of the separate multiply.
        scalar_t scaled = static_cast<scalar_t>(
            static_cast<acc_t>(input[i]) * multiplier);
        acc_t value = static_cast<acc_t>(scaled);
        output[i] = static_cast<scalar_t>(
            value > acc_t(0) ? value : value * negative_slope);
    }
}

torch::Tensor scale_leaky_relu_cuda(
    torch::Tensor input,
    double multiplier,
    double negative_slope)
{
    TORCH_CHECK(input.is_cuda(), "input must be a CUDA tensor");
    TORCH_CHECK(input.is_contiguous(), "input must be contiguous");

    const c10::cuda::CUDAGuard device_guard(input.device());
    auto output = torch::empty_like(input);
    const int64_t count = input.numel();

    if (count == 0) {
        return output;
    }

    constexpr int threads = 256;
    const cudaStream_t stream =
        at::cuda::getCurrentCUDAStream(input.get_device());

    if (input.scalar_type() == at::ScalarType::Float) {
        const int64_t work = (count + 3) / 4;
        const int blocks = static_cast<int>(
            std::min<int64_t>((work + threads - 1) / threads, 4096));

        scale_leaky_float4_kernel<<<blocks, threads, 0, stream>>>(
            input.data_ptr<float>(),
            output.data_ptr<float>(),
            count,
            static_cast<float>(multiplier),
            static_cast<float>(negative_slope));
    } else {
        const int blocks = static_cast<int>(
            std::min<int64_t>((count + threads - 1) / threads, 4096));

        AT_DISPATCH_FLOATING_TYPES_AND2(
            at::ScalarType::Half,
            at::ScalarType::BFloat16,
            input.scalar_type(),
            "scale_leaky_relu_cuda",
            [&] {
                using acc_t = at::opmath_type<scalar_t>;
                scale_leaky_kernel<scalar_t>
                    <<<blocks, threads, 0, stream>>>(
                        input.data_ptr<scalar_t>(),
                        output.data_ptr<scalar_t>(),
                        count,
                        static_cast<acc_t>(multiplier),
                        static_cast<acc_t>(negative_slope));
            });
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_fused_ops = load_inline(
    name="gemm_scale_leaky_relu_fused_v1",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["scale_leaky_relu_cuda"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    verbose=False,
)


class ModelNew(nn.Module):
    def __init__(
        self, in_features, out_features, multiplier, negative_slope
    ):
        super().__init__()
        self.gemm = nn.Linear(in_features, out_features)
        self.multiplier = multiplier
        self.leaky_relu = nn.LeakyReLU(negative_slope)

    def forward(self, x):
        x = self.gemm(x)
        if not x.is_cuda or (torch.is_grad_enabled() and x.requires_grad):
            return F.leaky_relu(
                x * self.multiplier,
                negative_slope=self.leaky_relu.negative_slope,
            )
        return _fused_ops.scale_leaky_relu_cuda(
            x.contiguous(),
            self.multiplier,
            self.leaky_relu.negative_slope,
        )
```