"""
Over virtual ranks R, for each (B, C) position take the maximum value; on ties choose the smallest rank index.

Output (B, C, 2): [:,:,0] = best value, [:,:,1] = winning rank as float32.
"""

import torch
import torch.nn as nn

from kernelbench.distributed_collectives import default_device

R = 4
B = 8
C = 28


class Model(nn.Module):
    def __init__(self):
        super().__init__()

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        # x: (R, B, C)
        best_val = x[0].clone()
        best_idx = torch.zeros(B, C, dtype=torch.long, device=x.device)
        for r in range(1, R):
            xv = x[r]
            improved = xv > best_val
            tied = (xv == best_val) & ~improved
            best_idx = torch.where(
                improved,
                torch.full_like(best_idx, r),
                torch.where(
                    tied,
                    torch.minimum(best_idx, torch.full_like(best_idx, r)),
                    best_idx,
                ),
            )
            best_val = torch.where(improved, xv, best_val)
        stacked = torch.stack([best_val, best_idx.to(best_val.dtype)], dim=-1)
        return stacked


def get_inputs():
    dev = default_device()
    p = popcorn_pri
    return [torch.randn(R, p.jitter_int(B), p.jitter_int(C), device=dev)]


def get_init_inputs():
    return []

import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline

_cpp_src = r"""
#include <torch/extension.h>

void max_rank_forward_out(torch::Tensor x, torch::Tensor out);
"""

_cuda_src = r"""
#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>

#define N_ELEMS 224

__global__ void max_rank_kernel(const float* __restrict__ x, float* __restrict__ out) {
    int i = threadIdx.x;
    float best = x[i];
    int idx = 0;
    float v = x[N_ELEMS + i];
    if (v > best) { best = v; idx = 1; }
    v = x[2 * N_ELEMS + i];
    if (v > best) { best = v; idx = 2; }
    v = x[3 * N_ELEMS + i];
    if (v > best) { best = v; idx = 3; }
    int o = i << 1;
    out[o] = best;
    out[o + 1] = (float)idx;
}

void max_rank_forward_out(torch::Tensor x, torch::Tensor out) {
    max_rank_kernel<<<1, N_ELEMS>>>(x.data_ptr<float>(), out.data_ptr<float>());
}
"""

_ext = load_inline(
    name="max_rank_ext_kb_224",
    cpp_sources=[_cpp_src],
    cuda_sources=[_cuda_src],
    functions=["max_rank_forward_out"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    verbose=False,
)

class ModelNew(nn.Module):
    def __init__(self):
        super().__init__()
        self._out = None
    def forward(self, x: torch.Tensor) -> torch.Tensor:
        out = self._out
        if out is None:
            out = torch.empty((8, 28, 2), device=x.device, dtype=x.dtype)
            self._out = out
        _ext.max_rank_forward_out(x, out)
        return out