```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: torch.diag(A) @ B is just row-wise scaling of B by A,
# i.e. C[i, j] = A[i] * B[i, j]. This replaces an O(N*N*M) matmul
# (plus an O(N*N) diag materialization) with an O(N*M) elementwise kernel.
diag_matmul_source = """
#include <torch/extension.h>
#include <cuda_runtime.h>

// Vectorized (float4) kernel using a 2D grid: blockIdx.y = row, avoids integer div.
__global__ void diag_matmul_vec4_kernel(const float* __restrict__ A,
                                        const float4* __restrict__ B,
                                        float4* __restrict__ C,
                                        int M4) {
    int row = blockIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (col < M4) {
        float a = A[row];
        int idx = row * M4 + col;
        float4 b = B[idx];
        float4 c;
        c.x = a * b.x;
        c.y = a * b.y;
        c.z = a * b.z;
        c.w = a * b.w;
        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,
                                          int M) {
    int row = blockIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (col < M) {
        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() == torch::kFloat32 && B.scalar_type() == torch::kFloat32,
                "Inputs must be float32");
    TORCH_CHECK(A.dim() == 1 && B.dim() == 2, "A must be 1D and 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();
    auto C = torch::empty_like(B_c);

    const int N = B_c.size(0);
    const int M = B_c.size(1);
    if (N == 0 || M == 0) {
        return C;
    }

    const int block_size = 256;
    const int max_grid_y = 65535;

    if (N <= max_grid_y) {
        if (M % 4 == 0) {
            const int M4 = M / 4;
            dim3 grid((M4 + block_size - 1) / block_size, N);
            diag_matmul_vec4_kernel<<<grid, block_size>>>(
                A_c.data_ptr<float>(),
                reinterpret_cast<const float4*>(B_c.data_ptr<float>()),
                reinterpret_cast<float4*>(C.data_ptr<float>()),
                M4);
        } else {
            dim3 grid((M + block_size - 1) / block_size, N);
            diag_matmul_scalar_kernel<<<grid, block_size>>>(
                A_c.data_ptr<float>(),
                B_c.data_ptr<float>(),
                C.data_ptr<float>(),
                M);
        }
    } else {
        // Fallback for extremely large N: flat 1D grid with integer division.
        const long long total = (long long)N * (long long)M;
        const long long num_blocks = (total + block_size - 1) / block_size;
        TORCH_CHECK(num_blocks <= 2147483647LL, "Tensor too large");
        // Reuse scalar kernel logic via a flat index is not needed for the target
        // sizes; fall back to batched row processing.
        for (int row0 = 0; row0 < N; row0 += max_grid_y) {
            int rows = std::min(N - row0, max_grid_y);
            dim3 grid((M + block_size - 1) / block_size, rows);
            diag_matmul_scalar_kernel<<<grid, block_size>>>(
                A_c.data_ptr<float>() + row0,
                B_c.data_ptr<float>() + (long long)row0 * M,
                C.data_ptr<float>() + (long long)row0 * M,
                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=True,
    extra_cflags=[""],
    extra_ldflags=[""],
)


class ModelNew(nn.Module):
    """
    Optimized model that performs diag(A) @ B as row-wise scaling with a custom CUDA kernel.
    C = diag(A) * B  <=>  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):
        return self.diag_matmul.diag_matmul_cuda(A, B)
```