import torch
import torch.nn as nn


class Model(nn.Module):
    """
    Row-wise softmax over edge scores stored in CSR format.

    This is a core primitive in graph attention networks and sparse attention on
    irregular neighborhoods. Each node owns a contiguous edge segment in
    `edge_scores[row_ptr[i]:row_ptr[i + 1]]`, and the kernel normalizes scores
    independently within each segment.
    """

    def __init__(self):
        super().__init__()

    def forward(
        self,
        row_ptr: torch.Tensor,
        edge_scores: torch.Tensor,
    ) -> torch.Tensor:
        """
        Args:
            row_ptr:      (num_nodes + 1,) CSR row offsets
            edge_scores:  (num_edges,) unnormalized attention logits
        Returns:
            (num_edges,) softmax-normalized edge weights per source node
        """
        num_nodes = row_ptr.numel() - 1
        out = torch.empty_like(edge_scores)

        for node in range(num_nodes):
            start = int(row_ptr[node].item())
            end = int(row_ptr[node + 1].item())
            if end > start:
                out[start:end] = torch.softmax(edge_scores[start:end], dim=0)

        return out


num_nodes = 512
avg_degree = 24
batchless_num_edges = num_nodes * avg_degree


def _make_row_ptr(num_nodes: int, avg_degree: int) -> torch.Tensor:
    base = torch.full((num_nodes,), avg_degree, dtype=torch.int32)
    return row_ptr


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



def get_init_inputs():
    return []

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

# Define the CUDA kernel for row-wise softmax on CSR format
cuda_source = """
#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <cmath>
#include <limits>

// Each thread processes one node's softmax computation
__global__ void csr_row_softmax_kernel(
    const int32_t* __restrict__ row_ptr,
    const float* __restrict__ edge_scores,
    float* __restrict__ out,
    int num_nodes
) {
    int node = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (node >= num_nodes) return;
    
    int32_t start = row_ptr[node];
    int32_t end = row_ptr[node + 1];
    int32_t length = end - start;
    
    // Handle empty rows
    if (length <= 0) return;
    
    const float* scores = edge_scores + start;
    float* out_segment = out + start;
    
    // Find max value for numerical stability
    float max_val = -std::numeric_limits<float>::infinity();
    for (int i = 0; i < length; i++) {
        float val = scores[i];
        if (val > max_val) max_val = val;
    }
    
    // Compute exp(x - max) and sum
    float sum = 0.0f;
    for (int i = 0; i < length; i++) {
        float exp_val = expf(scores[i] - max_val);
        out_segment[i] = exp_val;
        sum += exp_val;
    }
    
    // Normalize by sum
    float inv_sum = 1.0f / sum;
    for (int i = 0; i < length; i++) {
        out_segment[i] *= inv_sum;
    }
}

// Host function to launch the kernel
torch::Tensor csr_row_softmax_cuda(
    torch::Tensor row_ptr_contiguous,
    torch::Tensor edge_scores_contiguous
) {
    int num_nodes = row_ptr_contiguous.numel() - 1;
    auto out = torch::empty_like(edge_scores_contiguous);
    
    // Launch configuration
    int threads_per_block = 256;
    int num_blocks = (num_nodes + threads_per_block - 1) / threads_per_block;
    
    csr_row_softmax_kernel<<<num_blocks, threads_per_block>>>(
        row_ptr_contiguous.data_ptr<int32_t>(),
        edge_scores_contiguous.data_ptr<float>(),
        out.data_ptr<float>(),
        num_nodes
    );
    
    return out;
}
"""

cpp_source = """
torch::Tensor csr_row_softmax_cuda(torch::Tensor row_ptr_contiguous, torch::Tensor edge_scores_contiguous);
"""

# Load the CUDA extension
csr_softmax_ext = load_inline(
    name='csr_softmax_ext_v2',
    cpp_sources=cpp_source,
    cuda_sources=cuda_source,
    functions=['csr_row_softmax_cuda'],
    verbose=False,
    extra_cuda_cflags=['-O3']
)


class ModelNew(nn.Module):
    """
    Row-wise softmax over edge scores stored in CSR format.
    Uses a custom CUDA kernel for parallel computation.
    """

    def __init__(self):
        super().__init__()

    def forward(
        self,
        row_ptr: torch.Tensor,
        edge_scores: torch.Tensor,
    ) -> torch.Tensor:
        """
        Args:
            row_ptr:      (num_nodes + 1,) CSR row offsets
            edge_scores:  (num_edges,) unnormalized attention logits
        Returns:
            (num_edges,) softmax-normalized edge weights per source node
        """
        # Ensure contiguous tensors with correct types
        row_ptr_cont = row_ptr.contiguous().to(torch.int32)
        edge_scores_cont = edge_scores.contiguous().to(torch.float32)
        
        return csr_softmax_ext.csr_row_softmax_cuda(row_ptr_cont, edge_scores_cont)