```python
# Copyright (c) Meta Platforms, Inc. and affiliates.
# All rights reserved.
#
# This source code is licensed under the license found in the
# LICENSE file in the root directory of this source tree.

import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline

scalar_mul_cpp_source = "torch::Tensor scalar_mul_cuda(torch::Tensor A, double s);"

scalar_mul_source = """
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>
#include <cstdint>
#include <algorithm>

// Vectorized (128-bit loads/stores) grid-stride kernel for bulk elements
__global__ void scalar_mul_vec4_kernel(const float4* __restrict__ A,
                                       float4* __restrict__ C,
                                       float s,
                                       int64_t n4) {
    int64_t idx = (int64_t)blockIdx.x * blockDim.x + threadIdx.x;
    int64_t stride = (int64_t)gridDim.x * blockDim.x;
    for (int64_t i = idx; i < n4; i += stride) {
        float4 a = A[i];
        a.x *= s;
        a.y *= s;
        a.z *= s;
        a.w *= s;
        C[i] = a;
    }
}

// Tail kernel for the remaining (n % 4) elements
__global__ void scalar_mul_tail_kernel(const float* __restrict__ A,
                                       float* __restrict__ C,
                                       float s,
                                       int64_t n) {
    int64_t idx = (int64_t)blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        C[idx] = A[idx] * s;
    }
}

torch::Tensor scalar_mul_cuda(torch::Tensor A, double s) {
    TORCH_CHECK(A.is_cuda(), "A must be a CUDA tensor");
    TORCH_CHECK(A.scalar_type() == at::kFloat, "A must be float32");

    auto Ac = A.contiguous();
    auto C = torch::empty_like(Ac);
    int64_t n = Ac.numel();
    float sf = static_cast<float>(s);

    if (n == 0) {
        return C;
    }

    const int block_size = 256;
    int device = Ac.get_device();
    int num_sms = 0;
    cudaDeviceGetAttribute(&num_sms, cudaDevAttrMultiProcessorCount, device);
    int max_blocks = num_sms * 16;

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

    int64_t n4 = n / 4;
    if (n4 > 0) {
        int64_t blocks_needed = (n4 + block_size - 1) / block_size;
        int blocks = (int)std::min<int64_t>(blocks_needed, (int64_t)max_blocks);
        scalar_mul_vec4_kernel<<<blocks, block_size, 0, stream>>>(
            reinterpret_cast<const float4*>(Ac.data_ptr<float>()),
            reinterpret_cast<float4*>(C.data_ptr<float>()),
            sf, n4);
    }

    int64_t rem = n % 4;
    if (rem > 0) {
        int64_t start = n4 * 4;
        int blocks = (int)((rem + block_size - 1) / block_size);
        scalar_mul_tail_kernel<<<blocks, block_size, 0, stream>>>(
            Ac.data_ptr<float>() + start,
            C.data_ptr<float>() + start,
            sf, rem);
    }

    return C;
}
"""

scalar_mul_ext = load_inline(
    name="scalar_mul_ext",
    cpp_sources=scalar_mul_cpp_source,
    cuda_sources=scalar_mul_source,
    functions=["scalar_mul_cuda"],
    verbose=False,
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Optimized matrix-scalar multiplication using a custom CUDA kernel
    with vectorized float4 memory accesses to saturate memory bandwidth.
    """
    def __init__(self):
        super(ModelNew, self).__init__()
        self._scalar_mul = scalar_mul_ext

    def forward(self, A: torch.Tensor, s: float) -> torch.Tensor:
        if A.is_cuda and A.dtype == torch.float32:
            return self._scalar_mul.scalar_mul_cuda(A, float(s))
        return A * s
```