```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

# C = diag(A) @ B is equivalent to scaling row i of B by A[i]:
#     C[i, j] = A[i] * B[i, j]
# This avoids materializing the NxN diagonal matrix and reduces the work
# from O(N*N*M) to O(N*M), making the op purely memory-bound.
diag_matmul_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>

// Scalar kernel: one row per blockIdx.y, grid-stride loops for generality.
__global__ void diag_matmul_kernel(const float* __restrict__ A,
                                   const float* __restrict__ B,
                                   float* __restrict__ C,
                                   long N, long M) {
    const long stride_col = (long)gridDim.x * blockDim.x;
    for (long row = blockIdx.y; row < N; row += gridDim.y) {
        const float a = A[row];
        const long base = row * M;
        for (long col = (long)blockIdx.x * blockDim.x + threadIdx.x; col < M; col += stride_col) {
            C[base + col] = a * B[base + col];
        }
    }
}

// Vectorized (float4) kernel used when M is divisible by 4.
__global__ void diag_matmul_vec4_kernel(const float* __restrict__ A,
                                        const float4* __restrict__ B,
                                        float4* __restrict__ C,
                                        long N, long M4) {
    const long stride_col = (long)gridDim.x * blockDim.x;
    for (long row = blockIdx.y; row < N; row += gridDim.y) {
        const float a = A[row];
        const long base = row * M4;
        for (long col = (long)blockIdx.x * blockDim.x + threadIdx.x; col < M4; col += stride_col) {
            const float4 b = B[base + col];
            float4 c;
            c.x = a * b.x;
            c.y = a * b.y;
            c.z = a * b.z;
            c.w = a * b.w;
            C[base + col] = 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() == torch::kFloat32 && B.scalar_type() == torch::kFloat32,
                "A and B must be float32 tensors");
    TORCH_CHECK(B.dim() == 2, "B must be a 2D tensor");

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

    const long N = A_c.numel();
    const long M = B_c.size(1);
    TORCH_CHECK(B_c.size(0) == N, "Size mismatch: A has ", N, " elements but B has ", B_c.size(0), " rows");

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

    const int block_size = 256;
    const long max_grid = 65535;
    cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    if ((M & 3) == 0) {
        const long M4 = M >> 2;
        long blocks_x = (M4 + block_size - 1) / block_size;
        if (blocks_x > max_grid) blocks_x = max_grid;
        const long blocks_y = N < max_grid ? N : max_grid;
        dim3 grid((unsigned int)blocks_x, (unsigned int)blocks_y);
        diag_matmul_vec4_kernel<<<grid, block_size, 0, stream>>>(
            A_c.data_ptr<float>(),
            reinterpret_cast<const float4*>(B_c.data_ptr<float>()),
            reinterpret_cast<float4*>(C.data_ptr<float>()),
            N, M4);
    } else {
        long blocks_x = (M + block_size - 1) / block_size;
        if (blocks_x > max_grid) blocks_x = max_grid;
        const long blocks_y = N < max_grid ? N : max_grid;
        dim3 grid((unsigned int)blocks_x, (unsigned int)blocks_y);
        diag_matmul_kernel<<<grid, 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);"

# Compile the inline CUDA extension
diag_matmul_ext = 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 computes C = diag(A) @ B with a custom CUDA kernel
    performing row-wise scaling C[i, j] = A[i] * B[i, j].
    """
    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 and B.dim() == 2):
            return diag_matmul_ext.diag_matmul_cuda(A, B)
        # Fallback (CPU tensors or other dtypes): mathematically equivalent.
        return A.reshape(-1, 1) * B
```