import torch
import torch.nn as nn


class Model(nn.Module):
    """
    Degree-normalized neighbor aggregation.

    Computes a GCN-style sparse aggregation where each incoming message is
    scaled by 1 / sqrt(deg(dst) * deg(src)). This captures a common spectral
    graph normalization pattern over CSR adjacency.
    """

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

    def forward(
        self,
        row_ptr: torch.Tensor,
        col_idx: torch.Tensor,
        node_feat: torch.Tensor,
        degrees: torch.Tensor,
    ) -> torch.Tensor:
        num_nodes = row_ptr.numel() - 1
        feat_dim = node_feat.shape[1]
        out = torch.zeros(num_nodes, feat_dim, dtype=node_feat.dtype, device=node_feat.device)

        for dst in range(num_nodes):
            start = int(row_ptr[dst].item())
            end = int(row_ptr[dst + 1].item())
            if end <= start:
                continue
            src = col_idx[start:end].long()
            norm = 1.0 / torch.sqrt(degrees[dst] * degrees[src]).unsqueeze(-1)
            out[dst] = (node_feat[src] * norm).sum(dim=0)

        return out


num_nodes = 640
avg_degree = 16
feat_dim = 80




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



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>
#include <cmath>

__global__ void gcn_agg_kernel(
    const int* __restrict__ row_ptr,
    const int* __restrict__ col_idx,
    const float* __restrict__ node_feat,
    const float* __restrict__ degrees,
    float* __restrict__ out,
    int num_nodes,
    int feat_dim
) {
    int dst = blockIdx.x * blockDim.x + threadIdx.x;
    if (dst >= num_nodes) return;
    int start = row_ptr[dst];
    int end = row_ptr[dst + 1];
    float deg_dst = degrees[dst];
    for (int i = start; i < end; ++i) {
        int src = col_idx[i];
        float deg_src = degrees[src];
        float norm = 1.0f / sqrtf(deg_dst * deg_src);
        for (int f = 0; f < feat_dim; ++f) {
            out[dst * feat_dim + f] += node_feat[src * feat_dim + f] * norm;
        }
    }
}

torch::Tensor gcn_aggregate_cuda(
    torch::Tensor row_ptr,
    torch::Tensor col_idx,
    torch::Tensor node_feat,
    torch::Tensor degrees
) {
    row_ptr = row_ptr.contiguous();
    col_idx = col_idx.contiguous();
    node_feat = node_feat.contiguous();
    degrees = degrees.contiguous();
    int num_nodes = row_ptr.size(0) - 1;
    int feat_dim = node_feat.size(1);
    auto out = torch::zeros({num_nodes, feat_dim}, node_feat.options()).contiguous();
    int threads = 256;
    int blocks = (num_nodes + threads - 1) / threads;
    gcn_agg_kernel<<<blocks, threads>>>(
        row_ptr.data_ptr<int>(),
        col_idx.data_ptr<int>(),
        node_feat.data_ptr<float>(),
        degrees.data_ptr<float>(),
        out.data_ptr<float>(),
        num_nodes,
        feat_dim
    );
    cudaDeviceSynchronize();
    return out;
}
'''

cpp_source = '''
torch::Tensor gcn_aggregate_cuda(
    torch::Tensor row_ptr,
    torch::Tensor col_idx,
    torch::Tensor node_feat,
    torch::Tensor degrees
);
'''

module = load_inline(
    name='gcn_agg',
    cpp_sources=[cpp_source],
    cuda_sources=[cuda_source],
    functions=['gcn_aggregate_cuda'],
    verbose=False,
    extra_cuda_cflags=['-O3']
)

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

    def forward(self, row_ptr, col_idx, node_feat, degrees):
        return self.module.gcn_aggregate_cuda(row_ptr, col_idx, node_feat, degrees)