```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 fused CUDA kernel for C = diag(A) @ B.
#
# Key insight: (diag(A) @ B)[i, j] = A[i] * B[i, j], i.e. the operation is
# simply a per-row scaling of B. The original implementation materializes an
# (N, N) diagonal matrix (O(N^2) memory) and runs an O(N^2 * M) GEMM. This
# kernel computes the result directly in O(N * M) time with no extra memory,
# using vectorized 128-bit (float4) loads/stores when possible.
# ---------------------------------------------------------------------------

diag_matmul_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>
#include <algorithm>
#include <cstdint>

// Vectorized path (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__ C,
                                        int N, int Mv /* = M / 4 */) {
    for (int row = blockIdx.y; row < N; row += gridDim.y) {
        const float a = A[row];
        const float4* Brow = reinterpret_cast<const float4*>(B) + (int64_t)row * Mv;
        float4* Crow = reinterpret_cast<float4*>(C) + (int64_t)row * Mv;
        for (int i = blockIdx.x * blockDim.x + threadIdx.x;
             i < Mv;
             i += blockDim.x * gridDim.x) {
            const float4 b = Brow[i];
            float4 c;
            c.x = a * b.x;
            c.y = a * b.y;
            c.z = a * b.z;
            c.w = a * b.w;
            Crow[i] = c;
        }
    }
}

// Generic scalar path (any M).
__global__ void diag_matmul_scalar_kernel(const float* __restrict__ A,
                                          const float* __restrict__ B,
                                          float* __restrict__ C,
                                          int N, int M) {
    for (int row = blockIdx.y; row < N; row += gridDim.y) {
        const float a = A[row];
        const int64_t base = (int64_t)row * M;
        for (int col = blockIdx.x * blockDim.x + threadIdx.x;
             col < M;
             col += blockDim.x * gridDim.x) {
            C[base + col] = a * B[base + col];
        }
    }
}

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.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 must have one element per row of B");
    TORCH_CHECK(A.scalar_type() == at::kFloat && B.scalar_type() == at::kFloat,
                "Only float32 tensors are supported");
    TORCH_CHECK(A.size(0) <= 2147483647LL && B.size(1) <= 2147483647LL,
                "Tensor dimensions too large");

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

    const int N = (int)A_c.size(0);
    const int M = (int)B_c.size(1);
    if (N == 0 || M == 0) {
        return C;
    }

    const int threads = 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 int Mv = M >> 2;
        const int blocks_x = std::max(1, std::min((Mv + threads - 1) / threads, 1024));
        const int blocks_y = std::max(1, std::min(N, 65535));
        dim3 grid(blocks_x, blocks_y);
        diag_matmul_vec4_kernel<<<grid, threads, 0, stream>>>(A_ptr, B_ptr, C_ptr, N, Mv);
    } else {
        const int blocks_x = std::max(1, std::min((M + threads - 1) / threads, 1024));
        const int blocks_y = std::max(1, std::min(N, 65535));
        dim3 grid(blocks_x, blocks_y);
        diag_matmul_scalar_kernel<<<grid, threads, 0, stream>>>(A_ptr, B_ptr, C_ptr, N, M);
    }
    return C;
}
"""

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

# Compile the inline CUDA kernel.
diag_matmul = load_inline(
    name="diag_matmul_fused",
    cpp_sources=diag_matmul_cpp_source,
    cuda_sources=diag_matmul_cuda_source,
    functions=["diag_matmul_cuda"],
    verbose=False,
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Optimized version of Model that computes C = diag(A) @ B.

    Instead of building an (N, N) diagonal matrix and performing an
    O(N^2 * M) matrix multiplication, this module uses a fused custom CUDA
    kernel that directly evaluates C[i, j] = A[i] * B[i, j] in O(N * M),
    which is purely memory-bound and avoids materializing the diagonal
    matrix entirely.
    """

    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)
        # Generic fallback (mathematically identical to torch.diag(A) @ B).
        return A.unsqueeze(-1) * B
```