```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: C = diag(A) @ B is mathematically equivalent to a
# row-scaling: C[i, j] = A[i] * B[i, j]. This reduces an O(N^2 * M) matmul
# (plus an O(N^2) diag materialization) to a single memory-bound O(N*M) pass.
diag_matmul_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>
#include <cstdint>

// Vectorized kernel: each thread processes 4 consecutive floats (float4).
// Valid when M % 4 == 0 (rows stay 16-byte aligned).
__global__ void diag_matmul_vec4_kernel(const float* __restrict__ A,
                                        const float* __restrict__ B,
                                        float* __restrict__ C,
                                        int64_t total4,
                                        int64_t M,
                                        int M_shift) {
    int64_t i4 = (int64_t)blockIdx.x * blockDim.x + threadIdx.x;
    if (i4 < total4) {
        int64_t base = i4 << 2;
        int64_t row = (M_shift >= 0) ? (base >> M_shift) : (base / M);
        float a = __ldg(A + row);
        float4 b = *reinterpret_cast<const float4*>(B + base);
        float4 c;
        c.x = a * b.x;
        c.y = a * b.y;
        c.z = a * b.z;
        c.w = a * b.w;
        *reinterpret_cast<float4*>(C + base) = c;
    }
}

// Scalar fallback kernel for the general case.
__global__ void diag_matmul_kernel(const float* __restrict__ A,
                                   const float* __restrict__ B,
                                   float* __restrict__ C,
                                   int64_t total,
                                   int64_t M,
                                   int M_shift) {
    int64_t idx = (int64_t)blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < total) {
        int64_t row = (M_shift >= 0) ? (idx >> M_shift) : (idx / M);
        C[idx] = __ldg(A + row) * B[idx];
    }
}

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 tensors");
    TORCH_CHECK(A.dim() == 1 && B.dim() == 2, "A must be 1D, B must be 2D");
    TORCH_CHECK(A.size(0) == B.size(0), "size mismatch between A and B");

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

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

    // If M is a power of two, replace the integer division with a shift.
    int M_shift = -1;
    if ((M & (M - 1)) == 0) {
        M_shift = 0;
        while ((int64_t(1) << M_shift) < M) {
            ++M_shift;
        }
    }

    const int block_size = 256;
    const float* A_ptr = A_c.data_ptr<float>();
    const float* B_ptr = B_c.data_ptr<float>();
    float* C_ptr = C.data_ptr<float>();
    cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    if ((M & 3) == 0) {
        const int64_t total4 = total >> 2;
        const int64_t blocks = (total4 + block_size - 1) / block_size;
        diag_matmul_vec4_kernel<<<blocks, block_size, 0, stream>>>(
            A_ptr, B_ptr, C_ptr, total4, M, M_shift);
    } else {
        const int64_t blocks = (total + block_size - 1) / block_size;
        diag_matmul_kernel<<<blocks, block_size, 0, stream>>>(
            A_ptr, B_ptr, C_ptr, total, M, M_shift);
    }

    return C;
}
"""

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

# Compile the inline CUDA extension.
diag_matmul_ext = 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"],
)


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

    Instead of materializing torch.diag(A) (an N x N matrix) and running a
    full O(N^2 * M) matrix multiplication, we exploit the structure:
        (diag(A) @ B)[i, j] = A[i] * B[i, j]
    and perform a single fused, vectorized, memory-bound CUDA kernel.
    """
    def __init__(self):
        super(ModelNew, self).__init__()
        self.diag_matmul = diag_matmul_ext

    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)
        # Equivalent fallback (diag(A) @ B == row-wise scaling of B by A).
        return A.unsqueeze(1) * B
```