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


_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <c10/cuda/CUDAGuard.h>
#include <c10/cuda/CUDAStream.h>
#include <c10/cuda/CUDAException.h>
#include <cuda_runtime.h>
#include <algorithm>
#include <cstdint>

__device__ __forceinline__ float4 load_vector(const float4* ptr) {
    float4 v;
    asm volatile(
        "ld.global.cg.v4.f32 {%0, %1, %2, %3}, [%4];"
        : "=f"(v.x), "=f"(v.y), "=f"(v.z), "=f"(v.w)
        : "l"(ptr));
    return v;
}

__device__ __forceinline__ float4 multiply(float4 v, float s) {
    v.x *= s;
    v.y *= s;
    v.z *= s;
    v.w *= s;
    return v;
}

template<int THREADS>
__global__ void scalar_mul_chunks(
    const float* __restrict__ input,
    float* __restrict__ output,
    float scalar,
    int64_t count,
    int64_t chunk_vectors)
{
    const int64_t vectors = count / 4;
    const int64_t begin = static_cast<int64_t>(blockIdx.x) * chunk_vectors;
    const int64_t limit = begin + chunk_vectors;
    const int64_t end = limit < vectors ? limit : vectors;

    const float4* src = reinterpret_cast<const float4*>(input);
    float4* dst = reinterpret_cast<float4*>(output);
    int64_t i = begin + threadIdx.x;

    for (; i + 3 * THREADS < end; i += 4 * THREADS) {
        float4 a = load_vector(src + i);
        float4 b = load_vector(src + i + THREADS);
        float4 c = load_vector(src + i + 2 * THREADS);
        float4 d = load_vector(src + i + 3 * THREADS);

        dst[i] = multiply(a, scalar);
        dst[i + THREADS] = multiply(b, scalar);
        dst[i + 2 * THREADS] = multiply(c, scalar);
        dst[i + 3 * THREADS] = multiply(d, scalar);
    }

    for (; i < end; i += THREADS) {
        dst[i] = multiply(load_vector(src + i), scalar);
    }

    if (blockIdx.x == 0 && threadIdx.x < count % 4) {
        const int64_t tail = vectors * 4 + threadIdx.x;
        output[tail] = input[tail] * scalar;
    }
}

__global__ void scalar_mul_unaligned(
    const float* __restrict__ input,
    float* __restrict__ output,
    float scalar,
    int64_t count,
    int64_t chunk_elements)
{
    const int64_t begin = static_cast<int64_t>(blockIdx.x) * chunk_elements;
    const int64_t limit = begin + chunk_elements;
    const int64_t end = limit < count ? limit : count;

    for (int64_t i = begin + threadIdx.x; i < end; i += blockDim.x) {
        output[i] = input[i] * scalar;
    }
}

torch::Tensor scalar_mul_cuda(torch::Tensor input, double scalar) {
    TORCH_CHECK(input.is_cuda(), "Expected CUDA input");
    TORCH_CHECK(input.scalar_type() == at::kFloat, "Expected float32 input");
    TORCH_CHECK(input.is_contiguous(), "Expected contiguous input");

    const c10::cuda::CUDAGuard guard(input.device());
    auto output = torch::empty_like(input);
    const int64_t count = input.numel();
    if (count == 0) return output;

    constexpr int threads = 256;
    const int64_t max_blocks =
        static_cast<int64_t>(
            at::cuda::getCurrentDeviceProperties()->multiProcessorCount) * 8;

    const float* src = input.data_ptr<float>();
    float* dst = output.data_ptr<float>();
    const auto stream =
        c10::cuda::getCurrentCUDAStream(input.get_device()).stream();

    if ((reinterpret_cast<uintptr_t>(src) & 15) == 0) {
        constexpr int64_t tile = threads * 4;
        const int64_t vectors = count / 4;
        const int64_t tiles = std::max<int64_t>(
            1, (vectors + tile - 1) / tile);
        const int64_t target_blocks = std::min(tiles, max_blocks);
        const int64_t tiles_per_block =
            (tiles + target_blocks - 1) / target_blocks;
        const int64_t chunk = tiles_per_block * tile;
        const int blocks = static_cast<int>(
            (tiles + tiles_per_block - 1) / tiles_per_block);

        scalar_mul_chunks<threads><<<blocks, threads, 0, stream>>>(
            src, dst, static_cast<float>(scalar), count, chunk);
    } else {
        const int64_t tiles = (count + threads - 1) / threads;
        const int64_t target_blocks = std::min(tiles, max_blocks);
        const int64_t tiles_per_block =
            (tiles + target_blocks - 1) / target_blocks;
        const int64_t chunk = tiles_per_block * threads;
        const int blocks = static_cast<int>(
            (tiles + tiles_per_block - 1) / tiles_per_block);

        scalar_mul_unaligned<<<blocks, threads, 0, stream>>>(
            src, dst, static_cast<float>(scalar), count, chunk);
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_extension = load_inline(
    name="matrix_scalar_mul_contiguous_chunks_v3",
    cpp_sources="torch::Tensor scalar_mul_cuda(torch::Tensor input, double scalar);",
    cuda_sources=_cuda_source,
    functions=["scalar_mul_cuda"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    verbose=False,
)


class ModelNew(nn.Module):
    def __init__(self):
        super().__init__()

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