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

# diag(A) @ B is equivalent to scaling row i of B by A[i]: out[i, j] = A[i] * B[i, j].
# This is a purely memory-bound elementwise row-scaling, so we implement it as a
# single fused CUDA kernel (vectorized with float4 when possible) instead of
# materializing the N x N diagonal matrix and doing a full GEMM.
cuda_source = """
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>
#include <cstdint>

// Vectorized kernel: each thread handles 4 consecutive floats within one row.
// Requires M % 4 == 0 and N <= 65535.
__global__ void diag_matmul_vec4_kernel(const float* __restrict__ A,
                                        const float* __restrict__ B,
                                        float* __restrict__ out,
                                        int M4) {
    int row = blockIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (col < M4) {
        float a = __ldg(&A[row]);
        int64_t base = (int64_t)row * M4;
        float4 b = reinterpret_cast<const float4*>(B)[base + col];
        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] = o;
    }
}

// Generic fallback scalar kernel (any N, M).
__global__ void diag_matmul_scalar_kernel(const float* __restrict__ A,
                                          const float* __restrict__ B,
                                          float* __restrict__ out,
                                          int64_t N, int64_t M) {
    int64_t idx = blockIdx.x * (int64_t)blockDim.x + threadIdx.x;
    int64_t total = N * M;
    if (idx < total) {
        int64_t row = idx / M;
        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(), "Inputs must be CUDA tensors");
    TORCH_CHECK(A.scalar_type() == torch::kFloat32 && B.scalar_type() == torch::kFloat32,
                "Only float32 tensors are supported");
    auto A_c = A.contiguous();
    auto B_c = B.contiguous();
    const int64_t N = A_c.numel();
    const int64_t M = B_c.size(1);
    TORCH_CHECK(B_c.size(0) == N, "A must have length equal to B's first dimension");

    auto out = torch::empty(B_c.sizes(), B_c.options());
    cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    const int block_size = 256;
    if ((M % 4 == 0) && (N <= 65535)) {
        int M4 = (int)(M / 4);
        dim3 block(block_size);
        dim3 grid((M4 + block_size - 1) / block_size, (unsigned int)N);
        diag_matmul_vec4_kernel<<<grid, block, 0, stream>>>(
            A_c.data_ptr<float>(), B_c.data_ptr<float>(), out.data_ptr<float>(), M4);
    } else {
        int64_t total = N * M;
        int64_t 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>(), out.data_ptr<float>(), N, M);
    }
    return out;
}
"""

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

diag_matmul = load_inline(
    name="diag_matmul",
    cpp_sources=cpp_source,
    cuda_sources=cuda_source,
    functions=["diag_matmul_cuda"],
    verbose=False,
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Optimized version of Model that computes C = diag(A) * B with a custom
    fused CUDA kernel performing row-wise scaling of B by A, avoiding the
    explicit construction of the N x N diagonal matrix and the 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)
        # CPU fallback
        return torch.diag(A) @ B
```