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

cuda_source = r'''
#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <cstdint>

__device__ __forceinline__ float warp_reduce_max(float val) {
    for (int offset = 16; offset > 0; offset /= 2) {
        val = fmaxf(val, __shfl_down_sync(0xffffffff, val, offset));
    }
    return val;
}

__device__ __forceinline__ float warp_reduce_sum(float val) {
    for (int offset = 16; offset > 0; offset /= 2) {
        val += __shfl_down_sync(0xffffffff, val, offset);
    }
    return val;
}

__global__ void csr_softmax_kernel(
    const int32_t* __restrict__ row_ptr,
    const float* __restrict__ edge_scores,
    float* __restrict__ out,
    int num_nodes
) {
    int node = blockIdx.x;
    if (node >= num_nodes) return;
    int start = row_ptr[node];
    int end = row_ptr[node + 1];
    int len = end - start;
    if (len <= 0) return;
    int tid = threadIdx.x;
    float local_max = -INFINITY;
    for (int i = tid; i < len; i += 32) {
        local_max = fmaxf(local_max, edge_scores[start + i]);
    }
    float row_max = warp_reduce_max(local_max);
    row_max = __shfl_sync(0xffffffff, row_max, 0);
    float local_sum = 0.0f;
    for (int i = tid; i < len; i += 32) {
        float val = expf(edge_scores[start + i] - row_max);
        out[start + i] = val;
        local_sum += val;
    }
    float row_sum = warp_reduce_sum(local_sum);
    row_sum = __shfl_sync(0xffffffff, row_sum, 0);
    for (int i = tid; i < len; i += 32) {
        out[start + i] /= row_sum;
    }
}

torch::Tensor csr_softmax_cuda(torch::Tensor row_ptr, torch::Tensor edge_scores) {
    TORCH_CHECK(row_ptr.dtype() == torch::kInt32, "row_ptr must be int32");
    TORCH_CHECK(edge_scores.dtype() == torch::kFloat32, "edge_scores must be float32");
    TORCH_CHECK(row_ptr.is_cuda(), "row_ptr must be CUDA");
    int num_nodes = row_ptr.size(0) - 1;
    auto out = torch::empty_like(edge_scores);
    dim3 grid(num_nodes);
    dim3 block(32);
    csr_softmax_kernel<<<grid, block>>>(
        row_ptr.data_ptr<int32_t>(),
        edge_scores.data_ptr<float>(),
        out.data_ptr<float>(),
        num_nodes
    );
    return out;
}
'''

cpp_source = '''
torch::Tensor csr_softmax_cuda(torch::Tensor row_ptr, torch::Tensor edge_scores);
'''

class ModelNew(nn.Module):
    def __init__(self):
        super().__init__()
        self.op = load_inline(
            name='csr_softmax',
            cpp_sources=cpp_source,
            cuda_sources=cuda_source,
            functions=['csr_softmax_cuda'],
            verbose=False
        )

    def forward(self, row_ptr, edge_scores):
        row_ptr = row_ptr.to(torch.int32).cuda().contiguous()
        edge_scores = edge_scores.to(torch.float32).cuda().contiguous()
        return self.op.csr_softmax_cuda(row_ptr, edge_scores)
```