```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 optimization: diag(A) @ B is equivalent to scaling row i of B by A[i].
# This replaces an O(N^2 * M) matmul (plus materializing an N x N matrix)
# with an O(N * M) memory-bound elementwise kernel using vectorized float4 I/O.

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

// Vectorized kernel: each thread handles one float4 (4 columns) of a row.
__global__ void diag_matmul_vec4_kernel(const float* __restrict__ A,
                                        const float* __restrict__ B,
                                        float* __restrict__ C,
                                        const int N,
                                        const int M4) {
    const int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (col < M4) {
        for (int row = blockIdx.y; row < N; row += gridDim.y) {
            const float a = A[row];
            const long idx = (long)row * M4 + col;
            const float4 b = reinterpret_cast<const float4*>(B)[idx];
            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)[idx] = c;
        }
    }
}

// Scalar fallback kernel for M not divisible by 4.
__global__ void diag_matmul_scalar_kernel(const float* __restrict__ A,
                                          const float* __restrict__ B,
                                          float* __restrict__ C,
                                          const int M,
                                          const long total) {
    const long idx = (long)blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < total) {
        const int row = (int)(idx / M);
        C[idx] = A[row] * B[idx];
    }
}

torch::Tensor diag_matmul_cuda(torch::Tensor A, torch::Tensor B) {
    auto A_c = A.contiguous();
    auto B_c = B.contiguous();
    const int N = B_c.size(0);
    const int M = B_c.size(1);
    auto C = torch::empty_like(B_c);
    cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    if ((M & 3) == 0) {
        const int M4 = M >> 2;
        const int block_size = 256;
        const dim3 grid((M4 + block_size - 1) / block_size, N);
        diag_matmul_vec4_kernel<<<grid, block_size, 0, stream>>>(
            A_c.data_ptr<float>(), B_c.data_ptr<float>(),
            C.data_ptr<float>(), N, M4);
    } else {
        const long total = (long)N * M;
        const int block_size = 256;
        const long num_blocks = (total + 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>(),
            C.data_ptr<float>(), M, total);
    }
    return C;
}
"""

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

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


class ModelNew(nn.Module):
    """
    Optimized model that performs C = diag(A) * B via a custom CUDA kernel
    that scales each row of B by the corresponding element of A.
    """
    def __init__(self):
        super(ModelNew, self).__init__()

    def forward(self, A, B):
        if A.is_cuda and B.is_cuda and A.dtype == torch.float32 and B.dtype == torch.float32:
            return diag_matmul.diag_matmul_cuda(A, B)
        # Fallback (mathematically identical, avoids materializing diag(A)).
        return A.unsqueeze(1) * B
```