```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 for C = diag(A) @ B.
# Mathematically this is row-wise scaling: C[i, j] = A[i] * B[i, j].
# The kernel is purely memory-bandwidth bound: it reads B once (64 MB),
# writes C once (64 MB) and reads A (16 KB, cached). The implementation is
# optimized to reach peak DRAM bandwidth:
#   * 2D grid: blockIdx.x selects the row, so no integer division is needed
#     to recover the row index (the previous version used a costly 64-bit
#     division per thread).
#   * 128-bit vectorized loads/stores (float4) with 4x unrolling so each
#     thread issues 4 independent 16-byte transactions (64 B in flight),
#     which saturates memory bandwidth with high memory-level parallelism.
#   * Streaming cache hints (__ldg / __stcs) since B and C are touched once.
#   * A templated "even" path removes bounds checks when the row length is
#     an exact multiple of the block tile.
diag_matmul_source = """
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <c10/cuda/CUDAException.h>
#include <cuda_runtime.h>

#define BLOCK_SIZE 256
#define UNROLL 4

// Vectorized kernel: requires M % 4 == 0. One block-row per matrix row.
template <bool EVEN>
__global__ void diag_matmul_vec4_kernel(const float* __restrict__ A,
                                        const float* __restrict__ B,
                                        float* __restrict__ C,
                                        int vecM) {
    const long row = blockIdx.x;
    const float a = __ldg(A + row);
    const long offset = row * (long)vecM;
    const float4* Brow = reinterpret_cast<const float4*>(B) + offset;
    float4* Crow = reinterpret_cast<float4*>(C) + offset;

    const int base = blockIdx.y * (BLOCK_SIZE * UNROLL) + threadIdx.x;
    #pragma unroll
    for (int k = 0; k < UNROLL; ++k) {
        const int i = base + k * BLOCK_SIZE;
        if (EVEN || i < vecM) {
            const float4 b = __ldg(Brow + i);
            float4 c;
            c.x = a * b.x;
            c.y = a * b.y;
            c.z = a * b.z;
            c.w = a * b.w;
            __stcs(Crow + i, c);
        }
    }
}

// Scalar fallback kernel (handles any M).
__global__ void diag_matmul_scalar_kernel(const float* __restrict__ A,
                                          const float* __restrict__ B,
                                          float* __restrict__ C,
                                          long M) {
    const long row = blockIdx.x;
    const float a = __ldg(A + row);
    const long offset = row * M;
    const float* Brow = B + offset;
    float* Crow = C + offset;
    const long i = (long)blockIdx.y * blockDim.x + threadIdx.x;
    if (i < M) {
        __stcs(Crow + i, __ldg(Brow + i) * a);
    }
}

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(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);

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

    cudaStream_t stream = at::cuda::getCurrentCUDAStream();
    const dim3 block(BLOCK_SIZE);

    if ((M & 3) == 0) {
        const long vecM = M >> 2;
        const long tiles = (vecM + BLOCK_SIZE * UNROLL - 1) / (BLOCK_SIZE * UNROLL);
        const dim3 grid((unsigned)N, (unsigned)tiles);
        if (tiles * (BLOCK_SIZE * UNROLL) == vecM) {
            diag_matmul_vec4_kernel<true><<<grid, block, 0, stream>>>(
                A_c.data_ptr<float>(), B_c.data_ptr<float>(), C.data_ptr<float>(), (int)vecM);
        } else {
            diag_matmul_vec4_kernel<false><<<grid, block, 0, stream>>>(
                A_c.data_ptr<float>(), B_c.data_ptr<float>(), C.data_ptr<float>(), (int)vecM);
        }
    } else {
        const dim3 grid((unsigned)N, (unsigned)((M + BLOCK_SIZE - 1) / BLOCK_SIZE));
        diag_matmul_scalar_kernel<<<grid, block, 0, stream>>>(
            A_c.data_ptr<float>(), B_c.data_ptr<float>(), C.data_ptr<float>(), M);
    }
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    return C;
}
"""

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

# Compile the inline CUDA extension.
diag_matmul = load_inline(
    name="diag_matmul_v2",
    cpp_sources=diag_matmul_cpp_source,
    cuda_sources=diag_matmul_source,
    functions=["diag_matmul_cuda"],
    verbose=False,
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Optimized version of Model that computes C = diag(A) @ B as a fused
    row-scaling CUDA kernel (C[i, j] = A[i] * B[i, j]) in FP32.
    """
    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)
        # Fallback (e.g. CPU tensors): mathematically equivalent row scaling.
        return A.unsqueeze(-1) * B
```