```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

# Key insight: torch.diag(A) @ B simply scales row i of B by A[i]:
#     out[i, j] = A[i] * B[i, j]
# This fused CUDA kernel computes that directly in O(N*M) work and memory,
# avoiding the O(N^2) materialization of diag(A) and the O(N^2*M) matmul.
diag_matmul_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>
#include <limits>

template <typename index_t>
__global__ void diag_matmul_scalar_kernel(
    const float* __restrict__ A,
    const float* __restrict__ B,
    float* __restrict__ out,
    const index_t M,
    const index_t total)
{
    index_t idx = (index_t)blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < total) {
        index_t row = idx / M;
        out[idx] = __ldg(A + row) * B[idx];
    }
}

template <typename index_t>
__global__ void diag_matmul_vec4_kernel(
    const float* __restrict__ A,
    const float* __restrict__ B,
    float* __restrict__ out,
    const index_t M,
    const index_t total4)
{
    index_t idx4 = (index_t)blockIdx.x * blockDim.x + threadIdx.x;
    if (idx4 < total4) {
        index_t idx = idx4 << 2;
        index_t row = idx / M;  // M % 4 == 0, so a float4 never crosses rows
        float a = __ldg(A + row);
        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;
    }
}

torch::Tensor diag_matmul_cuda(torch::Tensor A, torch::Tensor B) {
    TORCH_CHECK(A.is_cuda() && B.is_cuda(), "A and B must be CUDA tensors");
    TORCH_CHECK(A.scalar_type() == at::kFloat && B.scalar_type() == at::kFloat,
                "A and B must be float32");
    TORCH_CHECK(A.dim() == 1 && B.dim() == 2, "A must be 1D, B must be 2D");
    TORCH_CHECK(A.size(0) == B.size(0), "A.size(0) must equal B.size(0)");

    auto A_c = A.contiguous();
    auto B_c = B.contiguous();
    auto out = torch::empty_like(B_c);

    const long N = (long)A_c.size(0);
    const long M = (long)B_c.size(1);
    const long total = N * M;
    if (total == 0) {
        return out;
    }

    const float* a_ptr = A_c.data_ptr<float>();
    const float* b_ptr = B_c.data_ptr<float>();
    float* o_ptr = out.data_ptr<float>();

    const int block_size = 256;
    cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    const bool use_int =
        total < (long)std::numeric_limits<int>::max() - (long)block_size;

    if ((M & 3L) == 0) {
        const long total4 = total >> 2;
        const long num_blocks = (total4 + block_size - 1) / block_size;
        if (use_int) {
            diag_matmul_vec4_kernel<int><<<num_blocks, block_size, 0, stream>>>(
                a_ptr, b_ptr, o_ptr, (int)M, (int)total4);
        } else {
            diag_matmul_vec4_kernel<long><<<num_blocks, block_size, 0, stream>>>(
                a_ptr, b_ptr, o_ptr, M, total4);
        }
    } else {
        const long num_blocks = (total + block_size - 1) / block_size;
        if (use_int) {
            diag_matmul_scalar_kernel<int><<<num_blocks, block_size, 0, stream>>>(
                a_ptr, b_ptr, o_ptr, (int)M, (int)total);
        } else {
            diag_matmul_scalar_kernel<long><<<num_blocks, block_size, 0, stream>>>(
                a_ptr, b_ptr, o_ptr, M, total);
        }
    }
    return out;
}
"""

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

diag_matmul_module = load_inline(
    name="diag_matmul_ext",
    cpp_sources=diag_matmul_cpp_source,
    cuda_sources=diag_matmul_cuda_source,
    functions=["diag_matmul_cuda"],
    verbose=False,
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3", "--use_fast_math"],
)


class ModelNew(nn.Module):
    """
    Optimized version of Model.

    torch.diag(A) @ B is mathematically equivalent to scaling row i of B by
    A[i], so the diag construction and full matmul are replaced by a single
    fused, memory-bound CUDA kernel (with a vectorized float4 path when the
    column count is a multiple of 4).
    """
    def __init__(self):
        super(ModelNew, self).__init__()
        self.diag_matmul = diag_matmul_module

    def forward(self, A, B):
        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 with identical math (diag(A) @ B == A[:, None] * B),
        # works for any device / dtype.
        return A.unsqueeze(1) * B
```