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

# Custom CUDA kernel for matrix-scalar multiplication (C = A * s).
# The op is purely memory-bound, so the kernel focuses on maximizing
# effective memory bandwidth:
#   * 128-bit vectorized loads/stores (float4) for coalesced access
#   * __restrict__ pointers so the compiler emits LDG / avoids aliasing
#   * grid-stride loop with an exact (capped) grid to cover any size
scalar_mul_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>
#include <cstdint>

namespace {

__global__ void scalar_mul_vec4_kernel(const float4* __restrict__ in,
                                       float4* __restrict__ out,
                                       const float s,
                                       const long n4) {
    long i = (long)blockIdx.x * blockDim.x + threadIdx.x;
    const long stride = (long)gridDim.x * blockDim.x;
    for (; i < n4; i += stride) {
        float4 v = in[i];
        v.x *= s;
        v.y *= s;
        v.z *= s;
        v.w *= s;
        out[i] = v;
    }
}

__global__ void scalar_mul_scalar_kernel(const float* __restrict__ in,
                                         float* __restrict__ out,
                                         const float s,
                                         const long n) {
    long i = (long)blockIdx.x * blockDim.x + threadIdx.x;
    const long stride = (long)gridDim.x * blockDim.x;
    for (; i < n; i += stride) {
        out[i] = in[i] * s;
    }
}

inline int pick_num_blocks(long work_items, int block_size) {
    long needed = (work_items + block_size - 1) / block_size;
    const long max_blocks = 2147483647L;  // gridDim.x limit on modern GPUs
    if (needed > max_blocks) needed = max_blocks;
    if (needed < 1) needed = 1;
    return (int)needed;
}

}  // namespace

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 A_contig = A.contiguous();
    auto out = torch::empty_like(A_contig);

    const long n = (long)A_contig.numel();
    if (n == 0) {
        return out;
    }

    const float sf = static_cast<float>(s);
    cudaStream_t stream = at::cuda::getCurrentCUDAStream();
    const int block_size = 256;

    const float* in_ptr = A_contig.data_ptr<float>();
    float* out_ptr = out.data_ptr<float>();

    const bool vec4_ok =
        ((n & 3L) == 0L) &&
        (reinterpret_cast<uintptr_t>(in_ptr) % 16 == 0) &&
        (reinterpret_cast<uintptr_t>(out_ptr) % 16 == 0);

    if (vec4_ok) {
        const long n4 = n >> 2;
        const int blocks = pick_num_blocks(n4, block_size);
        scalar_mul_vec4_kernel<<<blocks, block_size, 0, stream>>>(
            reinterpret_cast<const float4*>(in_ptr),
            reinterpret_cast<float4*>(out_ptr),
            sf, n4);
    } else {
        const int blocks = pick_num_blocks(n, block_size);
        scalar_mul_scalar_kernel<<<blocks, block_size, 0, stream>>>(
            in_ptr, out_ptr, sf, n);
    }

    return out;
}
"""

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

# Compile the inline CUDA extension
scalar_mul_ext = load_inline(
    name="scalar_mul_ext",
    cpp_sources=scalar_mul_cpp_source,
    cuda_sources=scalar_mul_cuda_source,
    functions=["scalar_mul_cuda"],
    verbose=False,
    extra_cuda_cflags=["-O3", "--use_fast_math"],
)


class ModelNew(nn.Module):
    """
    Optimized model that performs matrix-scalar multiplication (C = A * s)
    using a custom vectorized CUDA kernel.
    """

    def __init__(self):
        super(ModelNew, self).__init__()
        self._ext = scalar_mul_ext

    def forward(self, A: torch.Tensor, s: float) -> torch.Tensor:
        if A.is_cuda and A.dtype == torch.float32:
            return self._ext.scalar_mul_cuda(A, float(s))
        # Fallback for CPU tensors or non-float32 dtypes
        return A * s
```