```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 diag(A) @ B, which is mathematically equivalent to
# scaling each row i of B by A[i]. This avoids materializing the dense N x N
# diagonal matrix and the full O(N^2 * M) matmul, reducing it to an
# O(N * M) memory-bound elementwise operation with vectorized float4 access.
diag_matmul_source = """
#include <torch/extension.h>
#include <cuda_runtime.h>

// Vectorized kernel: each thread handles 4 consecutive floats (float4).
// blockIdx.y indexes the row (with a stride loop for very large N),
// threads along x cover M/4 float4 columns.
__global__ void diag_matmul_vec4_kernel(
    const float* __restrict__ A,
    const float* __restrict__ B,
    float* __restrict__ out,
    int N,
    int M4)
{
    int col4 = blockIdx.x * blockDim.x + threadIdx.x;
    if (col4 < M4) {
        for (int row = blockIdx.y; row < N; row += gridDim.y) {
            float a = A[row];
            int idx = row * M4 + col4;
            float4 b = reinterpret_cast<const float4*>(B)[idx];
            float4 o;
            o.x = a * b.x;
            o.y = a * b.y;
            o.z = a * b.z;
            o.w = a * b.w;
            reinterpret_cast<float4*>(out)[idx] = o;
        }
    }
}

// Scalar fallback kernel for the case where M is not divisible by 4.
__global__ void diag_matmul_scalar_kernel(
    const float* __restrict__ A,
    const float* __restrict__ B,
    float* __restrict__ out,
    int N,
    int M)
{
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (col < M) {
        for (int row = blockIdx.y; row < N; row += gridDim.y) {
            int idx = row * M + col;
            out[idx] = A[row] * B[idx];
        }
    }
}

torch::Tensor diag_matmul_cuda(torch::Tensor A, torch::Tensor B) {
    TORCH_CHECK(A.is_cuda() && B.is_cuda(), "Inputs must be CUDA tensors");
    TORCH_CHECK(A.dim() == 1, "A must be a 1D tensor");
    TORCH_CHECK(B.dim() == 2, "B must be a 2D tensor");
    TORCH_CHECK(A.size(0) == B.size(0), "A length must match B rows");
    TORCH_CHECK(A.scalar_type() == torch::kFloat32 &&
                B.scalar_type() == torch::kFloat32,
                "Inputs must be float32");

    auto A_c = A.contiguous();
    auto B_c = B.contiguous();

    const int64_t N = B_c.size(0);
    const int64_t M = B_c.size(1);

    auto out = torch::empty_like(B_c);
    if (N == 0 || M == 0) {
        return out;
    }

    const int block_size = 256;
    const unsigned int grid_y = (unsigned int)std::min<int64_t>(N, 65535);

    cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    if (M % 4 == 0) {
        const int M4 = (int)(M / 4);
        dim3 grid((M4 + block_size - 1) / block_size, grid_y);
        diag_matmul_vec4_kernel<<<grid, block_size, 0, stream>>>(
            A_c.data_ptr<float>(),
            B_c.data_ptr<float>(),
            out.data_ptr<float>(),
            (int)N,
            M4);
    } else {
        dim3 grid(((int)M + block_size - 1) / block_size, grid_y);
        diag_matmul_scalar_kernel<<<grid, block_size, 0, stream>>>(
            A_c.data_ptr<float>(),
            B_c.data_ptr<float>(),
            out.data_ptr<float>(),
            (int)N,
            (int)M);
    }

    return out;
}
"""

diag_matmul_cpp_source = (
    "torch::Tensor diag_matmul_cuda(torch::Tensor A, torch::Tensor B);"
)

# Compile the inline CUDA code for the fused diag-matmul (row scaling) kernel
diag_matmul = load_inline(
    name="diag_matmul",
    cpp_sources=diag_matmul_cpp_source,
    cuda_sources=diag_matmul_source,
    functions=["diag_matmul_cuda"],
    verbose=True,
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    extra_ldflags=[""],
)


class ModelNew(nn.Module):
    """
    Optimized model that performs diag(A) @ B via a custom CUDA kernel
    that scales each row of B by the corresponding element of A.
    """
    def __init__(self):
        super(ModelNew, self).__init__()
        self.diag_matmul = diag_matmul

    def forward(self, A, B):
        """
        Args:
            A (torch.Tensor): 1D tensor, diagonal entries. Shape: (N,).
            B (torch.Tensor): 2D tensor. Shape: (N, M).

        Returns:
            torch.Tensor: Result of diag(A) @ B. Shape: (N, M).
        """
        if A.is_cuda and B.is_cuda and A.dtype == torch.float32 and B.dtype == torch.float32:
            return self.diag_matmul.diag_matmul_cuda(A, B)
        # Fallback for CPU or non-float32 inputs (mathematically identical).
        return A.unsqueeze(1) * B
```