```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 is equivalent to row-wise scaling
# C[i, j] = A[i] * B[i, j], which avoids materializing the N x N diagonal
# matrix and reduces the work from O(N*N*M) to O(N*M) (pure memory-bound).
diag_matmul_source = """
#include <torch/extension.h>
#include <cuda_runtime.h>
#include <c10/cuda/CUDAStream.h>

__global__ void diag_matmul_vec4_kernel(const float* __restrict__ A,
                                        const float* __restrict__ B,
                                        float* __restrict__ C,
                                        int N, int M4) {
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (col < M4) {
        for (int row = blockIdx.y; row < N; row += gridDim.y) {
            int idx = row * M4 + col;
            float a = A[row];
            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;
        }
    }
}

__global__ void diag_matmul_kernel(const float* __restrict__ A,
                                   const float* __restrict__ B,
                                   float* __restrict__ C,
                                   int N, int M) {
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (col < M) {
        for (int row = blockIdx.y; row < N; row += gridDim.y) {
            int idx = row * M + col;
            C[idx] = A[row] * B[idx];
        }
    }
}

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() == at::kFloat && B.scalar_type() == at::kFloat,
                "Inputs must be float32");
    TORCH_CHECK(A.is_contiguous() && B.is_contiguous(), "Inputs must be contiguous");
    TORCH_CHECK(B.dim() == 2 && A.numel() == B.size(0),
                "A must have shape (N,) and B must have shape (N, M)");

    const int N = B.size(0);
    const int M = B.size(1);
    auto C = torch::empty_like(B);

    cudaStream_t stream = at::cuda::getCurrentCUDAStream();
    const int block_size = 256;
    const int max_grid_y = 65535;
    dim3 grid_y(std::min(N, max_grid_y));

    if (M % 4 == 0) {
        const int M4 = M / 4;
        dim3 grid((M4 + block_size - 1) / block_size, grid_y.y);
        diag_matmul_vec4_kernel<<<grid, block_size, 0, stream>>>(
            A.data_ptr<float>(), B.data_ptr<float>(), C.data_ptr<float>(), N, M4);
    } else {
        dim3 grid((M + block_size - 1) / block_size, grid_y.y);
        diag_matmul_kernel<<<grid, block_size, 0, stream>>>(
            A.data_ptr<float>(), B.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);"
)

# Compile the inline CUDA code
diag_matmul = load_inline(
    name="diag_matmul",
    cpp_sources=diag_matmul_cpp_source,
    cuda_sources=diag_matmul_source,
    functions=["diag_matmul_cuda"],
    verbose=False,
    extra_cflags=["-O3"],
    extra_ldflags=[""],
)


class ModelNew(nn.Module):
    """
    Optimized model that computes C = diag(A) * B with a custom CUDA kernel
    performing fused row-wise scaling: C[i, j] = A[i] * B[i, j].
    """
    def __init__(self):
        super(ModelNew, self).__init__()
        self.diag_matmul = diag_matmul

    def forward(self, A, B):
        if not A.is_cuda:
            A = A.cuda()
        if not B.is_cuda:
            B = B.cuda()
        A = A.contiguous()
        B = B.contiguous()
        return self.diag_matmul.diag_matmul_cuda(A, B)
```