```
# 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).
# This operation is purely memory-bandwidth bound, so the kernel uses
# 128-bit vectorized loads/stores (float4) to saturate DRAM bandwidth
# and minimize the number of memory transactions per element.

scalar_mul_cuda_source = """
#include <torch/extension.h>
#include <cuda_runtime.h>
#include <cstdint>

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

__global__ void scalar_mul_scalar_kernel(const float* __restrict__ A,
                                         float* __restrict__ C,
                                         const float s,
                                         const long start,
                                         const long size) {
    const long idx = start + blockIdx.x * (long)blockDim.x + threadIdx.x;
    if (idx < size) {
        C[idx] = A[idx] * s;
    }
}

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

    auto Ac = A.contiguous();
    auto C = torch::empty_like(Ac);

    const long size = Ac.numel();
    if (size == 0) {
        return C;
    }

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

    const int block_size = 256;

    const bool aligned =
        (reinterpret_cast<uintptr_t>(a_ptr) % 16 == 0) &&
        (reinterpret_cast<uintptr_t>(c_ptr) % 16 == 0);

    if (aligned) {
        // Vectorized path: process 4 floats per thread with float4.
        const long size4 = size >> 2;
        if (size4 > 0) {
            const long num_blocks = (size4 + block_size - 1) / block_size;
            scalar_mul_vec4_kernel<<<(unsigned int)num_blocks, block_size>>>(
                reinterpret_cast<const float4*>(a_ptr),
                reinterpret_cast<float4*>(c_ptr),
                sf, size4);
        }
        // Handle any remaining tail elements.
        const long tail_start = size4 << 2;
        if (tail_start < size) {
            const long rem = size - tail_start;
            const long num_blocks = (rem + block_size - 1) / block_size;
            scalar_mul_scalar_kernel<<<(unsigned int)num_blocks, block_size>>>(
                a_ptr, c_ptr, sf, tail_start, size);
        }
    } else {
        // Fallback scalar path for unaligned pointers.
        const long num_blocks = (size + block_size - 1) / block_size;
        scalar_mul_scalar_kernel<<<(unsigned int)num_blocks, block_size>>>(
            a_ptr, c_ptr, sf, 0L, size);
    }

    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=True,
    extra_cflags=[""],
    extra_ldflags=[""],
)


class ModelNew(nn.Module):
    """
    Optimized model that performs matrix-scalar multiplication (C = A * s)
    using a custom vectorized (float4) 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 CPU tensors or non-float32 dtypes.
        return A * s
```