```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 whole operator is one memory-bandwidth-bound streaming pass:
# read B once, write C once (A is tiny and broadcast).
#
# Tuning decisions:
#   * 2D grid (column-tile, row): blockIdx.y selects the row, so no per-thread
#     integer division is required, and consecutively scheduled blocks touch
#     consecutively increasing addresses -> perfect DRAM streaming locality.
#   * 128-bit vectorized loads/stores (float4), 4x unrolled per thread so each
#     thread has 4 independent 16-byte loads in flight (64 B of memory-level
#     parallelism per thread), enough to saturate DRAM at full occupancy.
#   * Streaming cache hints (__ldcs / __stcs): B and C are each touched exactly
#     once, so evict-first policy keeps L2 behavior clean.
#   * Templated EVEN path removes all bounds checks when the row length is an
#     exact multiple of the block tile (true for 4096 x 4096).
#   * __launch_bounds__ guarantees 100% occupancy (8 blocks x 256 threads/SM).
diag_matmul_source = r"""
#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
#define MAX_GRID_Y 65535
#define MAX_FLAT_BLOCKS 4194304L

// Fast path (requires M % 4 == 0, N <= 65535, vecM fits in int).
// grid = (tiles_per_row, N). blockIdx.y = row, blockIdx.x = column tile.
template <bool EVEN>
__global__ void __launch_bounds__(BLOCK_SIZE, 8)
diag_matmul_vec4_kernel(const float* __restrict__ A,
                        const float* __restrict__ B,
                        float* __restrict__ C,
                        const int vecM) {
    const float a = __ldg(A + blockIdx.y);
    const long offset = (long)blockIdx.y * vecM;
    const float4* Brow = reinterpret_cast<const float4*>(B) + offset;
    float4* Crow = reinterpret_cast<float4*>(C) + offset;

    const int base = blockIdx.x * (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 = __ldcs(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);
        }
    }
}

// Flat grid-stride vec4 fallback (M % 4 == 0, any N / vecM size).
__global__ void __launch_bounds__(BLOCK_SIZE)
diag_matmul_flat_vec4_kernel(const float* __restrict__ A,
                             const float* __restrict__ B,
                             float* __restrict__ C,
                             const long totalV,
                             const long vecM) {
    const long stride = (long)gridDim.x * BLOCK_SIZE;
    const float4* Bv = reinterpret_cast<const float4*>(B);
    float4* Cv = reinterpret_cast<float4*>(C);
    for (long idx = (long)blockIdx.x * BLOCK_SIZE + threadIdx.x;
         idx < totalV; idx += stride) {
        const float a = __ldg(A + idx / vecM);
        const float4 b = __ldcs(Bv + idx);
        float4 c;
        c.x = a * b.x;
        c.y = a * b.y;
        c.z = a * b.z;
        c.w = a * b.w;
        __stcs(Cv + idx, c);
    }
}

// Flat grid-stride scalar fallback (handles any M).
__global__ void __launch_bounds__(BLOCK_SIZE)
diag_matmul_flat_scalar_kernel(const float* __restrict__ A,
                               const float* __restrict__ B,
                               float* __restrict__ C,
                               const long total,
                               const long M) {
    const long stride = (long)gridDim.x * BLOCK_SIZE;
    for (long idx = (long)blockIdx.x * BLOCK_SIZE + threadIdx.x;
         idx < total; idx += stride) {
        C[idx] = __ldcs(B + idx) * __ldg(A + idx / M);
    }
}

torch::Tensor diag_matmul_cuda(const torch::Tensor& A, const 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.dim() == 1 && B.dim() == 2, "A must be 1D, B must be 2D");
    TORCH_CHECK(A.numel() == B.size(0),
                "A must have exactly B.size(0) elements");

    auto A_c = A.is_contiguous() ? A : A.contiguous();
    auto B_c = B.is_contiguous() ? B : B.contiguous();

    const long N = B_c.size(0);
    const long M = B_c.size(1);

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

    const float* a_ptr = A_c.data_ptr<float>();
    const float* b_ptr = B_c.data_ptr<float>();
    float* c_ptr = C.data_ptr<float>();

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

    const long vecM = M >> 2;
    if ((M & 3) == 0) {
        if (N <= MAX_GRID_Y && vecM <= 0x7fffffffL) {
            const long tiles =
                (vecM + BLOCK_SIZE * UNROLL - 1) / (BLOCK_SIZE * UNROLL);
            const dim3 grid((unsigned)tiles, (unsigned)N);
            if (tiles * (BLOCK_SIZE * UNROLL) == vecM) {
                diag_matmul_vec4_kernel<true><<<grid, block, 0, stream>>>(
                    a_ptr, b_ptr, c_ptr, (int)vecM);
            } else {
                diag_matmul_vec4_kernel<false><<<grid, block, 0, stream>>>(
                    a_ptr, b_ptr, c_ptr, (int)vecM);
            }
        } else {
            const long totalV = N * vecM;
            long blocks = (totalV + BLOCK_SIZE - 1) / BLOCK_SIZE;
            if (blocks > MAX_FLAT_BLOCKS) blocks = MAX_FLAT_BLOCKS;
            diag_matmul_flat_vec4_kernel<<<(unsigned)blocks, block, 0, stream>>>(
                a_ptr, b_ptr, c_ptr, totalV, vecM);
        }
    } else {
        const long total = N * M;
        long blocks = (total + BLOCK_SIZE - 1) / BLOCK_SIZE;
        if (blocks > MAX_FLAT_BLOCKS) blocks = MAX_FLAT_BLOCKS;
        diag_matmul_flat_scalar_kernel<<<(unsigned)blocks, block, 0, stream>>>(
            a_ptr, b_ptr, c_ptr, total, M);
    }
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    return C;
}
"""

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

# Compile the inline CUDA extension.
diag_matmul = load_inline(
    name="diag_matmul_v4",
    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._fn = diag_matmul.diag_matmul_cuda

    def forward(self, A, B):
        try:
            # Zero-overhead fast path: all validation lives in the C++ op.
            return self._fn(A, B)
        except RuntimeError:
            # Fallbacks (CPU tensors, other dtypes, or unexpected ranks),
            # mathematically equivalent to torch.diag(A) @ B.
            if A.dim() <= 1:
                return A.unsqueeze(-1) * B
            return torch.diag(A) @ B
```