```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.
# Key insight: diag(A) @ B is just row-wise scaling of B by A:
#   out[i, j] = A[i] * B[i, j]
# This avoids materializing the NxN diagonal matrix and the O(N^2 * M) matmul,
# reducing the operation to a purely memory-bound elementwise scaling.
diag_matmul_source = """
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>

// Scalar fallback kernel: one thread per output element.
__global__ void diag_matmul_scalar_kernel(const float* __restrict__ A,
                                          const float* __restrict__ B,
                                          float* __restrict__ out,
                                          const int N, const int M) {
    const int idx = blockIdx.x * blockDim.x + threadIdx.x;
    const int total = N * M;
    if (idx < total) {
        const int row = idx / M;
        out[idx] = A[row] * B[idx];
    }
}

// Vectorized kernel: each thread processes 4 consecutive floats (float4).
// Requires M % 4 == 0 so every row is 16-byte aligned.
__global__ void diag_matmul_vec4_kernel(const float* __restrict__ A,
                                        const float* __restrict__ B,
                                        float* __restrict__ out,
                                        const int N, const int M4) {
    const int idx = blockIdx.x * blockDim.x + threadIdx.x;
    const int total4 = N * M4;
    if (idx < total4) {
        const int row = idx / M4;
        const float a = A[row];
        const 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(), "Inputs must be CUDA tensors");
    TORCH_CHECK(A.scalar_type() == torch::kFloat32 && B.scalar_type() == torch::kFloat32,
                "Inputs must be float32");
    TORCH_CHECK(A.dim() == 1 && B.dim() == 2, "A must be 1D and B must be 2D");
    TORCH_CHECK(A.size(0) == B.size(0), "Size of A must match rows of B");

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

    const int64_t N = A_c.size(0);
    const int64_t M = B_c.size(1);
    const int64_t total = N * M;
    TORCH_CHECK(total <= INT32_MAX, "Tensor too large for this kernel");

    auto out = torch::empty_like(B_c);
    if (total == 0) {
        return out;
    }

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

    if (M % 4 == 0) {
        const int M4 = static_cast<int>(M / 4);
        const int total4 = static_cast<int>(N * M4);
        const int num_blocks = (total4 + block_size - 1) / block_size;
        diag_matmul_vec4_kernel<<<num_blocks, block_size, 0, stream>>>(
            A_c.data_ptr<float>(), B_c.data_ptr<float>(), out.data_ptr<float>(),
            static_cast<int>(N), M4);
    } else {
        const int total_i = static_cast<int>(total);
        const int num_blocks = (total_i + block_size - 1) / block_size;
        diag_matmul_scalar_kernel<<<num_blocks, block_size, 0, stream>>>(
            A_c.data_ptr<float>(), B_c.data_ptr<float>(), out.data_ptr<float>(),
            static_cast<int>(N), static_cast<int>(M));
    }

    return out;
}
"""

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

# Compile the inline CUDA extension.
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 version of Model.

    Instead of building the full NxN diagonal matrix and performing a dense
    matmul (O(N^2 * M) FLOPs and O(N^2) memory), this uses a custom CUDA
    kernel that directly computes out[i, j] = A[i] * B[i, j], which is
    mathematically identical to diag(A) @ B. A float4-vectorized path is used
    when the row length permits, making the kernel saturate memory bandwidth.
    """
    def __init__(self):
        super(ModelNew, self).__init__()
        self.diag_matmul = diag_matmul

    def forward(self, A, B):
        if (A.is_cuda and B.is_cuda
                and A.dtype == torch.float32 and B.dtype == torch.float32
                and A.dim() == 1 and B.dim() == 2):
            return self.diag_matmul.diag_matmul_cuda(A, B)
        # Fallback for CPU tensors or non-float32 dtypes.
        return torch.diag(A) @ B
```