```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).
#
# This operator is purely DRAM-bandwidth bound (reads 4 GiB, writes 4 GiB
# for the benchmark shape), so the kernel is engineered to saturate memory
# bandwidth with minimal instruction/launch overhead:
#   * 128-bit vectorized (float4) loads/stores, 8x unrolled per thread so
#     each thread keeps eight 16B transactions in flight (max MLP)
#   * streaming cache hints (__ldcs / __stcs): data has no reuse, so mark it
#     evict-first in L2 to reduce cache thrash and speed up write drain
#   * two-phase register staging: all loads are issued before any store
#   * templated "even" fast path with no bounds checks when the grid covers
#     the tensor exactly (true for the benchmark shape)
#   * 32-bit addressing fast path (covers any tensor < 64 GiB)
#   * compile-time block size via __launch_bounds__
#   * launch on the current CUDA stream (capture-safe), device guard
scalar_mul_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <c10/cuda/CUDAGuard.h>
#include <cuda_runtime.h>
#include <cstdint>

namespace {

constexpr int kBlock = 256;
constexpr int kUnroll = 8;

// Vectorized kernel. kEven == true means the grid covers n4 exactly, so all
// bounds checks are compiled out.
template <typename index_t, bool kEven>
__global__ void __launch_bounds__(kBlock)
scalar_mul_vec4_kernel(const float4* __restrict__ A,
                       float4* __restrict__ C,
                       const float s,
                       const index_t n4) {
    const index_t base = (index_t)blockIdx.x * (kBlock * kUnroll) + threadIdx.x;

    // Phase 1: issue all loads (independent, max memory-level parallelism).
    // __ldcs: streaming load, evict-first (no reuse of A).
    float4 v[kUnroll];
    #pragma unroll
    for (int j = 0; j < kUnroll; ++j) {
        const index_t i = base + (index_t)j * kBlock;
        if (kEven || i < n4) {
            v[j] = __ldcs(&A[i]);
        }
    }

    // Phase 2: multiply and store.
    // __stcs: streaming store, evict-first (C is never read by this kernel).
    #pragma unroll
    for (int j = 0; j < kUnroll; ++j) {
        const index_t i = base + (index_t)j * kBlock;
        if (kEven || i < n4) {
            float4 r;
            r.x = v[j].x * s;
            r.y = v[j].y * s;
            r.z = v[j].z * s;
            r.w = v[j].w * s;
            __stcs(&C[i], r);
        }
    }
}

// Handles the (n % 4) tail after the vectorized body.
__global__ void scalar_mul_tail_kernel(const float* __restrict__ A,
                                       float* __restrict__ C,
                                       const float s,
                                       const int64_t start,
                                       const int64_t n) {
    const int64_t idx = start + (int64_t)blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        __stcs(&C[idx], __ldcs(&A[idx]) * s);
    }
}

// Generic fallback for misaligned pointers (e.g. odd-offset contiguous views).
__global__ void __launch_bounds__(kBlock)
scalar_mul_kernel(const float* __restrict__ A,
                  float* __restrict__ C,
                  const float s,
                  const int64_t n) {
    const int64_t stride = (int64_t)gridDim.x * kBlock;
    for (int64_t i = (int64_t)blockIdx.x * kBlock + threadIdx.x; i < n; i += stride) {
        __stcs(&C[i], __ldcs(&A[i]) * s);
    }
}

}  // 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");

    const c10::cuda::CUDAGuard device_guard(A.device());

    auto Ac = A.is_contiguous() ? A : A.contiguous();
    auto C = at::empty_like(Ac);

    const int64_t n = Ac.numel();
    if (n == 0) {
        return C;
    }

    const float sf = static_cast<float>(s);
    const float* a_ptr = Ac.data_ptr<float>();
    float* c_ptr = C.data_ptr<float>();
    cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    const bool aligned =
        (((reinterpret_cast<uintptr_t>(a_ptr) |
           reinterpret_cast<uintptr_t>(c_ptr)) & 0xF) == 0);

    if (aligned && n >= 4) {
        const int64_t n4 = n >> 2;
        const int64_t per_block = (int64_t)kBlock * kUnroll;
        const int64_t blocks = (n4 + per_block - 1) / per_block;
        const bool even = (n4 % per_block) == 0;

        const float4* a4 = reinterpret_cast<const float4*>(a_ptr);
        float4* c4 = reinterpret_cast<float4*>(c_ptr);

        if (n4 <= (int64_t)0xFFFFFFFFu) {
            const unsigned int n4u = (unsigned int)n4;
            if (even) {
                scalar_mul_vec4_kernel<unsigned int, true>
                    <<<(unsigned int)blocks, kBlock, 0, stream>>>(a4, c4, sf, n4u);
            } else {
                scalar_mul_vec4_kernel<unsigned int, false>
                    <<<(unsigned int)blocks, kBlock, 0, stream>>>(a4, c4, sf, n4u);
            }
        } else {
            if (even) {
                scalar_mul_vec4_kernel<int64_t, true>
                    <<<(unsigned int)blocks, kBlock, 0, stream>>>(a4, c4, sf, n4);
            } else {
                scalar_mul_vec4_kernel<int64_t, false>
                    <<<(unsigned int)blocks, kBlock, 0, stream>>>(a4, c4, sf, n4);
            }
        }

        const int64_t rem = n & 3;
        if (rem > 0) {
            scalar_mul_tail_kernel<<<1, 32, 0, stream>>>(
                a_ptr, c_ptr, sf, n4 << 2, n);
        }
    } else {
        const int64_t full_blocks = (n + kBlock - 1) / kBlock;
        const int64_t blocks = full_blocks < 8192 ? full_blocks : 8192;
        scalar_mul_kernel<<<(unsigned int)blocks, kBlock, 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_v5",
    cpp_sources=scalar_mul_cpp_source,
    cuda_sources=scalar_mul_source,
    functions=["scalar_mul_cuda"],
    verbose=False,
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    extra_ldflags=[""],
)


class ModelNew(nn.Module):
    """
    Optimized model that performs matrix-scalar multiplication (C = A * s)
    using a custom vectorized, streaming-cache-hint CUDA kernel.
    """
    def __init__(self):
        super(ModelNew, self).__init__()
        self._fn = scalar_mul_ext.scalar_mul_cuda

    def forward(self, A: torch.Tensor, s: float) -> torch.Tensor:
        if A.is_cuda and A.dtype == torch.float32:
            return self._fn(A, float(s))
        # Fallback for non-CUDA / non-FP32 inputs
        return A * s
```