import torch
import torch.nn as nn


class Model(nn.Module):
    """
    Per-row top-k selection over CSR edge segments.

    Returns the top-k values and their original edge indices for each node's
    neighborhood. Useful for graph sparsification, hard attention, and
    candidate-pruning workloads.
    """

    def __init__(self, k):
        super().__init__()
        self.k = k

    def forward(
        self,
        row_ptr: torch.Tensor,
        edge_scores: torch.Tensor,
    ) -> tuple[torch.Tensor, torch.Tensor]:
        num_nodes = row_ptr.numel() - 1
        topk_vals = torch.full((num_nodes, self.k), float("-inf"), dtype=edge_scores.dtype, device=edge_scores.device)
        topk_idx = torch.full((num_nodes, self.k), -1, dtype=torch.int64, device=edge_scores.device)

        for node in range(num_nodes):
            start = int(row_ptr[node].item())
            end = int(row_ptr[node + 1].item())
            if end <= start:
                continue
            segment = edge_scores[start:end]
            curr_k = min(self.k, end - start)
            vals, idx = torch.topk(segment, k=curr_k, dim=0)
            topk_vals[node, :curr_k] = vals
            topk_idx[node, :curr_k] = idx.long() + start

        return topk_vals, topk_idx


num_nodes = 320
avg_degree = 12
k = 4




def get_inputs():
    return list(popcorn_pri.make_csr_scalar_edge_scores(num_nodes, avg_degree))



def get_init_inputs():
    return [k]

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

cuda_source = r'''
#include <cuda_runtime.h>
#include <cfloat>
#include <limits>
#include <tuple>
#include <cstdint>

__global__ void topk_csr_kernel(
    const int32_t* __restrict__ row_ptr,
    const float* __restrict__ edge_scores,
    float* __restrict__ topk_vals,
    int64_t* __restrict__ topk_idx,
    const int k,
    const int num_nodes
) {
    int node = blockIdx.x * blockDim.x + threadIdx.x;
    if (node >= num_nodes) return;

    int start = row_ptr[node];
    int end = row_ptr[node + 1];
    int degree = end - start;

    if (degree <= 0) {
        for (int i = 0; i < k; ++i) {
            topk_vals[node * k + i] = -FLT_MAX;
            topk_idx[node * k + i] = -1;
        }
        return;
    }

    int curr_k = min(k, degree);

    float vals[32];
    int orig_idx[32];
    for (int i = 0; i < degree; ++i) {
        vals[i] = edge_scores[start + i];
        orig_idx[i] = start + i;
    }

    for (int i = 0; i < curr_k; ++i) {
        int max_pos = i;
        for (int j = i + 1; j < degree; ++j) {
            if (vals[j] > vals[max_pos]) {
                max_pos = j;
            }
        }
        float tmpv = vals[i]; vals[i] = vals[max_pos]; vals[max_pos] = tmpv;
        int tmpi = orig_idx[i]; orig_idx[i] = orig_idx[max_pos]; orig_idx[max_pos] = tmpi;
    }

    for (int i = 0; i < curr_k; ++i) {
        topk_vals[node * k + i] = vals[i];
        topk_idx[node * k + i] = orig_idx[i];
    }
    for (int i = curr_k; i < k; ++i) {
        topk_vals[node * k + i] = -FLT_MAX;
        topk_idx[node * k + i] = -1;
    }
}

std::tuple<torch::Tensor, torch::Tensor> topk_csr_cuda(torch::Tensor row_ptr, torch::Tensor edge_scores, int k) {
    int num_nodes = row_ptr.size(0) - 1;
    auto options_vals = torch::TensorOptions().dtype(edge_scores.dtype()).device(edge_scores.device());
    auto options_idx = torch::TensorOptions().dtype(torch::kInt64).device(edge_scores.device());
    auto topk_vals = torch::full({num_nodes, k}, -std::numeric_limits<float>::infinity(), options_vals);
    auto topk_idx = torch::full({num_nodes, k}, -1, options_idx);

    int threads = 256;
    int blocks = (num_nodes + threads - 1) / threads;
    topk_csr_kernel<<<blocks, threads>>>(
        row_ptr.data_ptr<int32_t>(),
        edge_scores.data_ptr<float>(),
        topk_vals.data_ptr<float>(),
        topk_idx.data_ptr<int64_t>(),
        k,
        num_nodes
    );
    cudaDeviceSynchronize();
    return std::make_tuple(topk_vals, topk_idx);
}
'''

cpp_source = r'''
#include <torch/extension.h>
std::tuple<torch::Tensor, torch::Tensor> topk_csr_cuda(torch::Tensor row_ptr, torch::Tensor edge_scores, int k);
'''

module = load_inline(
    name='topk_csr',
    cpp_sources=[cpp_source],
    cuda_sources=[cuda_source],
    functions=['topk_csr_cuda'],
    verbose=True
)

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

    def forward(self, row_ptr: torch.Tensor, edge_scores: torch.Tensor):
        return self.module.topk_csr_cuda(row_ptr.contiguous(), edge_scores.contiguous(), self.k)