```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  <=>  C[i, j] = A[i] * B[i, j]
# This avoids materializing the N x N diagonal matrix and the O(N^2 * M) GEMM,
# reducing the whole operation to a single memory-bound element-wise pass.
diag_matmul_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>
#include <cstdint>

// Vectorized (float4) path, used when M % 4 == 0.
__global__ void diag_matmul_vec4_kernel(
    const float* __restrict__ A,
    const float* __restrict__ B,
    float* __restrict__ out,
    const int N,
    const int M) {
    const int M4 = M >> 2;
    const int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (col < M4) {
        for (int row = blockIdx.y; row < N; row += gridDim.y) {
            const int64_t base = (int64_t)row * M;
            const float a = A[row];
            const float4 b = *reinterpret_cast<const float4*>(B + base + (col << 2));
            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 + base + (col << 2)) = o;
        }
    }
}

// Scalar fallback path for arbitrary M.
__global__ void diag_matmul_kernel(
    const float* __restrict__ A,
    const float* __restrict__ B,
    float* __restrict__ out,
    const int N,
    const int M) {
    const int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (col < M) {
        for (int row = blockIdx.y; row < N; row += gridDim.y) {
            const int64_t idx = (int64_t)row * M + col;
            out[idx] = 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.size(0) == B.size(0),
                "A must have shape (N,) and B must have shape (N, M)");

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

    const int N = (int)A_c.size(0);
    const int M = (int)B_c.size(1);

    auto out = torch::empty_like(B_c);

    const int threads = 256;
    const int grid_y = N < 65535 ? N : 65535;
    cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    if ((M & 3) == 0) {
        const int M4 = M >> 2;
        dim3 grid((M4 + threads - 1) / threads, grid_y);
        diag_matmul_vec4_kernel<<<grid, threads, 0, stream>>>(
            A_c.data_ptr<float>(), B_c.data_ptr<float>(), out.data_ptr<float>(), N, M);
    } else {
        dim3 grid((M + threads - 1) / threads, grid_y);
        diag_matmul_kernel<<<grid, threads, 0, stream>>>(
            A_c.data_ptr<float>(), B_c.data_ptr<float>(), out.data_ptr<float>(), N, M);
    }

    return out;
}
"""

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

# Compile the inline CUDA code
diag_matmul = 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_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Optimized version of Model: replaces torch.diag(A) @ B with a single fused
    CUDA kernel computing C[i, j] = A[i] * B[i, j] (row-wise scaling of B),
    eliminating 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:
            return self.diag_matmul.diag_matmul_cuda(A, B)
        # Fallback for non-CUDA / non-float32 inputs (mathematically identical)
        return A.unsqueeze(1) * B
```