```python
import torch
import torch.nn as nn
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>
#include <cstdint>

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

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

    for (int64_t i = index; i < vectors; i += stride) {
        float4 value = input4[i];
        value.x *= scale;
        value.y *= scale;
        value.z *= scale;
        value.w *= scale;
        output4[i] = value;
    }

    const int64_t tail = vectors * 4 + index;
    if (tail < count) {
        output[tail] = input[tail] * scale;
    }
}

__global__ void scale_scalar(
    const float* __restrict__ input,
    float* __restrict__ output,
    const float scale,
    const int64_t count)
{
    const int64_t index =
        static_cast<int64_t>(blockIdx.x) * blockDim.x + threadIdx.x;
    const int64_t stride =
        static_cast<int64_t>(blockDim.x) * gridDim.x;

    for (int64_t i = index; i < count; i += stride) {
        output[i] = input[i] * scale;
    }
}

torch::Tensor matrix_scale_cuda(torch::Tensor input, double scale) {
    TORCH_CHECK(input.is_cuda(), "Expected a CUDA tensor");
    TORCH_CHECK(input.scalar_type() == at::kFloat, "Expected float32");
    TORCH_CHECK(input.is_contiguous(), "Expected contiguous input");

    const c10::cuda::CUDAGuard device_guard(input.device());
    auto output = torch::empty(input.sizes(), input.options());
    const int64_t count = input.numel();

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

    const float* input_ptr = input.data_ptr<float>();
    float* output_ptr = output.data_ptr<float>();
    const bool aligned =
        (reinterpret_cast<uintptr_t>(input_ptr) % 16 == 0) &&
        (reinterpret_cast<uintptr_t>(output_ptr) % 16 == 0);

    constexpr int threads = 256;
    const int64_t work = aligned ? (count + 3) / 4 : count;
    const int blocks = static_cast<int>(
        std::min<int64_t>((work + threads - 1) / threads, 8192));
    const cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    if (aligned) {
        scale_vectorized<<<blocks, threads, 0, stream>>>(
            input_ptr, output_ptr, static_cast<float>(scale), count);
    } else {
        scale_scalar<<<blocks, threads, 0, stream>>>(
            input_ptr, output_ptr, static_cast<float>(scale), count);
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_scale_extension = load_inline(
    name="matrix_scalar_scale_vectorized",
    cpp_sources=(
        "torch::Tensor matrix_scale_cuda(torch::Tensor input, double scale);"
    ),
    cuda_sources=_cuda_source,
    functions=["matrix_scale_cuda"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    verbose=False,
)


class ModelNew(nn.Module):
    def __init__(self):
        super().__init__()

    def forward(self, A: torch.Tensor, s: float) -> torch.Tensor:
        if (
            A.is_cuda
            and A.dtype == torch.float32
            and A.is_contiguous()
            and not (torch.is_grad_enabled() and A.requires_grad)
        ):
            return _scale_extension.matrix_scale_cuda(A, s)
        return A * s
```