```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/CUDAStream.h>
#include <c10/cuda/CUDAException.h>
#include <cuda_runtime.h>
#include <algorithm>
#include <cstdint>

__device__ __forceinline__ float4 load_streaming(const float4* p) {
    float4 v;
    asm volatile(
        "ld.global.cg.v4.f32 {%0, %1, %2, %3}, [%4];"
        : "=f"(v.x), "=f"(v.y), "=f"(v.z), "=f"(v.w)
        : "l"(p));
    return v;
}

__device__ __forceinline__ void store_scaled(
    float4* p, float4 v, float scalar)
{
    v.x *= scalar;
    v.y *= scalar;
    v.z *= scalar;
    v.w *= scalar;
    asm volatile(
        "st.global.cs.v4.f32 [%0], {%1, %2, %3, %4};"
        :
        : "l"(p), "f"(v.x), "f"(v.y), "f"(v.z), "f"(v.w)
        : "memory");
}

__global__ __launch_bounds__(256)
void scalar_mul_vectorized(
    const float* __restrict__ input,
    float* __restrict__ output,
    float scalar,
    int64_t count)
{
    constexpr int threads = 256;
    constexpr int tile = threads * 4;
    const int64_t vectors = count / 4;
    const int64_t stride = static_cast<int64_t>(gridDim.x) * tile;
    int64_t i = static_cast<int64_t>(blockIdx.x) * tile + threadIdx.x;

    const float4* src = reinterpret_cast<const float4*>(input);
    float4* dst = reinterpret_cast<float4*>(output);

    for (; i + 3 * threads < vectors; i += stride) {
        float4 a = load_streaming(src + i);
        float4 b = load_streaming(src + i + threads);
        float4 c = load_streaming(src + i + 2 * threads);
        float4 d = load_streaming(src + i + 3 * threads);

        store_scaled(dst + i, a, scalar);
        store_scaled(dst + i + threads, b, scalar);
        store_scaled(dst + i + 2 * threads, c, scalar);
        store_scaled(dst + i + 3 * threads, d, scalar);
    }

    #pragma unroll
    for (int k = 0; k < 4; ++k) {
        const int64_t j = i + k * threads;
        if (j < vectors) {
            store_scaled(dst + j, load_streaming(src + j), scalar);
        }
    }

    if (blockIdx.x == 0 && threadIdx.x < count % 4) {
        const int64_t j = vectors * 4 + threadIdx.x;
        output[j] = input[j] * scalar;
    }
}

__global__ void scalar_mul_unaligned(
    const float* __restrict__ input,
    float* __restrict__ output,
    float scalar,
    int64_t count)
{
    const int64_t stride = static_cast<int64_t>(blockDim.x) * gridDim.x;
    for (int64_t i = static_cast<int64_t>(blockIdx.x) * blockDim.x
                     + threadIdx.x;
         i < count;
         i += stride) {
        output[i] = input[i] * scalar;
    }
}

torch::Tensor scalar_mul_cuda(torch::Tensor input, double scalar) {
    TORCH_CHECK(input.is_cuda(), "Expected CUDA input");
    TORCH_CHECK(input.scalar_type() == at::kFloat, "Expected float32 input");
    TORCH_CHECK(input.is_contiguous(), "Expected contiguous input");

    const c10::cuda::CUDAGuard 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 int max_blocks =
        at::cuda::getCurrentDeviceProperties()->multiProcessorCount * 8;
    const auto stream =
        c10::cuda::getCurrentCUDAStream(input.get_device()).stream();

    const float* src = input.data_ptr<float>();
    float* dst = output.data_ptr<float>();

    if ((reinterpret_cast<uintptr_t>(src) & 15) == 0 &&
        (reinterpret_cast<uintptr_t>(dst) & 15) == 0) {
        constexpr int64_t elements_per_tile = threads * 4 * 4;
        const int blocks = static_cast<int>(std::min<int64_t>(
            max_blocks, (count + elements_per_tile - 1) / elements_per_tile));

        scalar_mul_vectorized<<<blocks, threads, 0, stream>>>(
            src, dst, static_cast<float>(scalar), count);
    } else {
        const int blocks = static_cast<int>(std::min<int64_t>(
            max_blocks, (count + threads - 1) / threads));

        scalar_mul_unaligned<<<blocks, threads, 0, stream>>>(
            src, dst, static_cast<float>(scalar), count);
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_extension = load_inline(
    name="matrix_scalar_mul_streaming_v4",
    cpp_sources="torch::Tensor scalar_mul_cuda(torch::Tensor input, double scalar);",
    cuda_sources=_cuda_source,
    functions=["scalar_mul_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 _extension.scalar_mul_cuda(A, s)
        return A * s
```