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


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

constexpr int THREADS = 256;
constexpr int VECTORS_PER_TILE = THREADS * 16;
constexpr int ELEMENTS_PER_TILE = VECTORS_PER_TILE * 4;

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

template<bool FULL_TILES>
__global__ __launch_bounds__(THREADS)
void scalar_mul_tiled(
    const float* __restrict__ input,
    float* __restrict__ output,
    const float scalar,
    const int64_t count)
{
    const int64_t tile_start =
        static_cast<int64_t>(blockIdx.x) * VECTORS_PER_TILE;
    const int64_t vector_count = count / 4;

    const float4* __restrict__ src =
        reinterpret_cast<const float4*>(input);
    float4* __restrict__ dst = reinterpret_cast<float4*>(output);

    if (FULL_TILES || tile_start + VECTORS_PER_TILE <= vector_count) {
        const float4* tile_src = src + tile_start + threadIdx.x;
        float4* tile_dst = dst + tile_start + threadIdx.x;

        #pragma unroll
        for (int group = 0; group < 4; ++group) {
            const int offset = group * 4 * THREADS;

            float4 a = __ldg(tile_src + offset);
            float4 b = __ldg(tile_src + offset + THREADS);
            float4 c = __ldg(tile_src + offset + 2 * THREADS);
            float4 d = __ldg(tile_src + offset + 3 * THREADS);

            tile_dst[offset] = scale(a, scalar);
            tile_dst[offset + THREADS] = scale(b, scalar);
            tile_dst[offset + 2 * THREADS] = scale(c, scalar);
            tile_dst[offset + 3 * THREADS] = scale(d, scalar);
        }
    } else {
        #pragma unroll
        for (int k = 0; k < 16; ++k) {
            const int64_t i = tile_start + threadIdx.x + k * THREADS;
            if (i < vector_count) {
                dst[i] = scale(__ldg(src + i), scalar);
            }
        }
    }

    if (!FULL_TILES && blockIdx.x == gridDim.x - 1) {
        const int64_t i = vector_count * 4 + threadIdx.x;
        if (i < count) {
            output[i] = input[i] * scalar;
        }
    }
}

__global__ void scalar_mul_unaligned(
    const float* __restrict__ input,
    float* __restrict__ output,
    const float scalar,
    const int64_t count)
{
    const int64_t start =
        static_cast<int64_t>(blockIdx.x) * THREADS * 8 + threadIdx.x;

    #pragma unroll
    for (int k = 0; k < 8; ++k) {
        const int64_t i = start + k * THREADS;
        if (i < count) {
            output[i] = input[i] * scalar;
        }
    }
}

torch::Tensor scalar_mul_cuda(torch::Tensor input, double scalar) {
    TORCH_CHECK(input.is_cuda(), "Expected a CUDA tensor");
    TORCH_CHECK(input.scalar_type() == at::kFloat, "Expected float32");
    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;
    }

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

    if (((reinterpret_cast<uintptr_t>(src) |
          reinterpret_cast<uintptr_t>(dst)) & 15) == 0) {
        const unsigned int blocks = static_cast<unsigned int>(
            (count + ELEMENTS_PER_TILE - 1) / ELEMENTS_PER_TILE);

        if (count % ELEMENTS_PER_TILE == 0) {
            scalar_mul_tiled<true><<<blocks, THREADS, 0, stream>>>(
                src, dst, s, count);
        } else {
            scalar_mul_tiled<false><<<blocks, THREADS, 0, stream>>>(
                src, dst, s, count);
        }
    } else {
        const unsigned int blocks = static_cast<unsigned int>(
            (count + THREADS * 8 - 1) / (THREADS * 8));
        scalar_mul_unaligned<<<blocks, THREADS, 0, stream>>>(
            src, dst, s, count);
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_extension = load_inline(
    name="matrix_scalar_mul_tiled_v5",
    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
```