```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>

constexpr int THREADS = 256;
constexpr int VECTORS_PER_THREAD = 8;
constexpr int TILE_VECTORS = THREADS * VECTORS_PER_THREAD;

__device__ __forceinline__ float4 streaming_load(const float4* p) {
    float4 v;
    asm volatile(
        "ld.global.cs.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 streaming_store(
    float4* p, float4 v, float s
) {
    v.x *= s;
    v.y *= s;
    v.z *= s;
    v.w *= s;
    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"
    );
}

template<bool FULL_TILES>
__global__ __launch_bounds__(THREADS)
void scalar_mul_streaming(
    const float* __restrict__ input,
    float* __restrict__ output,
    float scalar,
    int64_t count
) {
    const float4* __restrict__ src =
        reinterpret_cast<const float4*>(input);
    float4* __restrict__ dst = reinterpret_cast<float4*>(output);

    const int64_t vectors = count / 4;
    const int64_t step = int64_t(gridDim.x) * TILE_VECTORS;

    for (int64_t tile = int64_t(blockIdx.x) * TILE_VECTORS;
         tile < vectors;
         tile += step) {
        if (FULL_TILES || tile + TILE_VECTORS <= vectors) {
            const int64_t base = tile + threadIdx.x;

            #pragma unroll
            for (int group = 0; group < VECTORS_PER_THREAD / 4; ++group) {
                const int64_t i = base + group * 4 * THREADS;
                float4 a = streaming_load(src + i);
                float4 b = streaming_load(src + i + THREADS);
                float4 c = streaming_load(src + i + 2 * THREADS);
                float4 d = streaming_load(src + i + 3 * THREADS);

                streaming_store(dst + i, a, scalar);
                streaming_store(dst + i + THREADS, b, scalar);
                streaming_store(dst + i + 2 * THREADS, c, scalar);
                streaming_store(dst + i + 3 * THREADS, d, scalar);
            }
        } else {
            #pragma unroll
            for (int k = 0; k < VECTORS_PER_THREAD; ++k) {
                const int64_t i = tile + threadIdx.x + k * THREADS;
                if (i < vectors) {
                    streaming_store(dst + i, streaming_load(src + i), scalar);
                }
            }
        }
    }

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

__global__ __launch_bounds__(THREADS)
void scalar_mul_unaligned(
    const float* __restrict__ input,
    float* __restrict__ output,
    float scalar,
    int64_t count
) {
    const int64_t stride = int64_t(gridDim.x) * THREADS;
    for (int64_t i = int64_t(blockIdx.x) * THREADS + 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 a CUDA tensor");
    TORCH_CHECK(input.scalar_type() == at::kFloat, "Expected float32");
    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;
    }

    const float* src = input.data_ptr<float>();
    float* dst = output.data_ptr<float>();
    const float s = static_cast<float>(scalar);
    const auto stream =
        c10::cuda::getCurrentCUDAStream(input.get_device()).stream();
    const int max_blocks =
        at::cuda::getCurrentDeviceProperties()->multiProcessorCount * 8;

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

        if (count % TILE_ELEMENTS == 0) {
            scalar_mul_streaming<true><<<blocks, THREADS, 0, stream>>>(
                src, dst, s, count);
        } else {
            scalar_mul_streaming<false><<<blocks, THREADS, 0, stream>>>(
                src, dst, s, count);
        }
    } else {
        const int blocks = static_cast<int>(std::min<int64_t>(
            (count + THREADS - 1) / THREADS, max_blocks));
        scalar_mul_unaligned<<<blocks, THREADS, 0, stream>>>(
            src, dst, s, count);
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_extension = load_inline(
    name="matrix_scalar_mul_streaming_v6",
    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
```