```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: diag(A) @ B is just row-wise scaling: C[i, j] = A[i] * B[i, j].
# The original op materializes an (N, N) diagonal matrix and runs a full GEMM
# (O(N^2 * M) FLOPs, O(N^2) memory). This custom kernel does the same math in
# O(N * M) with a single memory-bound pass, using float4 vectorized I/O.
diag_matmul_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>
#include <cstdint>

// Scalar fallback kernel (any M, any alignment), grid-stride loop.
__global__ void diag_matmul_scalar_kernel(const float* __restrict__ A,
                                          const float* __restrict__ B,
                                          float* __restrict__ C,
                                          int64_t N, int64_t M) {
    int64_t total = N * M;
    int64_t stride = (int64_t)gridDim.x * blockDim.x;
    for (int64_t idx = (int64_t)blockIdx.x * blockDim.x + threadIdx.x;
         idx < total; idx += stride) {
        int64_t row = idx / M;
        C[idx] = A[row] * B[idx];
    }
}

// Vectorized kernel (requires M % 4 == 0 and 16-byte aligned pointers).
__global__ void diag_matmul_vec4_kernel(const float* __restrict__ A,
                                        const float* __restrict__ B,
                                        float* __restrict__ C,
                                        int64_t N, int64_t vecM) {
    int64_t total = N * vecM;
    int64_t stride = (int64_t)gridDim.x * blockDim.x;
    const float4* B4 = reinterpret_cast<const float4*>(B);
    float4* C4 = reinterpret_cast<float4*>(C);
    for (int64_t idx = (int64_t)blockIdx.x * blockDim.x + threadIdx.x;
         idx < total; idx += stride) {
        int64_t row = idx / vecM;
        float a = A[row];
        float4 b = B4[idx];
        float4 c;
        c.x = a * b.x;
        c.y = a * b.y;
        c.z = a * b.z;
        c.w = a * b.w;
        C4[idx] = c;
    }
}

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 length must equal B rows");

    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);

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

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

    const bool vec_ok =
        (M % 4 == 0) &&
        (reinterpret_cast<uintptr_t>(B_c.data_ptr<float>()) % 16 == 0) &&
        (reinterpret_cast<uintptr_t>(C.data_ptr<float>()) % 16 == 0);

    if (vec_ok) {
        const int64_t vecM = M / 4;
        const int64_t total = N * vecM;
        int64_t num_blocks = (total + block_size - 1) / block_size;
        if (num_blocks > max_blocks) num_blocks = max_blocks;
        diag_matmul_vec4_kernel<<<(unsigned int)num_blocks, block_size, 0, stream>>>(
            A_c.data_ptr<float>(), B_c.data_ptr<float>(), C.data_ptr<float>(), N, vecM);
    } else {
        const int64_t total = N * M;
        int64_t num_blocks = (total + block_size - 1) / block_size;
        if (num_blocks > max_blocks) num_blocks = max_blocks;
        diag_matmul_scalar_kernel<<<(unsigned int)num_blocks, block_size, 0, stream>>>(
            A_c.data_ptr<float>(), B_c.data_ptr<float>(), C.data_ptr<float>(), N, M);
    }

    return C;
}
"""

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

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_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Optimized model computing C = diag(A) * B via a fused custom CUDA kernel
    that performs row scaling directly, avoiding the O(N^2) diagonal matrix
    materialization and the O(N^2 * M) GEMM.
    """
    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)
        # Mathematically identical fallback without materializing diag(A).
        return A.unsqueeze(1) * B
```