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

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

__global__ void max_rank_kernel(const float* x, float* out, int R, int B, int C) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int total = B * C;
    if (idx >= total) return;
    int b = idx / C;
    int c = idx % C;

    float best_val = x[0 * B * C + b * C + c];
    int best_idx = 0;
    for (int r = 1; r < R; ++r) {
        float xv = x[r * B * C + b * C + c];
        if (xv > best_val) {
            best_val = xv;
            best_idx = r;
        } else if (xv == best_val && r < best_idx) {
            best_idx = r;
        }
    }
    out[idx * 2 + 0] = best_val;
    out[idx * 2 + 1] = (float)best_idx;
}

torch::Tensor max_rank_cuda(torch::Tensor x) {
    int R = x.size(0);
    int B = x.size(1);
    int C = x.size(2);
    auto out = torch::empty({B, C, 2}, x.options());
    int total = B * C;
    int threads = 256;
    int blocks = (total + threads - 1) / threads;
    max_rank_kernel<<<blocks, threads>>>(x.data_ptr<float>(), out.data_ptr<float>(), R, B, C);
    return out;
}
"""

cpp_source = """
torch::Tensor max_rank_cuda(torch::Tensor x);
"""

module = load_inline(
    name="max_rank_module",
    cpp_sources=[cpp_source],
    cuda_sources=[cuda_source],
    functions=["max_rank_cuda"],
    verbose=False,
    with_cuda=True,
)

class ModelNew(nn.Module):
    def __init__(self):
        super().__init__()
        self.module = module

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        return self.module.max_rank_cuda(x)
```