```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 operation is purely memory-bound, so the kernel is optimized with:
#   1. float4 vectorized loads/stores (128-bit) when the tensor size allows it.
#   2. A grid-stride loop to handle very large tensors safely (64-bit indexing).
scalar_mul_cuda_source = """
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>
#include <cstdint>

__global__ void scalar_mul_vec4_kernel(const float4* __restrict__ A,
                                       float4* __restrict__ C,
                                       float s,
                                       long n_vec) {
    long idx = (long)blockIdx.x * blockDim.x + threadIdx.x;
    long stride = (long)gridDim.x * blockDim.x;
    for (long i = idx; i < n_vec; i += stride) {
        float4 a = A[i];
        a.x *= s;
        a.y *= s;
        a.z *= s;
        a.w *= s;
        C[i] = a;
    }
}

__global__ void scalar_mul_kernel(const float* __restrict__ A,
                                  float* __restrict__ C,
                                  float s,
                                  long n) {
    long idx = (long)blockIdx.x * blockDim.x + threadIdx.x;
    long stride = (long)gridDim.x * blockDim.x;
    for (long i = idx; i < n; i += stride) {
        C[i] = A[i] * 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 A_contig = A.contiguous();
    auto C = torch::empty_like(A_contig);
    long n = (long)A_contig.numel();
    if (n == 0) {
        return C;
    }

    const float* a_ptr = A_contig.data_ptr<float>();
    float* c_ptr = C.data_ptr<float>();
    float sf = (float)s;

    const int block_size = 256;
    const long max_blocks = 1L << 30;  // well within the 2^31 - 1 x-grid limit
    cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    bool aligned = (((uintptr_t)a_ptr & 0xF) == 0) && (((uintptr_t)c_ptr & 0xF) == 0);

    if ((n & 3) == 0 && aligned) {
        long n_vec = n >> 2;
        long num_blocks = (n_vec + block_size - 1) / block_size;
        if (num_blocks > max_blocks) num_blocks = max_blocks;
        scalar_mul_vec4_kernel<<<(unsigned int)num_blocks, block_size, 0, stream>>>(
            reinterpret_cast<const float4*>(a_ptr),
            reinterpret_cast<float4*>(c_ptr),
            sf, n_vec);
    } else {
        long num_blocks = (n + block_size - 1) / block_size;
        if (num_blocks > max_blocks) num_blocks = max_blocks;
        scalar_mul_kernel<<<(unsigned int)num_blocks, block_size, 0, stream>>>(
            a_ptr, c_ptr, sf, n);
    }

    return C;
}
"""

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"],
)


class ModelNew(nn.Module):
    """
    Simple model that performs a matrix-scalar multiplication (C = A * s)
    using a custom vectorized CUDA kernel.
    """
    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))
        # Fallback for non-CUDA / non-float32 inputs
        return A * s
```