Repository: bfGraph/STGraph (language: python)

Issue:
🕸️ Implementing GPMA for Seastar
Seastar uses PCSR for representing and storing dynamic graphs. We will now try to represent dynamic graphs using a GPMA instead on the GPU for possibly faster operations on the graph.


Discussion / hints:

Fix patch:
diff --git a/.gitignore b/.gitignore
index ce6cc21..d3086a6 100644
--- a/.gitignore
+++ b/.gitignore
@@ -1,17 +1,27 @@
-exp/dataset/citeseer/
-exp/dataset/Coauthor_cs/
-exp/dataset/Coauthor_physics/
-exp/dataset/cora/
-exp/dataset/CoraFull/
-exp/dataset/pubmed/
-exp/dataset/reddit/
-exp/dataset/aifb/
-exp/dataset/mutag/
-python/egl.egg-info/
-python/egl/__pycache__/
-python/seastar.egg-info
-python/build/
-python/dist/
+benchmarking/dataset/citeseer/
+benchmarking/dataset/Coauthor_cs/
+benchmarking/dataset/Coauthor_physics/
+benchmarking/dataset/cora/
+benchmarking/dataset/CoraFull/
+benchmarking/dataset/pubmed/
+benchmarking/dataset/reddit/
+benchmarking/dataset/aifb/
+benchmarking/dataset/mutag/
+benchmarking/dataset/england_covid/
+benchmarking/dataset/bitcoin/
+*.egg-info/
+seastar/build/
+seastar/dist/
 docs/_build
 __pycache__/
 .DS_Store
+*.so
+build
+dist
+*.egg-info
+*.txt
+test_*.*
+*.csv
+*.prof
+.vscode/
+*.svg
\ No newline at end of file
diff --git a/benchmarking/dataset/doorah/doorah.py b/benchmarking/dataset/doorah/doorah.py
new file mode 100644
index 0000000..d2c66c7
--- /dev/null
+++ b/benchmarking/dataset/doorah/doorah.py
@@ -0,0 +1,141 @@
+# Toy Dynamic Dataset
+
+import numpy as np
+from rich.pretty import pprint
+from rich import inspect
+
+edges_list = []
+edges_list.append(
+    np.array([
+        [0, 1, 2, 4],
+        [1, 3, 4, 5]
+    ])
+)
+edges_list.append(
+    np.array([
+        [2, 1, 3, 2, 3],
+        [0, 3, 2, 4, 5]
+    ])
+)
+edges_list.append(
+    np.array([
+        [0, 1, 0, 3, 3],
+        [3, 3, 2, 4, 5]
+    ])
+)
+edges_list.append(
+    np.array([
+        [0, 2, 5],
+        [4, 3, 1]
+    ])
+)
+
+# ---------------------------------
+
+edges_weight_list = []
+edges_weight_list.append(
+    np.array(
+        [1, 2, 3, 4]
+    )
+)
+edges_weight_list.append(
+    np.array(
+        [1, 2, 3, 4, 5]
+    )
+)
+edges_weight_list.append(
+    np.array(
+        [1, 2, 3, 4, 5]
+    )
+)
+edges_weight_list.append(
+    np.array(
+        [1, 2, 3]
+    )
+)
+
+# ---------------------------------
+
+all_features_list = []
+all_features_list.append(
+    np.array(
+        [[4], [3], [5], [7], [3], [2]]
+    )
+)
+all_features_list.append(
+    np.array(
+        [[3], [2], [7], [4], [3], [2]]
+    )
+)
+all_features_list.append(
+    np.array(
+        [[1], [3], [1], [7], [2], [3]]
+    )
+)
+all_features_list.append(
+    np.array(
+        [[2], [3], [1], [5], [6], [7]]
+    )
+)
+
+# ---------------------------------
+
+all_targets_list = []
+all_targets_list.append(
+    np.array(
+        [[2], [1], [6], [2], [4], [3]]
+    )
+)
+all_targets_list.append(
+    np.array(
+        [[2], [6], [8], [10], [4], [4]]
+    )
+)
+all_targets_list.append(
+    np.array(
+        [[2], [4], [6], [8], [5], [8]]
+    )
+)
+all_targets_list.append(
+    np.array(
+        [[5], [8], [9], [6], [9], [10]]
+    )
+)
+
+# ---------------------------------
+
+# inspect(edges_list)
+# inspect(edges_weight_list)
+# inspect(all_features_list)
+# inspect(all_targets_list)
+
+def get_doorah_dataset():
+    return edges_list, edges_weight_list, all_features_list, all_targets_list
+
+def preprocess_graph_structure(edges):
+    # inspect(edges)
+    tmp_set = set()
+    for i in range(len(edges)):
+        tmp_set = set()
+        for j in range(len(edges[i][0])):
+            tmp_set.add(edges[i][0][j])
+            tmp_set.add(edges[i][1][j])
+    max_num_nodes = len(tmp_set)
+
+    edge_dict = {}
+    for i in range(len(edges)):
+        edge_set = set()
+        for j in range(len(edges[i][0])):
+            edge_set.add((edges[i][0][j],edges[i][1][j]))
+        edge_dict[str(i)] = edge_set
+    
+    edge_final_dict = {}
+    edge_final_dict["0"] = {"add": list(edge_dict["0"]),"delete": []}
+    for i in range(1,len(edges)):
+        edge_final_dict[str(i)] = {"add": list(edge_dict[str(i)].difference(edge_dict[str(i-1)])), "delete": list(edge_dict[str(i-1)].difference(edge_dict[str(i)]))}
+    
+    return edge_final_dict, max_num_nodes
+
+operations, max_num_nodes = preprocess_graph_structure(edges_list)
+# inspect(operations)
+# inspect(max_num_nodes)
\ No newline at end of file
diff --git a/exp/dataset/gen_dataset.py b/benchmarking/dataset/gen_dataset.py
similarity index 99%
rename from exp/dataset/gen_dataset.py
rename to benchmarking/dataset/gen_dataset.py
index 7bcfe75..47921db 100644
--- a/exp/dataset/gen_dataset.py
+++ b/benchmarking/dataset/gen_dataset.py
@@ -64,7 +64,9 @@ def extract_dataset():
     register_data_args(parser)
     args = parser.parse_args()
     dataset_name = ['cora', 'citeseer', 'pubmed', 'reddit', 'CoraFull', 'Coauthor_cs', 'Coauthor_physics', 'AmazonCoBuy_computers', 'AmazonCoBuy_photo']
-
+    
+    # TODO: Remove
+    dataset_name = ['cora'] 
     print("Now PATH IS ", os.getcwd())
     for name in dataset_name:
         '''
diff --git a/exp/dataset/readme.md b/benchmarking/dataset/readme.md
similarity index 100%
rename from exp/dataset/readme.md
rename to benchmarking/dataset/readme.md
diff --git a/exp/gat/dgl/model.py b/benchmarking/gat/dgl/model.py
similarity index 100%
rename from exp/gat/dgl/model.py
rename to benchmarking/gat/dgl/model.py
diff --git a/exp/gat/dgl/train.py b/benchmarking/gat/dgl/train.py
similarity index 100%
rename from exp/gat/dgl/train.py
rename to benchmarking/gat/dgl/train.py
diff --git a/exp/gat/seastar/egl_gat.py b/benchmarking/gat/seastar/egl_gat.py
similarity index 95%
rename from exp/gat/seastar/egl_gat.py
rename to benchmarking/gat/seastar/egl_gat.py
index abe2c2c..e573c7b 100644
--- a/exp/gat/seastar/egl_gat.py
+++ b/benchmarking/gat/seastar/egl_gat.py
@@ -5,10 +5,10 @@
 import torch
 import dgl
 from torch import nn
-from seastar import CtxManager
-from seastar.backend.pytorch_backend import run_egl
+from seastar import Seastar
+from seastar.backend.pytorch_backend import backend_cb
 
-from seastar import CtxManager
+from seastar import Seastar
 
 # pylint: enable=W0235
 class EglGATConv(nn.Module):
@@ -37,7 +37,7 @@ def __init__(self,
                     self._in_feats, num_heads * out_feats, bias=False)
         self.reset_parameters()
         self.activation = activation
-        self.cm = CtxManager(run_egl)
+        self.seastar = Seastar(backend_cb)
 
     def reset_parameters(self):
         """Reinitialize learnable parameters."""
@@ -56,7 +56,7 @@ def forward(self, graph, feat):
         er = (feat_dst * self.attn_r).sum(dim=-1).unsqueeze(-1)
 
         # Vertex-centric implementation.
-        @self.cm.zoomIn(nspace=[self, torch])
+        @self.seastar.compile(nspace=[self, torch])
         def nb_forward(v):
            embs = [nb.el + v.er for nb in v.innbs]
            coeff = [torch.exp(self.leaky_relu(emb - max(embs))) for emb in embs]
diff --git a/exp/gat/seastar/train.py b/benchmarking/gat/seastar/train.py
similarity index 100%
rename from exp/gat/seastar/train.py
rename to benchmarking/gat/seastar/train.py
diff --git a/exp/gat/seastar/utils.py b/benchmarking/gat/seastar/utils.py
similarity index 100%
rename from exp/gat/seastar/utils.py
rename to benchmarking/gat/seastar/utils.py
diff --git a/exp/gcn/seastar/gcn_spmv.py b/benchmarking/gcn/seastar/gcn_spmv.py
similarity index 92%
rename from exp/gcn/seastar/gcn_spmv.py
rename to benchmarking/gcn/seastar/gcn_spmv.py
index ed4dc48..f68fa70 100644
--- a/exp/gcn/seastar/gcn_spmv.py
+++ b/benchmarking/gcn/seastar/gcn_spmv.py
@@ -10,8 +10,8 @@
 import torch.nn as nn
 import dgl
 import dgl.function as fn
-from seastar import CtxManager
-from seastar.backend.pytorch_backend import run_egl
+from seastar.compiler import Seastar
+from seastar.compiler.backend.pytorch_backend import backend_cb
 
 class EglGCNLayer(nn.Module):
     def __init__(self,
@@ -35,7 +35,7 @@ def __init__(self,
         else:
             self.dropout = 0.
         self.reset_parameters()
-        self.cm = CtxManager(run_egl)
+        self.seastar = Seastar(backend_cb)
 
     def reset_parameters(self):
         stdv = 1. / math.sqrt(self.weight.size(1))
@@ -48,12 +48,14 @@ def forward(self, h):
             h = self.dropout(h)
 
         h = torch.mm(h, self.weight)
-        @self.cm.zoomIn(nspace=[self, torch])
+        
+        @self.seastar.compile(nspace=[self, torch])
         def nb_compute(v):
             h = sum([nb.h*nb.norm for nb in v.innbs])
             h = h * v.norm
             return h
         h = nb_compute(g=self.g, n_feats={'norm': self.norm, 'h' : h})
+        
         # bias
         if self.bias is not None:
             h = h + self.bias
diff --git a/exp/gcn/seastar/train.py b/benchmarking/gcn/seastar/train.py
similarity index 67%
rename from exp/gcn/seastar/train.py
rename to benchmarking/gcn/seastar/train.py
index b576155..95960a8 100644
--- a/exp/gcn/seastar/train.py
+++ b/benchmarking/gcn/seastar/train.py
@@ -4,9 +4,13 @@
 import torch
 import torch.nn as nn
 import torch.nn.functional as F
-import dgl
-from dgl import DGLGraph
-from dgl.data import register_data_args, load_data
+
+from rich import inspect
+from rich.pretty import pprint
+
+from seastar.graph.static.StaticGraph import StaticGraph
+from seastar.dataset.cora import CoraDataset
+
 import snoop
 
 from gcn_spmv import EglGCN
@@ -26,49 +30,33 @@ def evaluate(model, features, labels, mask):
         correct = torch.sum(indices == labels)
         return correct.item() * 1.0 / len(labels)
 
-def main(args):
-    # load and preprocess dataset
-    path = '../../dataset/' + str(args.dataset) + '/'
-    '''
-    edges = np.loadtxt(path + 'edges.txt')
-    edges = edges.astype(int)
-
-    features = np.loadtxt(path + 'features.txt')
-
-    train_mask = np.loadtxt(path + 'train_mask.txt')
-    train_mask = train_mask.astype(int)
-
-    labels = np.loadtxt(path + 'labels.txt')
-    labels = labels.astype(int)
-    '''
-    edges = np.load(path + 'edges.npy')
-    features = np.load(path + 'features.npy')
-    train_mask = np.load(path + 'train_mask.npy')
-    labels = np.load(path + 'labels.npy')
+# GPU | CPU
+def get_default_device():
     
-    num_edges = edges.shape[0]
-    num_nodes = features.shape[0]
-    num_feats = features.shape[1]
-    n_classes = int(max(labels) - min(labels) + 1)
-
-    assert train_mask.shape[0] == num_nodes
+    if torch.cuda.is_available():
+        return torch.device('cuda:0')
+    else:
+        return torch.device('cpu')
 
-    print('dataset {}'.format(args.dataset))
-    print('# of edges : {}'.format(num_edges))
-    print('# of nodes : {}'.format(num_nodes))
-    print('# of features : {}'.format(num_feats))
+def to_default_device(data):
     
-    features = torch.FloatTensor(features)
-    labels = torch.LongTensor(labels)
-
-    print("Features Shape")
-    print(features.shape)
+    if isinstance(data,(list,tuple)):
+        return [to_default_device(x,get_default_device()) for x in data]
+    
+    return data.to(get_default_device(),non_blocking = True)
 
-    if hasattr(torch, 'BoolTensor'):
-        train_mask = torch.BoolTensor(train_mask)
+def main(args):
+    
+    cora = CoraDataset(verbose=True)
+    
+    features = torch.FloatTensor(cora.get_all_features())
+    labels = torch.LongTensor(cora.get_all_targets())
+    
+    train_mask = cora.get_train_mask()
+    test_mask = cora.get_test_mask()
 
-    else:
-        train_mask = torch.ByteTensor(train_mask)
+    train_mask = torch.BoolTensor(train_mask)
+    test_mask = torch.BoolTensor(test_mask)
 
     if args.gpu < 0:
         cuda = False
@@ -78,33 +66,27 @@ def main(args):
         features = features.cuda()
         labels = labels.cuda()
         train_mask = train_mask.cuda()
+        test_mask = test_mask.cuda()
 
-    '''
-    # graph preprocess and calculate normalization factor
-    g = data.graph
-    # add self loop
-    if args.self_loop:
-        g.remove_edges_from(nx.selfloop_edges(g))
-        g.add_edges_from(zip(g.nodes(), g.nodes()))
-    g = DGLGraph(g)
-    n_edges = g.number_of_edges()
-    '''
-    u = edges[:,0]
-    v = edges[:,1]
-    g = dgl.graph((u,v), num_nodes=num_nodes)
+    g = StaticGraph(cora.get_edges())
+    
     # add self loop
-    if args.self_loop:
-        g = dgl.add_self_loop(g)
-    g = g.to(features.device)
+    # if args.self_loop:
+    #     g = dgl.add_self_loop(g)
+    
+    # g = g.to(features.device)
 
     # normalization
-    degs = g.in_degrees().float()
+    degs = torch.from_numpy(g.in_degrees()).type(torch.int32)
     norm = torch.pow(degs, -0.5)
     norm[torch.isinf(norm)] = 0
-    if cuda:
-        norm = norm.cuda()
+    
+    norm = to_default_device(norm)
     g.ndata['norm'] = norm.unsqueeze(1)
 
+    num_feats = features.shape[1]
+    n_classes = int(max(labels) - min(labels) + 1)
+
     model = EglGCN(g,
                 num_feats,
                 args.num_hidden,
@@ -160,14 +142,14 @@ def main(args):
 
 if __name__ == '__main__':
     parser = argparse.ArgumentParser(description='GCN')
-    #add_argument --dataset
-    register_data_args(parser)
 
     # COMMENT IF SNOOP IS TO BE ENABLED
     snoop.install(enabled=False)
 
     parser.add_argument("--dropout", type=float, default=0.5,
             help="dropout probability")
+    parser.add_argument("--dataset", type=str,
+            help="Datset to train your model")
     parser.add_argument("--gpu", type=int, default=0,
             help="gpu")
     parser.add_argument("--lr", type=float, default=1e-2,
diff --git a/exp/tgat/dgl/model.py b/benchmarking/tgat/dgl/model.py
similarity index 100%
rename from exp/tgat/dgl/model.py
rename to benchmarking/tgat/dgl/model.py
diff --git a/exp/tgat/dgl/train.py b/benchmarking/tgat/dgl/train.py
similarity index 100%
rename from exp/tgat/dgl/train.py
rename to benchmarking/tgat/dgl/train.py
diff --git a/exp/tgat/seastar/tgat.py b/benchmarking/tgat/seastar/tgat.py
similarity index 97%
rename from exp/tgat/seastar/tgat.py
rename to benchmarking/tgat/seastar/tgat.py
index a3b844a..8e799a9 100644
--- a/exp/tgat/seastar/tgat.py
+++ b/benchmarking/tgat/seastar/tgat.py
@@ -3,9 +3,9 @@
 import torch.nn as nn
 import dgl
 import dgl.function as fn
-from seastar import CtxManager
+from seastar import Seastar
 import torch.nn.functional as F
-from seastar.backend.pytorch_backend import run_egl
+from seastar.backend.pytorch_backend import backend_cb
 import snoop
 import inspect
 
@@ -40,7 +40,7 @@ def __init__(self,
             self.res_fc = None
         self.reset_parameters()
         self.activation = activation
-        self.cm = CtxManager(run_egl)
+        self.seastar = Seastar(backend_cb)
 
     def reset_parameters(self):
         """Reinitialize learnable parameters."""
@@ -58,7 +58,7 @@ def forward(self, feat):
         er = (feat_dst * self.attn_r).sum(dim=-1).unsqueeze(-1)
 
         # Vertex-centric implementation.
-        @self.cm.zoomIn(nspace=[self, torch])
+        @self.seastar.compile(nspace=[self, torch])
         def nb_forward(v):
            embs = [nb.el + v.er for nb in v.innbs]
            coeff = [torch.exp(self.leaky_relu(emb - max(embs))) for emb in embs]
diff --git a/exp/tgat/seastar/train.py b/benchmarking/tgat/seastar/train.py
similarity index 100%
rename from exp/tgat/seastar/train.py
rename to benchmarking/tgat/seastar/train.py
diff --git a/benchmarking/tgcn/dynamic/tgcn.py b/benchmarking/tgcn/dynamic/tgcn.py
new file mode 100644
index 0000000..7011812
--- /dev/null
+++ b/benchmarking/tgcn/dynamic/tgcn.py
@@ -0,0 +1,157 @@
+import math
+import torch
+import torch.nn as nn
+import dgl
+import dgl.function as fn
+from seastar.compiler import Seastar
+import torch.nn.functional as F
+from seastar.compiler.backend.pytorch_backend import backend_cb
+import snoop
+from rich import inspect
+
+class SeastarGCNLayer(nn.Module):
+    def __init__(self,
+                 in_feats,
+                 out_feats,
+                 activation = None,
+                 dropout = None,
+                 bias=True):
+        super(SeastarGCNLayer, self).__init__()
+        self.weight = nn.Parameter(torch.Tensor(in_feats, out_feats))
+        if bias:
+            self.bias = nn.Parameter(torch.Tensor(out_feats))
+        else:
+            self.bias = None
+        self.activation = activation
+        if dropout:
+            self.dropout = nn.Dropout(p=dropout)
+        else:
+            self.dropout = 0.
+        self.reset_parameters()
+        self.seastar = Seastar(backend_cb)
+
+    def reset_parameters(self):
+        stdv = 1. / math.sqrt(self.weight.size(1))
+        self.weight.data.uniform_(-stdv, stdv)
+        if self.bias is not None:
+            self.bias.data.uniform_(-stdv, stdv)
+
+    def forward(self, g, h, edge_weight=None):
+        if self.dropout:
+            h = self.dropout(h)
+
+        h = torch.mm(h, self.weight)
+
+        @self.seastar.compile(nspace=[self, torch])
+        def nb_compute(v):
+            # The nb_edge.src returns a list with one element, this element is an object of NbNode type
+            # hence the translation. Can be cleaned up later.
+            h = sum([nb_edge.src[0].norm * nb_edge.src[0].h * nb_edge.weight for nb_edge in v.inedges])
+            h = h * v.norm
+            return h
+        h = nb_compute(g=g, n_feats={'norm': g.ndata['norm'], 'h' : h}, e_feats={'weight': edge_weight})
+
+        # @self.seastar.compile(nspace=[self, torch])
+        # def nb_compute(v):
+        #     # The nb_edge.src returns a list with one element, this element is an object of NbNode type
+        #     # hence the translation. Can be cleaned up later.
+        #     h = sum([nb.h*nb.norm for nb in v.innbs])
+        #     h = h * v.norm
+        #     return h
+        # h = nb_compute(g=g, n_feats={'norm': g.ndata['norm'], 'h' : h})
+
+        # bias
+        if self.bias is not None:
+            h = h + self.bias
+        if self.activation is not None:
+            h = self.activation(h)
+        return h
+
+class SeastarTGCNCell(torch.nn.Module):
+
+    def __init__(
+        self,
+        in_channels: int,
+        out_channels: int,
+    ):
+        super(SeastarTGCNCell, self).__init__()
+
+        self.in_channels = in_channels
+        self.out_channels = out_channels
+
+        # Update GCN Layer
+        self.conv_z = SeastarGCNLayer(self.in_channels,self.out_channels)
+        
+        # Update linear layer
+        self.linear_z = torch.nn.Linear(2 * self.out_channels, self.out_channels)
+
+        # Reset GCN layer
+        self.conv_r = SeastarGCNLayer(self.in_channels,self.out_channels)
+        
+        # Reset linear layer
+        self.linear_r = torch.nn.Linear(2 * self.out_channels, self.out_channels)
+
+        # Candidate (Current Memory Content) GCN layer
+        self.conv_h = SeastarGCNLayer(self.in_channels,self.out_channels)
+
+        # Candidate linear layer
+        self.linear_h = torch.nn.Linear(2 * self.out_channels, self.out_channels)
+
+    def _set_hidden_state(self, X, H):
+        if H is None:
+            H = torch.zeros(X.shape[0], self.out_channels).to(X.device)
+        return H
+
+    def _calculate_update_gate(self, g, X, edge_weight, H):
+        h = self.conv_z(g, X, edge_weight=edge_weight)
+        # print(h.shape)
+        # print(H.shape)
+        Z = torch.cat((h, H), axis=1) # axis values need to be checked
+        Z = self.linear_z(Z)
+        Z = torch.sigmoid(Z)
+        return Z
+
+    def _calculate_reset_gate(self, g, X, edge_weight, H):
+        R = torch.cat((self.conv_r(g, X, edge_weight=edge_weight), H), axis=1) # axis values need to be checked
+        R = self.linear_r(R)
+        R = torch.sigmoid(R)
+        return R
+
+    def _calculate_candidate_state(self, g, X, edge_weight, H, R):
+        H_tilde = torch.cat((self.conv_h(g, X, edge_weight=edge_weight), H * R), axis=1) # axis values need to be checked
+        H_tilde = self.linear_h(H_tilde)
+        H_tilde = torch.tanh(H_tilde)
+        return H_tilde
+
+    def _calculate_hidden_state(self, Z, H, H_tilde):
+        H = Z * H + (1 - Z) * H_tilde
+        return H
+
+    def forward(
+        self,
+        g,
+        X: torch.Tensor,
+        edge_weight: torch.FloatTensor = None,
+        H: torch.FloatTensor = None,
+    ) -> torch.FloatTensor:
+
+        H = self._set_hidden_state(X, H)
+        Z = self._calculate_update_gate(g, X, edge_weight, H)
+        R = self._calculate_reset_gate(g, X, edge_weight, H)
+        H_tilde = self._calculate_candidate_state(g, X, edge_weight, H, R)
+        H = self._calculate_hidden_state(Z, H, H_tilde)
+        return H
+        
+
+class SeastarTGCN(torch.nn.Module):
+  def __init__(self, node_features):
+    super(SeastarTGCN, self).__init__()
+    self.temporal = SeastarTGCNCell(node_features, 32)
+    self.linear = torch.nn.Linear(32, 1)
+
+  def forward(self, g, node_feat, edge_weight, hidden_state):
+    h = self.temporal(g, node_feat, edge_weight, hidden_state)
+    y = F.relu(h)
+    y = self.linear(y)
+    
+    return y, h
\ No newline at end of file
diff --git a/benchmarking/tgcn/dynamic/train.py b/benchmarking/tgcn/dynamic/train.py
new file mode 100644
index 0000000..bc091f7
--- /dev/null
+++ b/benchmarking/tgcn/dynamic/train.py
@@ -0,0 +1,195 @@
+import argparse, time
+import numpy as np
+import networkx as nx
+import torch
+import torch.nn as nn
+import torch.nn.functional as F
+import json
+import urllib
+from tqdm import tqdm
+from tgcn import SeastarTGCN
+import snoop
+import os
+
+from rich import inspect
+from rich.pretty import pprint
+# from rich.traceback import install
+# install(show_locals=True)
+
+from seastar.graph.dynamic.gpma.GPMAGraph import GPMAGraph
+from seastar.graph.dynamic.pcsr.PCSRGraph import PCSRGraph
+from seastar.graph.dynamic.naive.NaiveGraph import NaiveGraph
+from seastar.dataset.EnglandCOVID import EnglandCOVID
+
+
+# GPU | CPU
+def get_default_device():
+    
+    if torch.cuda.is_available():
+        return torch.device('cuda:0')
+    else:
+        return torch.device('cpu')
+
+def to_default_device(data):
+    
+    if isinstance(data,(list,tuple)):
+        return [to_default_device(x,get_default_device()) for x in data]
+    
+    return data.to(get_default_device(),non_blocking = True)
+
+def main(args):
+
+    if torch.cuda.is_available():
+        print("🎉 CUDA is available")
+    else:
+        print("😔 CUDA is not available")
+    
+    eng_covid = EnglandCOVID(verbose=True)
+    
+    edge_list = eng_covid.get_edges()
+    edge_weight_list = eng_covid.get_edge_weights()
+    all_features = eng_covid.get_all_features()
+    all_targets = eng_covid.get_all_targets()
+    
+    all_features = to_default_device(torch.FloatTensor(np.array(all_features)))
+    all_targets = to_default_device(torch.FloatTensor(np.array(all_targets)))
+    
+    # Hyperparameters
+    train_test_split = 0.8
+    
+    # train_test_split for graph
+    train_edges_lst = edge_list[:int(len(edge_list) * train_test_split)]
+    test_edges_lst = edge_list[int(len(edge_list) * train_test_split):]
+    
+    train_edge_weights_lst = edge_weight_list[:int(len(edge_weight_list) * train_test_split)]
+    test_edge_weights_lst = edge_weight_list[int(len(edge_weight_list) * train_test_split):]
+
+    # train_test_split for features
+    train_features = all_features[:int(len(all_features) * train_test_split)]
+    train_targets = all_targets[:int(len(all_targets) * train_test_split)]
+    
+    test_features = all_features[int(len(all_features) * train_test_split):]
+    test_targets = all_targets[int(len(all_targets) * train_test_split):]
+
+    model = to_default_device(SeastarTGCN(8))
+
+    # use optimizer
+    optimizer = torch.optim.Adam(model.parameters(), lr=args.lr)
+
+    # metrics
+    dur = []
+    Used_memory = 0
+    cuda = True
+
+    G = GPMAGraph(train_edges_lst)
+    # G = PCSRGraph(train_edges_lst)
+    # G = NaiveGraph(train_edges_lst)
+
+    # train
+    print("Training...\n")
+    for epoch in range(args.num_epochs):
+        model.train()
+        if cuda:
+            torch.cuda.synchronize()
+        t0 = time.time()
+
+        cost = 0
+        hidden_state = None
+        optimizer.zero_grad()
+
+        # dyn_graph_index is dynamic graph index
+        for index in range(0,len(train_features)): 
+            
+            # Getting the graph for a particular timestamp
+            G.get_graph(index) 
+
+            # normalization
+            degs = torch.from_numpy(G.in_degrees()).type(torch.int32)
+            norm = torch.pow(degs, -0.5)
+            norm[torch.isinf(norm)] = 0
+            norm = to_default_device(norm)
+            G.ndata['norm'] = norm.unsqueeze(1)
+            edge_weight = to_default_device(torch.FloatTensor(train_edge_weights_lst[index]))
+            edge_weight = torch.unsqueeze(edge_weight,1)
+
+            # forward propagation
+            y_hat, hidden_state = model(G, train_features[index], edge_weight, hidden_state)
+            
+            cost = cost + torch.mean((y_hat-train_targets[index])**2)
+            
+    
+        cost = cost / (index+1)
+
+        now_mem = torch.cuda.max_memory_allocated(0)
+        Used_memory = max(now_mem, Used_memory)
+        
+        cost.backward()
+        optimizer.step()
+
+        if cuda:
+            torch.cuda.synchronize()
+
+        run_time_this_epoch = time.time() - t0
+
+        if epoch >= 3:
+            dur.append(run_time_this_epoch)
+
+        print('Epoch {:05d} | Time(s) {:.4f} | MSE {:.6f} | Used_Memory {:.6f} mb'.format(
+            epoch, run_time_this_epoch, cost, (now_mem * 1.0 / (1024**2))
+        ))
+
+    Used_memory /= (1024**3)
+    print('^^^{:6f}^^^{:6f}'.format(Used_memory, np.mean(dur)))
+
+    # evaluate
+    # print("Evaluating...\n")
+    # model.eval()
+    # cost = 0
+
+    # test_graph_log_dict, test_max_num_nodes = preprocess_graph_structure(test_edges_lst)
+    # G = GPMAGraph(test_graph_log_dict,test_max_num_nodes)
+    
+    # # G = PCSRGraph(test_graph_log_dict,test_max_num_nodes)
+
+    # predictions = []
+    # true_y = []
+    # hidden_state=None
+    # # dyn_graph_index is dynamic graph index
+    # for index in range(len(test_features)):
+    #     # normalization
+    #     # degs = G.in_degrees().float()
+    #     degs = torch.from_numpy(G.in_degrees())
+    #     norm = torch.pow(degs, -0.5)
+    #     norm[torch.isinf(norm)] = 0
+    #     norm = to_default_device(norm)
+    #     G.ndata['norm'] = norm.unsqueeze(1)
+    #     edge_weight = to_default_device(torch.FloatTensor(test_edge_weights_lst[index]))
+    #     edge_weight = torch.unsqueeze(edge_weight,1)
+        
+    #     # forward propagation
+    #     y_hat, hidden_state = model(G, test_features[index], edge_weight, hidden_state)
+    #     cost = cost + torch.mean((y_hat-test_targets[index])**2)
+    #     predictions.append(y_hat)
+    #     true_y.append(test_targets[index])
+
+    # cost = cost / (index+1)
+    # cost = cost.item()
+    # print("MSE: {:.4f}".format(cost))
+
+
+
+if __name__ == '__main__':
+    parser = argparse.ArgumentParser(description='GCN')
+
+    # COMMENT IF SNOOP IS TO BE ENABLED
+    snoop.install(enabled=False)
+
+
+    parser.add_argument("--lr", type=float, default=1e-2,
+            help="learning rate")
+    parser.add_argument("--num_epochs", type=int, default=1,
+            help="number of training epochs")
+    args = parser.parse_args()
+    print(args)
+
+    main(args)
\ No newline at end of file
diff --git a/exp/tgcn/seastar/tgcn.py b/benchmarking/tgcn/seastar/tgcn.py
similarity index 96%
rename from exp/tgcn/seastar/tgcn.py
rename to benchmarking/tgcn/seastar/tgcn.py
index e411548..b6b73f3 100644
--- a/exp/tgcn/seastar/tgcn.py
+++ b/benchmarking/tgcn/seastar/tgcn.py
@@ -3,9 +3,9 @@
 import torch.nn as nn
 import dgl
 import dgl.function as fn
-from seastar import CtxManager
+from seastar import Seastar
 import torch.nn.functional as F
-from seastar.backend.pytorch_backend import run_egl
+from seastar.backend.pytorch_backend import backend_cb
 import snoop
 import inspect
 
@@ -31,7 +31,7 @@ def __init__(self,
         else:
             self.dropout = 0.
         self.reset_parameters()
-        self.cm = CtxManager(run_egl)
+        self.seastar = Seastar(backend_cb)
 
     def reset_parameters(self):
         stdv = 1. / math.sqrt(self.weight.size(1))
@@ -45,7 +45,7 @@ def forward(self, h, edge_weight=None):
 
         h = torch.mm(h, self.weight)
 
-        @self.cm.zoomIn(nspace=[self, torch])
+        @self.seastar.compile(nspace=[self, torch])
         def nb_compute(v):
             # The nb_edge.src returns a list with one element, this element is an object of NbNode type
             # hence the translation. Can be cleaned up later.
diff --git a/exp/tgcn/seastar/train.py b/benchmarking/tgcn/seastar/train.py
similarity index 100%
rename from exp/tgcn/seastar/train.py
rename to benchmarking/tgcn/seastar/train.py
diff --git a/exp/gcn/seastar/__pycache__/gcn_spmv.cpython-38.pyc b/exp/gcn/seastar/__pycache__/gcn_spmv.cpython-38.pyc
deleted file mode 100644
index d10b544..0000000
Binary files a/exp/gcn/seastar/__pycache__/gcn_spmv.cpython-38.pyc and /dev/null differ
diff --git a/python/dynamic_graph/PCSR/README.md b/python/dynamic_graph/PCSR/README.md
deleted file mode 100644
index 1716c7f..0000000
--- a/python/dynamic_graph/PCSR/README.md
+++ /dev/null
@@ -1,40 +0,0 @@
-# PCSR Python Module
-
-The ```pcsr``` python module is written using C++. The API is connected using PyBind11.
-
-The C++ code used for implementing PCSR is used from ![Packed CSR](https://github.com/wheatman/Packed-Compressed-Sparse-Row).
-
-## Making changes to the API
-
-After making any changes to the API in the ```pcsr.cpp``` file, run the following command to create the ```pcsr.cpython-310-x86_64-linux-gnu.so``` file.
-
-```
-c++ -O3 -Wall -shared -std=c++11 -fPIC $(python3 -m pybind11 --includes) pcsr.cpp -g -o pcsr$(python3-config --extension-suffix)
-```
-
-Make sure to PyBind11 installed in whichever environment you are working on.
-
-## Using the pcsr module
-
-Like any other module, you can import ```pcsr``` as follows. Run the following code to get started.
-
-```
-import pcsr
-
-graph = pcsr.PCSR()
-graph.init_graph("./datasets/test_graph.txt")
-graph.print_graph()
-```
-
-Output:
-
-```
-   0   1   2
-0  •       •
-1          •
-2  •   •   •
-```
-
-## Methods provided by pcsr
-
-TODO: Add documentation for each method
\ No newline at end of file
diff --git a/python/dynamic_graph/PCSR/main.py b/python/dynamic_graph/PCSR/main.py
deleted file mode 100644
index 64741da..0000000
--- a/python/dynamic_graph/PCSR/main.py
+++ /dev/null
@@ -1,5 +0,0 @@
-import pcsr
-
-graph = pcsr.PCSR()
-graph.init_graph("./datasets/test_graph.txt")
-graph.print_graph()
\ No newline at end of file
diff --git a/python/dynamic_graph/PCSR/pcsr.cpp b/python/dynamic_graph/PCSR/pcsr.cpp
deleted file mode 100644
index a6a97ea..0000000
--- a/python/dynamic_graph/PCSR/pcsr.cpp
+++ /dev/null
@@ -1,696 +0,0 @@
-#include <pybind11/pybind11.h>
-#include <pybind11/stl.h>
-
-#include <algorithm>
-#include <queue>
-#include <stdint.h>
-#include <stdio.h>
-#include <stdlib.h>
-#include <tuple>
-#include <vector>
-#include <string>
-#include <cstring>
-#include <iostream>
-
-namespace py = pybind11;
-using namespace pybind11::literals;
-using namespace std;
-
-////////////////////////////////////////////////////////////////////////////
-
-typedef struct _node {
-  // beginning and end of the associated region in the edge list
-  uint32_t beginning;     // deleted = max int
-  uint32_t end;           // end pointer is exclusive
-  uint32_t num_neighbors; // number of edges with this node as source
-
-  _node(int beg=0, int _end=0, int num_neigh=0) {
-    beginning = beg;
-    end = _end;
-    num_neighbors = num_neigh;
-  }
-} node_t;
-
-typedef struct _edge {
-  uint32_t dest;    // destination of this edge in the graph, MAX_INT if this is a sentinel
-  uint32_t value;   // edge value of zero means it a null since we don't store 0 edges
-
-  _edge(int _dest=0, int _value=0) {
-    dest = _dest;
-    value = _value;
-  }
-} edge_t;
-
-typedef struct edge_list {
-  int N;
-  int H;
-  int logN;
-  vector<edge_t> items;
-
-  edge_list() {
-    N = 0;
-    H = 0;
-    logN = 0;
-
-    vector<edge_t> temp(0,0);
-    items = temp;
-  }
-} edge_list_t;
-
-typedef struct _pair_int {
-    int x; // length in array
-    int y; // depth
-
-    _pair_int(int _x=0, int _y=0){
-        x = _x;
-        y = _y;
-    }
-} pair_int;
-
-typedef struct _pair_double {
-    double x;
-    double y;
-
-    _pair_double(double _x=0, double _y=0){
-        x = _x;
-        y = _y;
-    }
-} pair_double;
-
-////////////////////////////////////////////////////////////////////////////
-
-static inline int bsf_word(int word) {
-  int result;
-  __asm__ volatile("bsf %1, %0" : "=r"(result) : "r"(word));
-  return result;
-}
-
-static inline int bsr_word(int word) {
-  int result;
-  __asm__ volatile("bsr %1, %0" : "=r"(result) : "r"(word));
-  return result;
-}
-
-// given index, return the starting index of the leaf it is in
-int find_leaf(edge_list_t *list, int index) {
-  return (index / list->logN) * list->logN;
-}
-
-bool is_null(edge_t e) { return e.value == 0; }
-
-bool is_sentinel(edge_t e) {
-  return e.dest == UINT32_MAX || e.value == UINT32_MAX;
-}
-
-uint32_t binary_search(edge_list_t *list, edge_t *elem, uint32_t start,
-                       uint32_t end) {
-  while (start + 1 < end) {
-    uint32_t mid = (start + end) / 2;
-
-    edge_t item = list->items[mid];
-    uint32_t change = 1;
-    uint32_t check = mid;
-
-    bool flag = true;
-    while (is_null(item) && flag) {
-      flag = false;
-      check = mid + change;
-      if (check < end) {
-        flag = true;
-        if (check <= end) {
-          item = list->items[check];
-          if (!is_null(item)) {
-            break;
-          } else if (check == end) {
-            break;
-          }
-        }
-      }
-      check = mid - change;
-      if (check >= start) {
-        flag = true;
-        item = list->items[check];
-      }
-      change++;
-    }
-
-    if (is_null(item) || start == check || end == check) {
-      if (!is_null(item) && start == check && elem->dest <= item.dest) {
-        return check;
-      }
-      return mid;
-    }
-
-    // if we found it, return
-    if (elem->dest == item.dest) {
-      return check;
-    } else if (elem->dest < item.dest) {
-      end =
-          check; // if the searched for item is less than current item, set end
-    } else {
-      start = check;
-      // otherwise, searched for item is more than current and we set start
-    }
-  }
-  if (end < start) {
-    start = end;
-  }
-  // handling the case where there is one element left
-  // if you are leq, return start (index where elt is)
-  // otherwise, return end (no element greater than you in the range)
-  // printf("start = %d, end = %d, n = %d\n", start,end, list->N);
-  if (elem->dest <= list->items[start].dest && !is_null(list->items[start])) {
-    return start;
-  }
-  return end;
-}
-
-// get density of a node
-double get_density(edge_list_t *list, int index, int len) {
-  int full = 0;
-  for (int i = index; i < index + len; i++) {
-    full += (!is_null(list->items[i]));
-  }
-  double full_d = (double)full;
-  return full_d / len;
-}
-
-int find_node(int index, int len) { return (index / len) * len; }
-
-pair_double density_bound(edge_list_t *list, int depth) {
-  pair_double pair;
-
-  // between 1/4 and 1/2
-  // pair.x = 1.0/2.0 - (( .25*depth)/list->H);
-  // between 1/8 and 1/4
-  pair.x = 1.0 / 4.0 - ((.125 * depth) / list->H);
-  pair.y = 3.0 / 4.0 + ((.25 * depth) / list->H);
-  return pair;
-}
-
-bool edge_equals(edge_t e1, edge_t e2) {
-  return e1.dest == e2.dest && e1.value == e2.value;
-}
-
-uint32_t find_elem_pointer(edge_list_t *list, uint32_t index, edge_t elem) {
-  edge_t item = list->items[index];
-  while (!edge_equals(item, elem)) {
-    item = list->items[++index];
-  }
-  return index;
-}
-
-////////////////////////////////////////////////////////////////////////////
-
-class PCSR {
-    public:
-      // data members
-      std::vector<node_t> nodes;
-      edge_list_t edges;
-
-      // member functions
-      PCSR(uint32_t init_n);
-      void init_graph(std::string graph_file_path);
-
-      void add_node();
-      void add_edge(uint32_t src, uint32_t dest, uint32_t value);
-      void add_edge_update(uint32_t src, uint32_t dest, uint32_t value);
-      void print_graph();
-      void print_array();
-      py::dict get_csr_arrays();
-
-      uint32_t insert(uint32_t index, edge_t elem, uint32_t src);
-      void double_list();
-      int slide_right(int index);
-      void slide_left(int index);
-      void fix_sentinel(int32_t node_index, int in);
-      void redistribute(int index, int len);
-};
-
-////////////////////////////////////////////////////////////////////////////
-
-PCSR::PCSR(uint32_t init_n=0) {
-  if (init_n != 0){
-    edges.N = 2 << bsr_word(init_n);
-    edges.logN = (1 << bsr_word(bsr_word(edges.N) + 1));
-    edges.H = bsr_word(edges.N / edges.logN);
-
-    for (int i = 0; i < edges.N; i++) {
-        edge_t new_edge(0,0);
-        edges.items.push_back(new_edge);
-    }
-
-    for (int i = 0; i < init_n; i++) {
-        add_node();
-    }
-  }
-}
-
-void PCSR::init_graph(std::string graph_file_path) {
-
-  FILE *fp;
-  fp = fopen(graph_file_path.c_str(), "r");
-  if (!fp) {
-      printf("Open file failed.\n");
-      exit(-1);
-  }
-
-  // loading the number of nodes and edges present in
-  // the graph into the variables node_size and edge_size respectively
-  int node_size, edge_size;
-  fscanf(fp, "%d %d", &node_size, &edge_size);
-
-  edges.N = 2 << bsr_word(node_size);
-  edges.logN = (1 << bsr_word(bsr_word(edges.N) + 1));
-  edges.H = bsr_word(edges.N / edges.logN);
-
-  for (int i = 0; i < edges.N; i++) {
-    edge_t new_edge(0,0);
-    edges.items.push_back(new_edge);
-  }
-
-  for (int i = 0; i < node_size; i++) {
-    add_node();
-  }
-
-  for (int i = 0; i < edge_size; i++) {
-    int source, dest, edge_value;
-    fscanf(fp, "%d %d %d", &source, &dest, &edge_value);
-    add_edge(source, dest, edge_value);
-  }
-}
-
-// add a node to the graph
-void PCSR::add_node() {
-    node_t node;
-    int len = nodes.size();
-
-    edge_t sentinel;
-    sentinel.dest = UINT32_MAX; // placeholder
-    sentinel.value = len;       // back pointer
-
-    if (len > 0) {
-        node.beginning = nodes[len - 1].end;
-        node.end = node.beginning + 1;
-    } else {
-        node.beginning = 0;
-        node.end = 1;
-        sentinel.value = UINT32_MAX;
-    }
-    node.num_neighbors = 0;
-
-    nodes.push_back(node);
-    insert(node.beginning, sentinel, nodes.size() - 1);
-}
-
-uint32_t PCSR::insert(uint32_t index, edge_t elem, uint32_t src) {
-    int node_index = find_leaf(&edges, index);
-    int level = edges.H;
-    int len = edges.logN;
-
-    // always deposit on the left
-    if (is_null(edges.items[index])) {
-        edges.items[index].value = elem.value;
-        edges.items[index].dest = elem.dest;
-    } else {
-        // if the edge already exists in the graph, update its value
-        // do not make another edge
-        // return index of the edge that already exists
-        if (!is_sentinel(elem) && edges.items[index].dest == elem.dest) {
-        edges.items[index].value = elem.value;
-        return index;
-        }
-        if (index == edges.N - 1) {
-        // when adding to the end double then add edge
-        double_list();
-        node_t node = nodes[src];
-        uint32_t loc_to_add =
-            binary_search(&edges, &elem, node.beginning + 1, node.end);
-        return insert(loc_to_add, elem, src);
-        } else {
-        if (slide_right(index) == -1) {
-            index -= 1;
-            slide_left(index);
-        }
-        }
-        edges.items[index].value = elem.value;
-        edges.items[index].dest = elem.dest;
-    }
-
-    double density = get_density(&edges, node_index, len);
-
-    // spill over into next level up, node is completely full.
-    if (density == 1) {
-        node_index = find_node(node_index, len * 2);
-        redistribute(node_index, len * 2);
-    } else {
-        // makes the last slot in a section empty so you can always slide right
-        redistribute(node_index, len);
-    }
-
-    // get density of the leaf you are in
-    pair_double density_b = density_bound(&edges, level);
-    density = get_density(&edges, node_index, len);
-
-    // while density too high, go up the implicit tree
-    // go up to the biggest node above the density bound
-    while (density >= density_b.y) {
-        len *= 2;
-        if (len <= edges.N) {
-        level--;
-        node_index = find_node(node_index, len);
-        density_b = density_bound(&edges, level);
-        density = get_density(&edges, node_index, len);
-        } else {
-        // if you reach the root, double the list
-        double_list();
-
-        // search from the beginning because list was doubled
-        return find_elem_pointer(&edges, 0, elem);
-        }
-    }
-    redistribute(node_index, len);
-
-    return find_elem_pointer(&edges, node_index, elem);
-}
-
-void PCSR::double_list() {
-    edges.N *= 2;
-    edges.logN = (1 << bsr_word(bsr_word(edges.N) + 1));
-    edges.H = bsr_word(edges.N / edges.logN);
-
-    for (int i = edges.N / 2; i < edges.N; i++) {
-        edge_t new_edge(0,0);
-        edges.items.push_back(new_edge);
-    }
-
-    redistribute(0, edges.N);
-}
-
-int PCSR::slide_right(int index) {
-  int rval = 0;
-  edge_t el = edges.items[index];
-  edges.items[index].dest = 0;
-  edges.items[index].value = 0;
-  index++;
-  while (index < edges.N && !is_null(edges.items[index])) {
-    edge_t temp = edges.items[index];
-    edges.items[index] = el;
-    if (!is_null(el) && is_sentinel(el)) {
-      // fixing pointer of node that goes to this sentinel
-      uint32_t node_index = el.value;
-      if (node_index == UINT32_MAX) {
-        node_index = 0;
-      }
-      fix_sentinel(node_index, index);
-    }
-    el = temp;
-    index++;
-  }
-  if (!is_null(el) && is_sentinel(el)) {
-    // fixing pointer of node that goes to this sentinel
-    uint32_t node_index = el.value;
-    if (node_index == UINT32_MAX) {
-      node_index = 0;
-    }
-    fix_sentinel(node_index, index);
-  }
-  // TODO There might be an issue with this going of the end sometimes
-  if (index == edges.N) {
-    index--;
-    slide_left(index);
-    rval = -1;
-    printf("slide off the end on the right, should be rare\n");
-  }
-  edges.items[index] = el;
-  return rval;
-}
-
-void PCSR::slide_left(int index) {
-  edge_t el = edges.items[index];
-  edges.items[index].dest = 0;
-  edges.items[index].value = 0;
-
-  index--;
-  while (index >= 0 && !is_null(edges.items[index])) {
-    edge_t temp = edges.items[index];
-    edges.items[index] = el;
-    if (!is_null(el) && is_sentinel(el)) {
-      // fixing pointer of node that goes to this sentinel
-      uint32_t node_index = el.value;
-      if (node_index == UINT32_MAX) {
-        node_index = 0;
-      }
-
-      fix_sentinel(node_index, index);
-    }
-    el = temp;
-    index--;
-  }
-
-  if (index == -1) {
-    double_list();
-
-    slide_right(0);
-    index = 0;
-  }
-  if (!is_null(el) && is_sentinel(el)) {
-    // fixing pointer of node that goes to this sentinel
-    uint32_t node_index = el.value;
-    if (node_index == UINT32_MAX) {
-      node_index = 0;
-    }
-    fix_sentinel(node_index, index);
-  }
-
-  edges.items[index] = el;
-}
-
-void PCSR::fix_sentinel(int32_t node_index, int in) {
-  nodes[node_index].beginning = in;
-  if (node_index > 0) {
-    nodes[node_index - 1].end = in;
-  }
-  if (node_index == nodes.size() - 1) {
-    nodes[node_index].end = edges.N - 1;
-  }
-}
-
-void PCSR::redistribute(int index, int len) {
-  // printf("REDISTRIBUTE: \n");
-  // print_array();
-  // std::vector<edge_t> space(len); //
-  edge_t new_edge;
-  vector<edge_t> space(len, new_edge);
-  
-  int j = 0;
-
-  // move all items in ofm in the range into
-  // a temp array
-  for (int i = index; i < index + len; i++) {
-    space[j] = edges.items[i];
-    // counting non-null edges
-    j += (!is_null(edges.items[i]));
-    // setting section to null
-    edges.items[i].value = 0;
-    edges.items[i].dest = 0;
-  }
-
-  // evenly redistribute for a uniform density
-  double index_d = index;
-  double step = ((double)len) / j;
-  for (int i = 0; i < j; i++) {
-    int in = index_d;
-
-    edges.items[in] = space[i];
-    if (is_sentinel(space[i])) {
-      // fixing pointer of node that goes to this sentinel
-      uint32_t node_index = space[i].value;
-      if (node_index == UINT32_MAX) {
-        node_index = 0;
-      }
-      fix_sentinel(node_index, in);
-    }
-    index_d += step;
-  }
-//   free(space);
-}
-
-void PCSR::add_edge(uint32_t src, uint32_t dest, uint32_t value) {
-  if (value != 0) {
-    node_t node = nodes[src];
-    nodes[src].num_neighbors++;
-
-    edge_t e;
-    e.dest = dest;
-    e.value = value;
-
-    uint32_t loc_to_add =
-        binary_search(&edges, &e, node.beginning + 1, node.end);
-    insert(loc_to_add, e, src);
-  }
-}
-
-void PCSR::add_edge_update(uint32_t src, uint32_t dest, uint32_t value) {
-  if (value != 0) {
-    node_t node = nodes[src];
-
-    edge_t e;
-    e.dest = dest;
-    e.value = value;
-
-    uint32_t loc_to_add =
-        binary_search(&edges, &e, node.beginning + 1, node.end);
-    if (edges.items[loc_to_add].dest == dest) {
-      edges.items[loc_to_add].value = value;
-      return;
-    }
-    nodes[src].num_neighbors++;
-    insert(loc_to_add, e, src);
-  }
-}
-
-void PCSR::print_graph() {
-  int num_vertices = nodes.size();
-
-  // printing the graph matrix column indices
-  for (int i = 0; i < num_vertices; ++i)
-    printf("   %d", i);
-  
-  printf("\n");
-
-  for (int i = 0; i < num_vertices; i++) {
-    // +1 to avoid sentinel
-
-    // printing the graph matrix row indices
-    printf("%d ", i);
-    int matrix_index = 0;
-
-    for (uint32_t j = nodes[i].beginning + 1; j < nodes[i].end; j++) {
-      if (!is_null(edges.items[j])) {
-        while (matrix_index < edges.items[j].dest) {
-          printf("    ");
-          matrix_index++;
-        }
-        // printf("%03d ", edges.items[j].value);
-        printf(" •  ");
-        matrix_index++;
-      }
-    }
-    for (uint32_t j = matrix_index; j < num_vertices; j++) {
-      printf("    ");
-    }
-    printf("\n");
-  }
-}
-
-void PCSR::print_array() {
-  for (int i = 0; i < edges.N; i++) {
-    if (is_null(edges.items[i])) {
-      printf("%d-x ", i);
-    } else if (is_sentinel(edges.items[i])) {
-      uint32_t value = edges.items[i].value;
-      if (value == UINT32_MAX) {
-        value = 0;
-      }
-      printf("\n%d-s(%u):(%d, %d) ", i, value, nodes[value].beginning,
-             nodes[value].end);
-    } else {
-      printf("%d-(%d, %u) ", i, edges.items[i].dest, edges.items[i].value);
-    }
-  }
-  printf("\n\n");
-}
-
-py::dict PCSR::get_csr_arrays() {
-  std::vector<int> beg_array;
-  std::vector<int> end_array;
-  std::vector<int> num_neighbors_array;
-  std::vector<int> dest_array;
-  std::vector<int> value_array;
-
-  std::vector<int> row_offset;
-  std::vector<int> column_indices;
-
-  int row_offset_size = nodes.size() + 1;
-  int column_indices_size = edges.items.size();
-
-  for(int i=0; i<row_offset_size-1; ++i){
-    row_offset.push_back(nodes[i].beginning);
-    num_neighbors_array.push_back(nodes[i].num_neighbors);
-  }
-
-  row_offset.push_back(nodes[row_offset_size-2].end);
-
-  for(int i=0; i<column_indices_size; ++i) {
-    if(is_sentinel(edges.items[i]) || is_null(edges.items[i])){
-      column_indices.push_back(-1);
-    }
-    else{
-      column_indices.push_back(edges.items[i].dest);
-    }
-  }
-
-  py::list beg_list = py::cast(beg_array);
-  py::list end_list = py::cast(end_array);
-  py::list num_neigh_list = py::cast(num_neighbors_array);
-
-  py::dict csr_arrays("row_offset"_a = row_offset, "column_indices"_a = column_indices);
-  return csr_arrays;
-}
-
-////////////////////////////////////////////////////////////////////////////
-
-// PCSR Python Module
-
-PYBIND11_MODULE(pcsr, m) {
-    m.doc() = "PCSR Dynamic Graph Representation";
-
-    py::class_<node_t>(m, "Node")
-        .def(py::init<int, int, int>(), py::arg("beg")=0, py::arg("_end")=0, py::arg("num_neigh")=0)
-        .def_readwrite("beginning", &_node::beginning)
-        .def_readwrite("end", &_node::end)
-        .def_readwrite("num_neighbors", &_node::num_neighbors);
-
-    py::class_<edge_t>(m, "Edge")
-        .def(py::init<int, int>(), py::arg("_dest")=0, py::arg("_value")=0)
-        .def_readwrite("dest", &edge_t::dest)
-        .def_readwrite("value", &edge_t::value);
-
-    py::class_<edge_list_t>(m, "EdgeList")
-        .def(py::init<>())
-        .def_readwrite("N", &edge_list_t::N)
-        .def_readwrite("H", &edge_list_t::H)
-        .def_readwrite("logN", &edge_list_t::logN)
-        .def_readwrite("items", &edge_list_t::items);
-
-    py::class_<pair_int>(m, "PairInt")
-        .def(py::init<int, int>())
-        .def_readwrite("x", &pair_int::x)
-        .def_readwrite("y", &pair_int::y);
-
-    py::class_<pair_double>(m, "PairDouble")
-        .def(py::init<double, double>())
-        .def_readwrite("x", &pair_double::x)
-        .def_readwrite("y", &pair_double::y);
-
-    py::class_<PCSR>(m, "PCSR")
-        .def(py::init<int>(), py::arg("init_n")=0)
-        .def("init_graph", &PCSR::init_graph)
-        .def("add_node", &PCSR::add_node)
-        .def("insert", &PCSR::insert)
-        .def("double_list", &PCSR::double_list)
-        .def("slide_right", &PCSR::slide_right)
-        .def("slide_left", &PCSR::slide_left)
-        .def("fix_sentinel", &PCSR::fix_sentinel)
-        .def("redistribute", &PCSR::redistribute)
-        .def("print_graph", &PCSR::print_graph)
-        .def("add_edge", &PCSR::add_edge)
-        .def("add_edge_update", &PCSR::add_edge_update)
-        .def("print_array", &PCSR::print_array)
-        .def("get_csr_arrays", &PCSR::get_csr_arrays)
-        .def_readwrite("nodes", &PCSR::nodes)
-        .def_readwrite("edges", &PCSR::edges);
-}
-
-// c++ -O3 -Wall -shared -std=c++11 -fPIC $(python3 -m pybind11 --includes) pcsr.cpp -o pcsr$(python3-config --extension-suffix)
\ No newline at end of file
diff --git a/python/dynamic_graph/PCSR/pcsr.cpython-310-x86_64-linux-gnu.so b/python/dynamic_graph/PCSR/pcsr.cpython-310-x86_64-linux-gnu.so
deleted file mode 100644
index 7753816..0000000
Binary files a/python/dynamic_graph/PCSR/pcsr.cpython-310-x86_64-linux-gnu.so and /dev/null differ
diff --git a/python/dynamic_graph/__pycache__/dynamic_graph.cpython-310.pyc b/python/dynamic_graph/__pycache__/dynamic_graph.cpython-310.pyc
deleted file mode 100644
index 7249224..0000000
Binary files a/python/dynamic_graph/__pycache__/dynamic_graph.cpython-310.pyc and /dev/null differ
diff --git a/python/dynamic_graph/dynamic_graph.py b/python/dynamic_graph/dynamic_graph.py
deleted file mode 100644
index 7b18dd7..0000000
--- a/python/dynamic_graph/dynamic_graph.py
+++ /dev/null
@@ -1,50 +0,0 @@
-from PCSR import pcsr
-
-from rich.pretty import pprint
-from rich import inspect
-
-def pretty_print_pcsr(graph):
-    print("Edges:")
-    inspect(graph.edges)
-
-    print("Items:")
-    for edge in graph.edges.items:
-        inspect(edge)
-
-    print("Nodes:")
-    inspect(graph.nodes)
-
-    for node in graph.nodes:
-        inspect(node)
-
-class DynamicGraph:
-    def __init__(self, graph_type: str, graph_file_path: str):
-        self.graph_type = graph_type
-        if self.graph_type == "pcsr":
-            self.graph = pcsr.PCSR()
-            self.graph.init_graph(graph_file_path)
-
-    def update_graph(self, operations_list):
-        for operation in operations_list:
-            action, src, dst, value = operation
-
-            # add an edge
-            if action == "ae":
-                if self.graph_type == "pcsr":
-                    self.graph.add_edge(src, dst, value)
-
-    def print_graph(self):
-        if self.graph_type == "pcsr":
-            self.graph.print_graph()
-
-    def print_array(self):
-        if self.graph_type == "pcsr":
-            self.graph.print_array()
-
-    def get_csr_arrays(self):
-        if self.graph_type == "pcsr":
-            return self.graph.get_csr_arrays()
-
-    def debug_graph(self):
-        if self.graph_type == "pcsr":
-            pretty_print_pcsr(self.graph)
\ No newline at end of file
diff --git a/python/dynamic_graph/main.py b/python/dynamic_graph/main.py
deleted file mode 100644
index a6f95de..0000000
--- a/python/dynamic_graph/main.py
+++ /dev/null
@@ -1,21 +0,0 @@
-from timeit import default_timer as timer
-
-from rich.pretty import pprint
-from dynamic_graph import DynamicGraph
-
-
-graph = DynamicGraph(
-    graph_type="pcsr",
-    graph_file_path="./PCSR/datasets/test_graph.txt"
-)
-
-print("\nIntial Graph\n")
-graph.print_graph()
-graph.print_array()
-
-start = timer()
-graph_dict = graph.get_csr_arrays()
-end = timer()
-pprint(graph_dict)
-
-print(f'\nTime taken to create CSR Arrays: {end-start}\n')
\ No newline at end of file
diff --git a/python/seastar/code_gen/template.py b/python/seastar/code_gen/template.py
deleted file mode 100644
index a44f4dd..0000000
--- a/python/seastar/code_gen/template.py
+++ /dev/null
@@ -1,121 +0,0 @@
-from jinja2 import Template
-
-from collections import namedtuple
-
-from .compiler import compile_cuda
-
-header_tpl = Template(
-"""
-{%for header in headers%}
-{{'#include ' + header}}{%endfor%}
-"""
-)
-
-tpl_v2 = Template(
-"""
-extern "C" __global__ void {{kernel_name}}({%for arg in args%}{{arg.type}} {{'*' if arg.is_ptr}}{{arg.name}}, {% endfor %}
-  {{index_type}} *row_offsets,
-  {{index_type}} *eids,
-  {{index_type}} *column_indices,
-  {{index_type}} num_nodes,
-  {{index_type}} max_dimx,
-  {{index_type}} max_dimy,
-  {{index_type}} tile_sizex,
-  {{index_type}} tile_sizey) {
-    for ({{index_type}} {{row_offset}} = blockIdx.y; {{row_offset}} < num_nodes; {{row_offset}} += gridDim.y) {
-        {{index_type}} start_off = row_offsets[{{row_offset}}];
-        {{index_type}} end_off = row_offsets[{{row_offset}} + 1];
-        {{index_type}} blk_tid = threadIdx.x + threadIdx.y * blockDim.x;
-        {{index_type}} tid = blk_tid + blockIdx.x * blockDim.x * blockDim.y;
-        {{index_type}} warp_id = tid/32;
-        {{index_type}} warp_tid = tid%32;
-        {{index_type}} num_warp_x = max_dimx/tile_sizex;
-        {{index_type}} tx = warp_tid % tile_sizex + warp_id % num_warp_x * tile_sizex; 
-        {{index_type}} ty = warp_tid / tile_sizex + warp_id / num_warp_x * tile_sizey;
-        for (;tx < max_dimx; tx += blockDim.x*gridDim.x) {
-            for (;ty < max_dimy; ty += blockDim.y) {
-                {%for agg_stmt in aggs%}{{agg_stmt.init}}{%endfor%}
-                {{init_outter_offset}}
-                for ({{index_type}} e=start_off;e<end_off;++e) {
-                    {{index_type}} {{col_index}} = column_indices[e];
-                    {{index_type}} eid = eids[e];
-                    {{init_inner_offset}}
-                    {%for edge_stmt in edges%}
-                    {{edge_stmt.load}}
-                    {{edge_stmt.compute}}
-                    {{edge_stmt.inner_write}}{%endfor%}
-                    {%for agg_stmt in aggs%}
-                    {{agg_stmt.compute}}
-                    {{agg_stmt.inner_write}}{%endfor%}
-                }
-                {%for agg_stmt in aggs%}
-                {{agg_stmt.outter_write}}{%endfor%}
-                {%for node_stmt in nodes%}
-                {{node_stmt.load}}{{node_stmt.compute}}{{node_stmt.inner_write}}{%endfor%}
-            }
-        }
-    }
-}
-"""
-)
-
-tpl_fa = Template(
-"""
-extern "C" __global__ void {{kernel_name}}({%for arg in args%}{{arg.type}} {{'*' if arg.is_ptr}}{{arg.name}}, {% endfor %}
-  {{index_type}} *row_offsets,
-  {{index_type}} *eids,
-  {{index_type}} *column_indices,
-  {{index_type}} num_nodes,
-  {{index_type}} max_dimx,
-  {{index_type}} max_dimy,
-  {{index_type}} thrs_per_group,
-  {{index_type}} nodes_per_block) {
-    {{index_type}} {{row_offset}} = nodes_per_block*blockIdx.x + threadIdx.x/thrs_per_group;;
-    if ({{row_offset}} < num_nodes) {
-        {{index_type}} feat_len = max_dimx * max_dimy;
-        {{index_type}} beg = __ldg(row_offsets + {{row_offset}});
-        {{index_type}} end = __ldg(row_offsets + {{row_offset}} + 1);
-        {{index_type}} tx = threadIdx.x % thrs_per_group;
-        for (; tx<feat_len; tx+=blockDim.x) {
-            {%for agg_stmt in aggs%}{{agg_stmt.init}}{%endfor%}
-            {{init_outter_offset}}
-            for ({{index_type}} e=beg;e<end;++e) {
-                {{index_type}} {{col_index}} = __ldg(column_indices + e);
-                {{index_type}} eid = __ldg(eids + e);
-                {{init_inner_offset}}
-                {%for edge_stmt in edges%}
-                {{edge_stmt.load}}
-                {{edge_stmt.compute}}
-                {{edge_stmt.inner_write}}{%endfor%}
-                {%for agg_stmt in aggs%}
-                {{agg_stmt.compute}}
-                {{agg_stmt.inner_write}}{%endfor%}
-            }
-            {%for agg_stmt in aggs%}
-            {{agg_stmt.outter_write}}{%endfor%}
-            {%for node_stmt in nodes%}
-            {{node_stmt.load}}{{node_stmt.compute}}{{node_stmt.inner_write}}{%endfor%}
-        }
-    }
-}
-"""
-)
-
-EdgeInfo = namedtuple('EdgeInfo', ['load', 'compute', 'inner_write'])
-NodeInfo = namedtuple('NodeInfo', ['load', 'compute', 'inner_write'])
-ArgInfo = namedtuple('ArgInfo', ['name', 'type', 'is_ptr'])
-AggInfo = namedtuple('AggInfo', ['init', 'compute', 'inner_write', 'outter_write'])
-
-def gen_cuda(configs):
-    h = ''
-    for config in configs:
-        if config['template_name'] == 'fa':
-            rendered_tpl = tpl_fa.render(**config)
-        elif config['template_name'] == 'v2':
-            rendered_tpl = tpl_v2.render(**config)
-        else:
-            raise NotImplementedError('Have not implement template for', configs['template_name'])
-        h += rendered_tpl
-    # UNCOMMENT IF KERNEL IS TO BE PRINTED    
-    print(h)
-    return compile_cuda(h)
diff --git a/python/seastar/passes/.gitignore b/python/seastar/passes/.gitignore
deleted file mode 100644
index bee8a64..0000000
--- a/python/seastar/passes/.gitignore
+++ /dev/null
@@ -1,1 +0,0 @@
-__pycache__
diff --git a/python/seastar/zoom.py b/python/seastar/zoom.py
deleted file mode 100644
index 8b08e12..0000000
--- a/python/seastar/zoom.py
+++ /dev/null
@@ -1,273 +0,0 @@
-import torch
-from collections import defaultdict, namedtuple
-from collections.abc import Iterable
-import sys
-
-from .node import CentralNode
-from .val import create_edge_val, create_src_node_val, create_dst_node_val, create_param_val
-from .op import create_op, AggMaxOp, AggMinOp, AggMeanOp
-from .program import Var, Stmt, Program
-from .passes import optimize, CF, fuse, visualize
-from .schema import Schema
-from .autodiff import diff
-from .code_gen import code_gen 
-from .executor import Executor
-from .utils import var_prefix, cen_attr_postfix, inb_attr_postfix
-
-class SkipContextException(Exception):
-    pass
-
-class zoomIn(object):
-    def __init__(self, namespace, input_map, node_feats, edge_feats, fprog):
-        self._namespace = namespace
-        self._backend = self._find_backend(namespace)
-        self._nfeats = node_feats 
-        self._efeats = edge_feats
-        self._v = CentralNode()
-        self._old_mapping = defaultdict(dict)
-        self.inputs = input_map
-        self.fprog = fprog
-        self.entry_count = 0
-    
-    def clear_cache(self):
-        if self.inputs:
-            del self._namespace
-            del self._backend
-            del self._nfeats
-            del self._efeats
-            del self._v
-            del self._old_mapping
-            del self.inputs
-            del self.fprog
-
-    def __enter__(self):
-        if self.entry_count == 0:
-            print('entering local context')
-            self._symbolization()
-            self.entry_count += 1
-            return self._v
-        else:
-            sys.settrace(lambda *args, **kargs: None)
-            frame = sys._getframe(1)
-            frame.f_trace = self.trace
-        
-    def trace(self, frame, event, arg):
-        raise SkipContextException
-
-    def __exit__(self, ctx_type, ctx_value, ctx_traceback):
-        if ctx_type == SkipContextException:
-            return True
-        self._desymbolization()
-        print('TracedProgram:\n' + str(self.fprog))
-    
-    def mapping_key(self, name_space_id, original_key):
-        return str(name_space_id) + str(original_key)
-
-    def _symbolization(self):
-        """Symbolizing central node and its innbs and inedges""" 
-        var_name = ''
-        for k, v in self._nfeats.items():
-            setattr(self._v, k, create_dst_node_val(v, self._backend, id=k+cen_attr_postfix, fprog=self.fprog))
-            var_name += '%' + k  + ','
-            for n in self._v.innbs:
-                setattr(n, k, create_src_node_val(v, self._backend, id=k+inb_attr_postfix, fprog=self.fprog))
-        for k, v in self._efeats.items():
-            var_name += '%' + k  + ','
-            for e in self._v.inedges:
-                setattr(e, k, create_edge_val(v, self._backend, id=k, fprog=self.fprog))
-        self.fprog.set_input(var_name.strip(','))
-        if self._backend[0] == 'torch':
-            for i, nspace in enumerate(self._namespace):
-                if '__name__' in nspace.__dict__:
-                    # symbolizing functions for torch namespace
-                    k = self.mapping_key(i, 'function')
-                    for key in nspace.__dict__:
-                        m = nspace.__dict__[key]
-                        if 'function' in str(type(m)):
-                            if key in self._old_mapping[k]:
-                                raise KeyError('Found', key, ' already in _old_mapping')
-                            self._old_mapping[k][key] = m
-                            nspace.__dict__[key] = create_op(m, self._backend[0], fprog=self.fprog)
-                else:
-                    ## Dealing with module
-                    for key in nspace.__dict__.keys():
-                        # symbolizing parameters for self namespace
-                        m = nspace.__dict__[key]
-                        k = self.mapping_key(i, key)
-                        if key.startswith('_parameters'):
-                            for mkey in m.keys():
-                                if mkey in self._old_mapping[k]:
-                                    raise KeyError('Found', key, ' already in _old_mapping')
-                                self._old_mapping[k][mkey] = m[mkey] 
-                                self.inputs[var_prefix+mkey] = m[mkey]
-                                m[mkey] = create_param_val(m[mkey], self._backend, id=mkey, fprog=self.fprog)
-
-                        # symbolizing buffers for self namespace
-                        if key.startswith('_buffers'):
-                            for mkey in m.keys():
-                                if mkey in self._old_mapping[k]:
-                                    raise KeyError('Found', key, 'already in _old_mapping')
-                                self._old_mapping[k][mkey] = m[mkey]
-                                self.inputs[var_prefix+mkey] = m[mkey]
-                                m[mkey] = create_param_val(m[mkey], self._backend, id=mkey, fprog=self.fprog)
-
-                        # symbolizing modules for self namespace
-                        if key.startswith('_modules'):
-                            for mkey in m.keys():
-                                if mkey in self._old_mapping[k]:
-                                    raise KeyError('Found', key, ' already in _old_mapping')
-                                self._old_mapping[k][mkey] = m[mkey]
-                                m[mkey] = create_op(m[mkey], self._backend[0], fprog=self.fprog)
-        else:
-            raise NotImplementedError('Backend ' + self._backend[0] + ' is not supported yet!') 
-
-    def _desymbolization(self):
-        if self._backend[0] == 'torch':
-            for i, nspace in enumerate(self._namespace):
-                if '__name__' in nspace.__dict__:
-                    # desymbolizing functions for torch namespace
-                    if 'torch' in nspace.__name__.lower():
-                        k = self.mapping_key(i, 'function')
-                        for key in self._old_mapping[k]:
-                            nspace.__dict__[key] = self._old_mapping[k][key]
-                else:
-                    for key in nspace.__dict__.keys():
-                        # desymbolizing parameters for self namespace
-                        m = nspace.__dict__[key]
-                        k = self.mapping_key(i, key) 
-                        if key.startswith('_parameters') or key.startswith('_modules') or key.startswith('_buffers'):
-                            for mkey in self._old_mapping[k]:
-                                m[mkey] = self._old_mapping[k][mkey]
-
-        else:
-            raise NotImplementedError('Backend ' + self._backend[0] + ' is not supported yet!') 
-
-    def _find_backend(self, namespace):
-        k =  '__name__'
-        for n in namespace:
-            if k in n.__dict__:
-                name = n.__dict__[k].lower()
-                if 'torch' in name:
-                    return ('torch', n)
-                elif 'tensorflow' in name:
-                    return ('tensorflow', n)
-                elif 'mxnet' in name:
-                    return ('mxnet', n)
-                else:
-                    raise NotImplementedError("Backend support for " + name + " is not implemnted yet")
-    
-    def outputs(self):
-        return self._v.outputs
-
-def zoomOut(out_set, fprog, graph):
-    optimize(fprog)
-    vars = []
-    for var in out_set:
-        vars.append(var)
-    forward_exe_units = fuse([fprog], vars)
-    grads = []
-    for var in vars:
-        grads.append(Var.create_var(var_shape=var.var_shape, var_dtype=var.var_dtype, val_type=var.val_type, device=var.device))
-    backward_exe_units, full_grad_map = diff(vars, grads, forward_exe_units)
-    visualize.plot_exec_units(forward_exe_units + backward_exe_units)
-    compiled_module = code_gen.gen_code(forward_exe_units + backward_exe_units, 'int' if graph.nbits == 32 else 'long long int')
-    ex = Executor(graph, forward_exe_units, backward_exe_units, compiled_module, full_grad_map, vars)
-    return ex 
-
-class ContextManager(object):
-    GraphInfo = namedtuple('GraphInfo', ['number_of_nodes',
-                                         'number_of_edges',
-                                         'in_row_offsets',
-                                         'in_col_indices',
-                                         'in_eids',
-                                         'out_row_offsets',
-                                         'out_col_indices',
-                                         'out_eids',
-                                         'nbits'])
-    graph_info = None
-    cached_graph = None
-    reset = False
-
-    def __init__(self, run_egl):
-        self._executor = None
-        self._ctx = None
-        self.input_map = {}
-        self._fprog = Program()
-        self.run_egl_cb = run_egl
-        self.outputs = set()
-
-    def zoomIn(self, namespace=None, graph=None, node_feats={}, edge_feats={}):
-        if ContextManager.graph_info == None or id(graph) != id(ContextManager.cached_graph):
-            ContextManager.cached_graph = graph
-            ContextManager.reset = True
-            ## Assuming different context are working on the same graph
-            # This may not be true for mini-batch training
-            in_csr = graph.get_in_csr()
-            out_csr = graph.get_out_csr()
-            ContextManager.graph_info = ContextManager.GraphInfo(
-                graph.number_of_nodes(),
-                graph.number_of_edges(),
-                in_csr(0).copy_to_gpu(0),
-                in_csr(1).copy_to_gpu(0),
-                in_csr(2).copy_to_gpu(0),
-                out_csr(0).copy_to_gpu(0),
-                out_csr(1).copy_to_gpu(0),
-                out_csr(2).copy_to_gpu(0),
-                graph.nbits()
-            )
-        for k, v in node_feats.items():
-            self.input_map[var_prefix + k + cen_attr_postfix] = v
-            self.input_map[var_prefix + k + inb_attr_postfix] = v
-        for k, v in edge_feats.items():
-            self.input_map[var_prefix+k] = v
-        if not self._ctx:
-            self._ctx = zoomIn(namespace, self.input_map, node_feats, edge_feats, self._fprog)
-        return self._ctx
-
-    def zoomOut(self):
-        if not self._executor:
-            self._executor = zoomOut(self.outputs, self._fprog, ContextManager.graph_info)
-            ContextManager.reset = False
-        self._executor.restart(self.input_map, ContextManager.graph_info if ContextManager.reset else None)
-        ret = self.run_egl_cb(self._executor)
-        if len(ret) == 1:
-            return ret[0]
-        self.input_map.clear()
-        self._ctx.clear_cache()
-        return ret
-
-    def collect_output(self, val):
-        if isinstance(val, list):
-            for v in val:
-                self.outputs.add(v.var)
-        else:
-            self.outputs.add(val.var)
-    
-    def update_allnode(self, dict):
-        cen=self._ctx._v
-        for k,v in dict.items():
-            if k in cen.__dict__:
-                cen.__dict__[k] = v
-            for nb in cen.innbs:
-                if k in nb.__dict__:
-                    # Here we re-use the var id of central node to indicate the fact that
-                    # they are the same tensor with different access modifier(src and dst)
-                    nb.__dict__[k] = create_src_node_val(v._v, v.backend, id=str(v.var.int_id), fprog=self._fprog, reduce_dim=False)
-                    nb.__dict__[k].var._stmt = Stmt.create_stmt(op_schema=Schema('GTypeCast'), args=[v.var], ret=nb.__dict__[k].var) 
-                    self._fprog.append_stmt(nb.__dict__[k].var._stmt)
-                    nb.__dict__[k].var._requires_grad = v.var._requires_grad
-
-    def max(self, args):
-        assert isinstance(args, list), 'Can only aggregate feature list'
-        max_op = AggMaxOp()
-        return max_op(self._fprog, args)
-
-    def min(self, args):
-        assert isinstance(args, list), 'Can only aggregate feature list'
-        min_op = AggMinOp()
-        return min_op(self._fprog, args)
-
-    def mean(self, args):
-        assert isinstance(args, list), 'Can only aggregate feature list'
-        mean_op = AggMeanOp()
-        return mean_op(self._fprog, args)
diff --git a/seastar/__init__.py b/seastar/__init__.py
new file mode 100644
index 0000000..43e8e03
--- /dev/null
+++ b/seastar/__init__.py
@@ -0,0 +1,1 @@
+'''Seastar Module'''
\ No newline at end of file
diff --git a/python/seastar/__init__.py b/seastar/compiler/__init__.py
similarity index 50%
rename from python/seastar/__init__.py
rename to seastar/compiler/__init__.py
index 599b9fe..f47f15e 100644
--- a/python/seastar/__init__.py
+++ b/seastar/compiler/__init__.py
@@ -1,4 +1,3 @@
-from .zoom import *
-from .zoom_v2 import *
+from .seastar import *
 from .program import *
 from .schema import *
\ No newline at end of file
diff --git a/python/seastar/autodiff.py b/seastar/compiler/autodiff.py
similarity index 100%
rename from python/seastar/autodiff.py
rename to seastar/compiler/autodiff.py
diff --git a/python/seastar/backend/__init__.py b/seastar/compiler/backend/__init__.py
similarity index 100%
rename from python/seastar/backend/__init__.py
rename to seastar/compiler/backend/__init__.py
diff --git a/python/seastar/backend/pytorch_backend.py b/seastar/compiler/backend/pytorch_backend.py
similarity index 97%
rename from python/seastar/backend/pytorch_backend.py
rename to seastar/compiler/backend/pytorch_backend.py
index 7334a5a..5fb33c2 100644
--- a/python/seastar/backend/pytorch_backend.py
+++ b/seastar/compiler/backend/pytorch_backend.py
@@ -13,7 +13,7 @@ def backward(ctx, *gradout):
         executor, kid = ctx.backward_cache
         return (None, None, None, None) +  executor.backward_cb(kid, gradout)
 
-def run_egl(executor):
+def backend_cb(executor):
 
     def new_zeros_call_back(size, dtype, device, requires_grad=True):
         return torch.zeros(size=size, dtype=dtype, device=device, requires_grad=requires_grad)
diff --git a/python/seastar/code_gen/__init__.py b/seastar/compiler/code_gen/__init__.py
similarity index 100%
rename from python/seastar/code_gen/__init__.py
rename to seastar/compiler/code_gen/__init__.py
diff --git a/python/seastar/code_gen/code_gen.py b/seastar/compiler/code_gen/code_gen.py
similarity index 60%
rename from python/seastar/code_gen/code_gen.py
rename to seastar/compiler/code_gen/code_gen.py
index 0fbc298..c16c17b 100644
--- a/python/seastar/code_gen/code_gen.py
+++ b/seastar/compiler/code_gen/code_gen.py
@@ -1,5 +1,12 @@
-from .template import gen_cuda, EdgeInfo, AggInfo, ArgInfo, NodeInfo
 from ..utils import is_const_scalar, ParallelMode
+from collections import namedtuple
+from .compiler import compile_cuda
+from jinja2 import Environment, PackageLoader
+
+EdgeInfo = namedtuple('EdgeInfo', ['load', 'compute', 'inner_write'])
+NodeInfo = namedtuple('NodeInfo', ['load', 'compute', 'inner_write'])
+ArgInfo = namedtuple('ArgInfo', ['name', 'type', 'is_ptr'])
+AggInfo = namedtuple('AggInfo', ['init', 'compute', 'inner_write', 'outter_write'])
 
 const_id = 0
 def gen_arg_info(arg):
@@ -29,7 +36,7 @@ def gen_node_info(stmt, ctx):
         raise NotImplementedError('Cannot generate code for', stmt)
     return NodeInfo(**m)
 
-def gen_code(exe_units, index_type):
+def gen_code(exe_units, index_type, graph_type):
     '''Generating cuda code by instantiate code template'''
     if not isinstance(exe_units, list):
         exe_units = [exe_units]
@@ -71,5 +78,34 @@ def gen_code(exe_units, index_type):
             'col_index': 'src_id' if dst_parallel else 'dst_id',
             'init_inner_offset': (ctx.src_var_offset_init if dst_parallel else ctx.dst_var_offset_init) + ctx.edge_var_offset_init,
             'template_name': ctx.template_name,
+            'graph_type': graph_type
         })
     return gen_cuda(configs)
+
+def render_template(config, template_name):
+    env = Environment(
+    loader=PackageLoader("seastar.compiler.code_gen"),
+    )
+    tpl = env.get_template("fa/{}.jinja".format(template_name))
+    return tpl.render(**config)
+
+def gen_cuda(configs):
+    h = ''
+    for config in configs:
+        if config['template_name'] == 'fa':
+            if config['graph_type'] == 'csr':
+                rendered_tpl = render_template(config, "tpl_fa_csr")
+            elif config['graph_type'] == 'pcsr': 
+                rendered_tpl = render_template(config, "tpl_fa_pcsr")
+            elif config['graph_type'] == 'gpma':   
+                rendered_tpl = render_template(config, "tpl_fa_gpma")
+            else:
+                raise NotImplementedError('{} Template for {} is not supported'.format(config['template_name'],config['graph_type']))
+        elif config['template_name'] == 'v2':
+            raise NotImplementedError('{} Template for {} is not supported'.format(config['template_name'],config['graph_type']))
+        else:
+            raise NotImplementedError('{} Template not supported'.format(config['template_name']))
+        h += rendered_tpl
+    # UNCOMMENT IF KERNEL IS TO BE PRINTED    
+    print(h)
+    return compile_cuda(h)
diff --git a/python/seastar/code_gen/compiler.py b/seastar/compiler/code_gen/compiler.py
similarity index 100%
rename from python/seastar/code_gen/compiler.py
rename to seastar/compiler/code_gen/compiler.py
diff --git a/python/seastar/code_gen/cuda_driver.py b/seastar/compiler/code_gen/cuda_driver.py
similarity index 97%
rename from python/seastar/code_gen/cuda_driver.py
rename to seastar/compiler/code_gen/cuda_driver.py
index d565181..aa84ed1 100644
--- a/python/seastar/code_gen/cuda_driver.py
+++ b/seastar/compiler/code_gen/cuda_driver.py
@@ -3,7 +3,7 @@
 import sys
 
 if 'linux' in sys.platform:
-	cuda = CDLL('libcuda.so')
+	cuda = CDLL('/usr/lib/wsl/lib/libcuda.so')
 elif 'win' in sys.platform:
 	cuda = CDLL('nvcuda.dll')
 
diff --git a/python/seastar/code_gen/device_info.py b/seastar/compiler/code_gen/device_info.py
similarity index 97%
rename from python/seastar/code_gen/device_info.py
rename to seastar/compiler/code_gen/device_info.py
index 58538f3..e3dde11 100644
--- a/python/seastar/code_gen/device_info.py
+++ b/seastar/compiler/code_gen/device_info.py
@@ -15,7 +15,7 @@
 class deviceinfo():
 
     def __init__(self):
-        libnames = ('libcuda.so', 'libcuda.dylib', 'cuda.dll')
+        libnames = ('/usr/lib/wsl/lib/libcuda.so', 'libcuda.dylib', 'cuda.dll')
         for libname in libnames:
             try:
                 cuda = ctypes.CDLL(libname)
diff --git a/python/seastar/code_gen/kernel_context.py b/seastar/compiler/code_gen/kernel_context.py
similarity index 100%
rename from python/seastar/code_gen/kernel_context.py
rename to seastar/compiler/code_gen/kernel_context.py
diff --git a/python/seastar/debugging/__init__.py b/seastar/compiler/code_gen/templates/__init__.py
similarity index 100%
rename from python/seastar/debugging/__init__.py
rename to seastar/compiler/code_gen/templates/__init__.py
diff --git a/seastar/compiler/code_gen/templates/fa/__init__.py b/seastar/compiler/code_gen/templates/fa/__init__.py
new file mode 100644
index 0000000..e69de29
diff --git a/seastar/compiler/code_gen/templates/fa/tpl_fa_csr.jinja b/seastar/compiler/code_gen/templates/fa/tpl_fa_csr.jinja
new file mode 100644
index 0000000..1d97dcd
--- /dev/null
+++ b/seastar/compiler/code_gen/templates/fa/tpl_fa_csr.jinja
@@ -0,0 +1,56 @@
+extern "C" __global__ void {{kernel_name}}
+({%for arg in args%}{{arg.type}} {{'*' if arg.is_ptr}}{{arg.name}}, {% endfor %}
+  {{index_type}} *row_offsets,
+  {{index_type}} *eids,
+  {{index_type}} *column_indices,
+  {{index_type}} num_nodes,
+  {{index_type}} max_dimx,
+  {{index_type}} max_dimy,
+  {{index_type}} thrs_per_group,
+  {{index_type}} nodes_per_block) {
+      
+    {{index_type}} {{row_offset}} = nodes_per_block*blockIdx.x + threadIdx.x/thrs_per_group;
+    
+    if ({{row_offset}} < num_nodes) {
+        
+        {{index_type}} feat_len = max_dimx * max_dimy;
+        {{index_type}} beg = __ldg(row_offsets + {{row_offset}});
+        {{index_type}} end = __ldg(row_offsets + {{row_offset}} + 1);
+        {{index_type}} tx = threadIdx.x % thrs_per_group;
+        
+        for (; tx<feat_len; tx+=blockDim.x) {
+            
+            {%for agg_stmt in aggs%}{{agg_stmt.init}}{%endfor%}
+            {{init_outter_offset}}
+            
+            for ({{index_type}} e=beg;e<end;++e) {
+                
+                {{index_type}} {{col_index}} = __ldg(column_indices + e);
+                {{index_type}} eid = __ldg(eids + e);
+                
+                {{init_inner_offset}}
+                
+                {%for edge_stmt in edges%}
+                {{edge_stmt.load}}
+                {{edge_stmt.compute}}
+                {{edge_stmt.inner_write}}
+                {%endfor%}
+                
+                {%for agg_stmt in aggs%}
+                {{agg_stmt.compute}}
+                {{agg_stmt.inner_write}}
+                {%endfor%}
+            }
+            
+            {%for agg_stmt in aggs%}
+            {{agg_stmt.outter_write}}
+            {%endfor%}
+            
+            {%for node_stmt in nodes%}
+            {{node_stmt.load}}
+            {{node_stmt.compute}}
+            {{node_stmt.inner_write}}
+            {%endfor%}
+        }
+    }
+}
\ No newline at end of file
diff --git a/seastar/compiler/code_gen/templates/fa/tpl_fa_gpma.jinja b/seastar/compiler/code_gen/templates/fa/tpl_fa_gpma.jinja
new file mode 100644
index 0000000..6a9c7bc
--- /dev/null
+++ b/seastar/compiler/code_gen/templates/fa/tpl_fa_gpma.jinja
@@ -0,0 +1,65 @@
+extern "C" __global__ void {{kernel_name}}(
+  {%for arg in args%}{{arg.type}} {{'*' if arg.is_ptr}}{{arg.name}}, {% endfor %}
+  unsigned int *row_offsets,
+  unsigned int *eids,
+  unsigned long long *column_indices,
+  {{index_type}} num_nodes,
+  {{index_type}} max_dimx,
+  {{index_type}} max_dimy,
+  {{index_type}} thrs_per_group,
+  {{index_type}} nodes_per_block) {
+
+    {{index_type}} {{row_offset}} = nodes_per_block*blockIdx.x + threadIdx.x/thrs_per_group;
+
+    if ({{row_offset}} < num_nodes) {
+
+        {{index_type}} feat_len = max_dimx * max_dimy;
+        unsigned int beg = __ldg(row_offsets + {{row_offset}});
+        unsigned int end = __ldg(row_offsets + {{row_offset}} + 1);
+        {{index_type}} tx = threadIdx.x % thrs_per_group;
+
+        for (; tx<feat_len; tx+=blockDim.x) {
+
+            {%for agg_stmt in aggs%}{{agg_stmt.init}}{%endfor%}
+            {{init_outter_offset}}
+
+            for ({{index_type}} e=beg;e<end;++e) {
+
+                unsigned long long {{col_index}} = __ldg(column_indices + e);
+
+                // GPMA indexes edges starting from 1
+                // Seastar requires edgs to be indexed from 0
+                unsigned int eid = __ldg(eids + e) - 1;
+                
+                unsigned long long mask = (unsigned long long){{row_offset}} << 32;
+                unsigned int dst_check = ({{col_index}} - mask);
+                {{col_index}} = ({{col_index}} - mask);
+
+                if(dst_check != 0xFFFFFFFF && eid != 0){
+                    {{init_inner_offset}}
+
+                    {%for edge_stmt in edges%}
+                    {{edge_stmt.load}}
+                    {{edge_stmt.compute}}
+                    {{edge_stmt.inner_write}}
+                    {%endfor%}
+
+                    {%for agg_stmt in aggs%}
+                    {{agg_stmt.compute}}
+                    {{agg_stmt.inner_write}}
+                    {%endfor%}    
+                }
+            }
+
+            {%for agg_stmt in aggs%}
+            {{agg_stmt.outter_write}}
+            {%endfor%}
+
+            {%for node_stmt in nodes%}
+            {{node_stmt.load}}
+            {{node_stmt.compute}}
+            {{node_stmt.inner_write}}
+            {%endfor%}
+        }
+    }
+}
\ No newline at end of file
diff --git a/seastar/compiler/code_gen/templates/fa/tpl_fa_pcsr.jinja b/seastar/compiler/code_gen/templates/fa/tpl_fa_pcsr.jinja
new file mode 100644
index 0000000..9285593
--- /dev/null
+++ b/seastar/compiler/code_gen/templates/fa/tpl_fa_pcsr.jinja
@@ -0,0 +1,59 @@
+extern "C" __global__ void {{kernel_name}}
+({%for arg in args%}{{arg.type}} {{'*' if arg.is_ptr}}{{arg.name}}, {% endfor %}
+  {{index_type}} *row_offsets,
+  {{index_type}} *eids,
+  {{index_type}} *column_indices,
+  {{index_type}} num_nodes,
+  {{index_type}} max_dimx,
+  {{index_type}} max_dimy,
+  {{index_type}} thrs_per_group,
+  {{index_type}} nodes_per_block) {
+      
+    {{index_type}} {{row_offset}} = nodes_per_block*blockIdx.x + threadIdx.x/thrs_per_group;
+    
+    if ({{row_offset}} < num_nodes) {
+        
+        {{index_type}} feat_len = max_dimx * max_dimy;
+        {{index_type}} beg = __ldg(row_offsets + {{row_offset}});
+        {{index_type}} end = __ldg(row_offsets + {{row_offset}} + 1);
+        {{index_type}} tx = threadIdx.x % thrs_per_group;
+        
+        for (; tx<feat_len; tx+=blockDim.x) {
+            
+            {%for agg_stmt in aggs%}{{agg_stmt.init}}{%endfor%}
+            {{init_outter_offset}}
+            
+            for ({{index_type}} e=beg;e<end;++e) {
+                
+                {{index_type}} {{col_index}} = __ldg(column_indices + e);
+
+                // PCSR indexes edges starting from 1
+                // Seastar requires edgs to be indexed from 0
+                {{index_type}} eid = __ldg(eids + e) - 1;
+                
+                {{init_inner_offset}}
+                
+                {%for edge_stmt in edges%}
+                {{edge_stmt.load}}
+                {{edge_stmt.compute}}
+                {{edge_stmt.inner_write}}
+                {%endfor%}
+                
+                {%for agg_stmt in aggs%}
+                {{agg_stmt.compute}}
+                {{agg_stmt.inner_write}}
+                {%endfor%}
+            }
+            
+            {%for agg_stmt in aggs%}
+            {{agg_stmt.outter_write}}
+            {%endfor%}
+            
+            {%for node_stmt in nodes%}
+            {{node_stmt.load}}
+            {{node_stmt.compute}}
+            {{node_stmt.inner_write}}
+            {%endfor%}
+        }
+    }
+}
\ No newline at end of file
diff --git a/seastar/compiler/code_gen/templates/tpl_header.jinja b/seastar/compiler/code_gen/templates/tpl_header.jinja
new file mode 100644
index 0000000..b4c3f8d
--- /dev/null
+++ b/seastar/compiler/code_gen/templates/tpl_header.jinja
@@ -0,0 +1,3 @@
+{%for header in headers%}
+{{'#include ' + header}}
+{%endfor%}
\ No newline at end of file
diff --git a/seastar/compiler/code_gen/templates/v2/__init__.py b/seastar/compiler/code_gen/templates/v2/__init__.py
new file mode 100644
index 0000000..e69de29
diff --git a/seastar/compiler/debugging/__init__.py b/seastar/compiler/debugging/__init__.py
new file mode 100644
index 0000000..e69de29
diff --git a/python/seastar/debugging/pretty_printers.py b/seastar/compiler/debugging/pretty_printers.py
similarity index 100%
rename from python/seastar/debugging/pretty_printers.py
rename to seastar/compiler/debugging/pretty_printers.py
diff --git a/python/seastar/execution_unit.py b/seastar/compiler/execution_unit.py
similarity index 86%
rename from python/seastar/execution_unit.py
rename to seastar/compiler/execution_unit.py
index c34ebae..77a107d 100644
--- a/python/seastar/execution_unit.py
+++ b/seastar/compiler/execution_unit.py
@@ -235,15 +235,17 @@ def add_ret_val(self, ret_val):
     def max_ret_id(self):
         return sorted([ret.int_id for ret in self.unit_rets()])[-1]
 
-    def prepare_compiled_kernel(self, graph_info, compiled_module):
+    def prepare_compiled_kernel(self, graph, compiled_module):
         if self.parallel_mode() == ParallelMode.DstParallel:
-            row_offsets = graph_info.in_row_offsets.data
-            col_indices = graph_info.in_col_indices.data
-            eids = graph_info.in_eids.data
+            row_offsets_ptr = graph.fwd_row_offset_ptr
+            col_indices_ptr = graph.fwd_column_indices_ptr
+            eids_ptr = graph.fwd_eids_ptr
         else:
-            row_offsets = graph_info.out_row_offsets.data
-            col_indices = graph_info.out_col_indices.data
-            eids = graph_info.out_eids.data
+            #TODO: Will probably have to change this so that this accesses 
+            #backward row_offset, col_indices, eids
+            row_offsets_ptr = graph.bwd_row_offset_ptr
+            col_indices_ptr = graph.bwd_column_indices_ptr
+            eids_ptr = graph.bwd_eids_ptr
         max_dims = [1, 1]
         if len(self.max_dims()) == 1:
             max_dims[-1] = self.max_dims()[-1]
@@ -251,26 +253,28 @@ def prepare_compiled_kernel(self, graph_info, compiled_module):
             max_dims = self.max_dims()
         else:
             raise NotImplementedError('Feature dimension larger than 2 are not supported.')
-        num_nodes = graph_info.number_of_nodes
+        num_nodes = graph.get_num_nodes()
         if self.use_fa_tmpl():
             launch_config = self.calculate_kernel_params_fa(num_nodes)
             print('template name:', self._template_name, 'number of nodes:', num_nodes, 'launch_config', launch_config)
-            self._K = FeatureAdaptiveKernel(num_nodes, row_offsets, col_indices, eids, max_dims, self._kernel_name, compiled_module, launch_config)
+            self._K = FeatureAdaptiveKernel(num_nodes, row_offsets_ptr, col_indices_ptr, eids_ptr, max_dims, self._kernel_name, compiled_module, launch_config)
         else:
             launch_config, tile_sizes = self.calculate_kernel_params(num_nodes)
             print('template name:', self._template_name, 'number of nodes:', num_nodes, 'launch_config', launch_config, 'tile_sizes', tile_sizes, 'max_dims', self.max_dims())
-            self._K = V2Kernel(num_nodes, row_offsets, col_indices, eids, max_dims, self._kernel_name, compiled_module, launch_config, tile_sizes)
+            self._K = V2Kernel(num_nodes, row_offsets_ptr, col_indices_ptr, eids_ptr, max_dims, self._kernel_name, compiled_module, launch_config, tile_sizes)
 
-    def reset_graph_info(self, graph_info):
+    def reset_graph_info(self, graph):
         if self.parallel_mode() == ParallelMode.DstParallel:
-            row_offsets = graph_info.in_row_offsets.data
-            col_indices = graph_info.in_col_indices.data
-            eids = graph_info.in_eids.data
+            row_offsets_ptr = graph.fwd_row_offset_ptr
+            col_indices_ptr = graph.fwd_column_indices_ptr
+            eids_ptr = graph.fwd_eids_ptr
         else:
-            row_offsets = graph_info.out_row_offsets.data
-            col_indices = graph_info.out_col_indices.data
-            eids = graph_info.out_eids.data
-        self._K.reset_graph_info(graph_info.number_of_nodes, row_offsets, col_indices, eids)
+            #TODO: Will probably have to change this so that this accesses 
+            #backward row_offset, col_indices, eids
+            row_offsets_ptr = graph.bwd_row_offset_ptr
+            col_indices_ptr = graph.bwd_column_indices_ptr
+            eids_ptr = graph.bwd_eids_ptr
+        self._K.reset_graph_info(graph.get_num_nodes(), row_offsets_ptr, col_indices_ptr, eids_ptr)
 
     def kernel_run(self, tensor_list):
         assert self._K, 'Must call prepare_compiled_kernel before call kernel_run.'
@@ -337,11 +341,11 @@ def kernel_name(self):
         return self._kernel_name
 
 class Kernel():
-    def reset_graph_info(self, number_of_nodes, row_offsets, col_indices, eids):
-        self.const_kernel_args[0] = c_void_p(row_offsets.data_ptr())
-        self.const_kernel_args[1] = c_void_p(eids.data_ptr())
-        self.const_kernel_args[2] = c_void_p(col_indices.data_ptr())
-        self.const_kernel_args[3] = c_int(number_of_nodes)
+    def reset_graph_info(self, num_nodes, row_offsets_ptr, col_indices_ptr, eids_ptr):
+        self.const_kernel_args[0] = c_void_p(row_offsets_ptr)
+        self.const_kernel_args[1] = c_void_p(eids_ptr)
+        self.const_kernel_args[2] = c_void_p(col_indices_ptr)
+        self.const_kernel_args[3] = c_int(num_nodes)
 
         # for row offsets, eids, col_indices
         for i in range(4):
@@ -366,9 +370,9 @@ def run(self, tensor_list):
             raise e
 
 class V2Kernel(Kernel):
-    def __init__(self, num_nodes, row_offsets, col_indices, eids, max_dims, kernel_name, compiled_module, launch_config, tile_sizes):
+    def __init__(self, num_nodes, row_offsets_ptr, col_indices_ptr, eids_ptr, max_dims, kernel_name, compiled_module, launch_config, tile_sizes):
         self.scalar_args = [c_int(num_nodes), c_int(max_dims[1]), c_int(max_dims[0]), c_int(tile_sizes[0]), c_int(tile_sizes[1])]
-        self.const_kernel_args =  [c_void_p(row_offsets.data_ptr()), c_void_p(eids.data_ptr()), c_void_p(col_indices.data_ptr())] + self.scalar_args
+        self.const_kernel_args =  [c_void_p(row_offsets_ptr), c_void_p(eids_ptr), c_void_p(col_indices_ptr)] + self.scalar_args
         self.const_kernel_ptrs = [c_void_p(addressof(v)) for v in self.const_kernel_args]
         self.K = c_void_p(0)
         ret = cuModuleGetFunction(byref(self.K), compiled_module, c_char_p(kernel_name.encode()))
@@ -377,9 +381,9 @@ def __init__(self, num_nodes, row_offsets, col_indices, eids, max_dims, kernel_n
         self.launch_config = launch_config[0],launch_config[1], 1, launch_config[2], launch_config[3],1
 
 class FeatureAdaptiveKernel(Kernel):
-    def __init__(self, num_nodes, row_offsets, col_indices, eids, max_dims, kernel_name, compiled_module, launch_config):
+    def __init__(self, num_nodes, row_offsets_ptr, col_indices_ptr, eids_ptr, max_dims, kernel_name, compiled_module, launch_config):
         self.scalar_args = [c_int(num_nodes), c_int(max_dims[1]), c_int(max_dims[0]), c_int(launch_config[2]), c_int(launch_config[3])]
-        self.const_kernel_args =  [c_void_p(row_offsets.data_ptr()), c_void_p(eids.data_ptr()), c_void_p(col_indices.data_ptr())] + self.scalar_args
+        self.const_kernel_args =  [c_void_p(row_offsets_ptr), c_void_p(eids_ptr), c_void_p(col_indices_ptr)] + self.scalar_args
         self.const_kernel_ptrs = [c_void_p(addressof(v)) for v in self.const_kernel_args]
 
         self.K = c_void_p(0)
diff --git a/python/seastar/executor.py b/seastar/compiler/executor.py
similarity index 85%
rename from python/seastar/executor.py
rename to seastar/compiler/executor.py
index 51f6007..ea76e7a 100644
--- a/python/seastar/executor.py
+++ b/seastar/compiler/executor.py
@@ -1,7 +1,7 @@
-import torch
 from .utils import is_const_scalar, ParallelMode
 import snoop
 from collections import deque
+from ..graph.dynamic.DynamicGraph import DynamicGraph
 
 class Stack:
     def __init__(self, val=None):
@@ -28,6 +28,9 @@ def __init__(self):
 
         # contains tensors for all previous execution of nb_compute
         self.tensor_map_stack = Stack()
+        
+        # contains timestamps of graphs that were forward propagated
+        self.graph_timestamp_stack = Stack()
 
         # contains arg tensors for the current execution of nb_compute only
         self.current_tensor_map = {}
@@ -155,7 +158,7 @@ def  __iter__(self):
             yield unit
 
 class Executor(object):
-    def __init__(self, graph_info, forward_exec_units, backward_exec_units, compiled_module, rets):
+    def __init__(self, graph, forward_exec_units, backward_exec_units, compiled_module, rets):
         self.forward_exec_units = self.merge_units(forward_exec_units)
         self.bulist = backward_exec_units
         self.var2bu = self.construct_backward_mappping(self.forward_exec_units,backward_exec_units)
@@ -163,15 +166,16 @@ def __init__(self, graph_info, forward_exec_units, backward_exec_units, compiled
         self.ts = ExeState()
         self.new_zeros = None
         self.raw_ptr = None
-        self.num_nodes = graph_info.number_of_nodes
-        self.num_edges = graph_info.number_of_edges
+        self.num_nodes = graph.get_num_nodes()
+        self.num_edges = graph.get_num_edges()
+        self.graph = graph
         for mu in self.forward_exec_units:
             for u in mu:
                 if u.compiled:
-                    u.prepare_compiled_kernel(graph_info, compiled_module)
+                    u.prepare_compiled_kernel(graph, compiled_module)
         for u in self.bulist:
             if u.compiled:
-                u.prepare_compiled_kernel(graph_info, compiled_module)
+                u.prepare_compiled_kernel(graph, compiled_module)
     
     def construct_backward_mappping(self, funits, bunits):
         ret = {}
@@ -198,19 +202,29 @@ def merge_units(self, exec_units):
         print('merged units', len(grouped_unit), grouped_unit)
         return grouped_unit
   
-    def restart(self, input_map, graph_info=None):
+    def restart(self, input_map, graph=None):
         self.ts.reset(input_map, self.forward_exec_units, self.bulist)
-        if graph_info != None:
+        if graph != None:
+            
+            # TODO: REMOVE
+            # TODO: getting graph of current timestamp, probably better to move
+            # this outside the compiler
+            # current_timestamp = self.ts.tensor_map_stack.len() 
+            # self.graph.get_forward_graph_for_timestamp(current_timestamp)
+            
             for mu in self.forward_exec_units:
                 for u in mu:
                     if u.compiled:
                         # TODO: (Joel) Feel like this is going to be problematic for dynamic graphs
-                        u.reset_graph_info(graph_info)
-            for u in self.bulist:
-                if u.compiled:
-                    u.reset_graph_info(graph_info)
-            self.num_nodes = graph_info.number_of_nodes
-            self.num_edges = graph_info.number_of_edges
+                        u.reset_graph_info(graph)
+            
+            # NOTE: COMMENTED OUT NOW SINCE THIS IS HANDLED IN BACKWARD_CB
+            # for u in self.bulist:
+            #     if u.compiled:
+            #         u.reset_graph_info(graph)
+            
+            self.num_nodes = graph.get_num_nodes()
+            self.num_edges = graph.get_num_edges()
 
     
     def set_raw_ptr_cb(self, cb):
@@ -235,7 +249,9 @@ def execute(self, FuncWrapper):
         # bytes_list = [v.numel() *4 for k,v in self.ts.tensor_map.items()]
         #print('after forward', self.ts.tensor_map.keys(), ' bytes ', bytes_list, sum(bytes_list))
 
-        self.ts.tensor_map_stack.push(self.ts.current_tensor_map)
+        # Old position
+        # self.ts.tensor_map_stack.push(self.ts.current_tensor_map)
+        # self.ts.graph_timestamp_stack.push(self.graph.current_timestamp)
 
         # print("🔴 After ForwardProp status of tensor_map")
         # for index in range(len(self.ts.tensor_map_stack.content)):
@@ -284,6 +300,11 @@ def forward_cb(self, uid, kernel_args, rets, tensor_list):
         units = self.forward_exec_units[uid]
         for i,unit in enumerate(units):
             self.execute_unit(unit, [tensor_list[tidx] for tidx in kernel_args[i]])
+            
+        self.ts.tensor_map_stack.push(self.ts.current_tensor_map)
+        
+        if isinstance(self.graph,DynamicGraph):
+            self.ts.graph_timestamp_stack.push(self.graph.current_timestamp)
 
         return tuple([self.ts.current_tensor_map[ret.id] for ret in rets])
 
@@ -298,15 +319,24 @@ def backward_cb(self, kid, grad_list):
         inputs = funits.joint_inputs()
         ret_grads = [ret._grad for ret in rets] # ret_grads corresponds vars in grad_list
         tensor_map = self.ts.tensor_map_stack.top()
+        
+        if isinstance(self.graph,DynamicGraph):
+            current_timestamp = self.ts.graph_timestamp_stack.top()
+            self.graph.get_backward_graph(current_timestamp)
 
         for i,grad in enumerate(ret_grads):
             # We track the ret_grads as its value is fixed to grad_list
             tensor_map[grad.id] = grad_list[i]
         arg_grads = [arg._grad if arg in inputs and arg.requires_grad else None for arg in args] # arg_grads corresponds to the grads of funit.unit_args
         for bu in self.bulist:
+
             if bu.compiled:
                 # if self.ts.is_executed_bu(bu):
                 #     continue
+                
+                # NOTE: Added to fix gpma
+                bu.reset_graph_info(self.graph)
+                
                 tensor_map = self.create_tensor_for_grad_vars(bu.unit_rets(),tensor_map)
 
                 self.execute_unit(bu, [tensor_map[arg.id] for arg in bu.kernel_args()])
@@ -320,11 +350,10 @@ def backward_cb(self, kid, grad_list):
         
         del tensor_map
         self.ts.tensor_map_stack.pop()
-
-        # print("🔴 After Backprop status of tensor_map")
-        # for index in range(len(self.ts.tensor_map_stack.content)):
-        #     print("Index: {}".format(index))
-        #     print(self.ts.tensor_map_stack.content[index])
+        
+        if isinstance(self.graph,DynamicGraph):
+            self.ts.graph_timestamp_stack.pop()
+        
         return ret
 
     def execute_prog(self, units):
diff --git a/python/seastar/node.py b/seastar/compiler/node.py
similarity index 100%
rename from python/seastar/node.py
rename to seastar/compiler/node.py
diff --git a/python/seastar/op.py b/seastar/compiler/op.py
similarity index 100%
rename from python/seastar/op.py
rename to seastar/compiler/op.py
diff --git a/python/seastar/passes/__init__.py b/seastar/compiler/passes/__init__.py
similarity index 100%
rename from python/seastar/passes/__init__.py
rename to seastar/compiler/passes/__init__.py
diff --git a/python/seastar/passes/cf.py b/seastar/compiler/passes/cf.py
similarity index 100%
rename from python/seastar/passes/cf.py
rename to seastar/compiler/passes/cf.py
diff --git a/python/seastar/passes/cse.py b/seastar/compiler/passes/cse.py
similarity index 100%
rename from python/seastar/passes/cse.py
rename to seastar/compiler/passes/cse.py
diff --git a/python/seastar/passes/dce.py b/seastar/compiler/passes/dce.py
similarity index 100%
rename from python/seastar/passes/dce.py
rename to seastar/compiler/passes/dce.py
diff --git a/python/seastar/passes/dependency_analysis.py b/seastar/compiler/passes/dependency_analysis.py
similarity index 100%
rename from python/seastar/passes/dependency_analysis.py
rename to seastar/compiler/passes/dependency_analysis.py
diff --git a/python/seastar/passes/fusion.py b/seastar/compiler/passes/fusion.py
similarity index 100%
rename from python/seastar/passes/fusion.py
rename to seastar/compiler/passes/fusion.py
diff --git a/python/seastar/passes/mem_planning.py b/seastar/compiler/passes/mem_planning.py
similarity index 100%
rename from python/seastar/passes/mem_planning.py
rename to seastar/compiler/passes/mem_planning.py
diff --git a/python/seastar/passes/peephole.py b/seastar/compiler/passes/peephole.py
similarity index 100%
rename from python/seastar/passes/peephole.py
rename to seastar/compiler/passes/peephole.py
diff --git a/python/seastar/passes/visualize.py b/seastar/compiler/passes/visualize.py
similarity index 100%
rename from python/seastar/passes/visualize.py
rename to seastar/compiler/passes/visualize.py
diff --git a/python/seastar/program.py b/seastar/compiler/program.py
similarity index 100%
rename from python/seastar/program.py
rename to seastar/compiler/program.py
diff --git a/python/seastar/registry.py b/seastar/compiler/registry.py
similarity index 100%
rename from python/seastar/registry.py
rename to seastar/compiler/registry.py
diff --git a/python/seastar/schema.py b/seastar/compiler/schema.py
similarity index 100%
rename from python/seastar/schema.py
rename to seastar/compiler/schema.py
diff --git a/python/seastar/zoom_v2.py b/seastar/compiler/seastar.py
similarity index 81%
rename from python/seastar/zoom_v2.py
rename to seastar/compiler/seastar.py
index 2316b86..702ea3b 100644
--- a/python/seastar/zoom_v2.py
+++ b/seastar/compiler/seastar.py
@@ -15,19 +15,8 @@
 from .debugging.pretty_printers import pretty_print_GIR, pretty_print_Central_Node
 import snoop
 
-# REMOVE THIS
-import torch
 
 class Context():
-    GraphInfo = namedtuple('GraphInfo', ['number_of_nodes',
-                                         'number_of_edges',
-                                         'in_row_offsets',
-                                         'in_col_indices',
-                                         'in_eids',
-                                         'out_row_offsets',
-                                         'out_col_indices',
-                                         'out_eids',
-                                         'nbits'])
     def __init__(self, func, nspace, run_cb):
         functools.update_wrapper(self, func)
         self._f = func
@@ -52,13 +41,12 @@ def _setup_executor(self, **kwargs):
         edge_feats = kwargs.get('e_feats', {})
         if not graph:
             raise NameError('Need to provide the graph as one of keyward arguments')
-        graph_info, need_reset = self._update_graph_info(graph)
         if self._entry_count == 0:
             fprog = Program()
             ret = self._trace(node_feats, edge_feats, self._input_cache, fprog)
             # print('TracedProgram' + str(fprog), 'Ret value:', ret)
             # pretty_print_GIR(fprog,"TGCN GIR")
-            self._executor_cache = self._diff_then_compile(ret, fprog, graph_info)
+            self._executor_cache = self._diff_then_compile(ret, fprog, graph)
         for k, v in node_feats.items():
             self._input_cache[var_prefix + k + cen_attr_postfix] = v
             self._input_cache[var_prefix + k + inb_attr_postfix] = v
@@ -66,31 +54,10 @@ def _setup_executor(self, **kwargs):
             self._input_cache[var_prefix+k] = v
         # print("🔴 Input Cache")
         # print(self._input_cache)
-        self._executor_cache.restart(self._input_cache, graph_info if need_reset else None)
+        self._executor_cache.restart(self._input_cache, graph)
         self._entry_count += 1
         return self._executor_cache
 
-    def _update_graph_info(self, graph):
-        reset = False
-        if not (self._graph_info_cache != None
-                and self._graph_info_cache.number_of_nodes == graph.number_of_nodes()
-                and self._graph_info_cache.number_of_edges == graph.number_of_edges()):
-            out_csr = graph.adj_sparse(fmt="csr")
-            in_csr = graph.reverse().adj_sparse(fmt="csr")
-
-            self._graph_info_cache = Context.GraphInfo(
-                graph.number_of_nodes(),
-                graph.number_of_edges(),
-                in_csr[0].type(torch.int32),
-                in_csr[1].type(torch.int32),
-                in_csr[2].type(torch.int32),
-                out_csr[0].type(torch.int32),
-                out_csr[1].type(torch.int32),
-                out_csr[2].type(torch.int32),
-                graph._graph.bits_needed(0)) #TODO: Setting this to graph._graph.bits_needed(0), since function not there in DGL and the number plays no role in seastar, but this wont work for homogenous
-            reset = True
-        return self._graph_info_cache, reset
-
     def _trace(self, nfeats, efeats, input_cache, fprog):
         backend = self.find_backend(self._nspace)
         central_node = self._init_central_node(nfeats, efeats, fprog, backend)
@@ -113,8 +80,11 @@ def _diff_then_compile(self, out_set, fprog, graph):
         for var in vars:
             grads.append(Var.create_var(var_shape=var.var_shape, var_dtype=var.var_dtype, val_type=var.val_type, device=var.device))
         backward_exe_units = diff(vars, grads, forward_exe_units, fprog)
-        visualize.plot_exec_units(forward_exe_units + backward_exe_units)
-        compiled_module = code_gen.gen_code(forward_exe_units + backward_exe_units, 'int' if graph.nbits == 32 else 'long long int')
+        # visualize.plot_exec_units(forward_exe_units + backward_exe_units)
+        
+        # NOTE: The last parameter here was ('int' if graph.nbits == 32 else 'long long int') but we changed
+        # it to just 'int' since that should be sufficient for all use case that we can think of now
+        compiled_module = code_gen.gen_code(forward_exe_units + backward_exe_units, 'int', graph.graph_type())
         return Executor(graph, forward_exe_units, backward_exe_units, compiled_module, vars)
         
     def _init_central_node(self, nfeats, efeats, fprog, backend):
@@ -216,12 +186,12 @@ def _mapping_key(self, name_space_id, original_key):
         return str(name_space_id) + str(original_key)
 
 
-class CtxManager():
+class Seastar():
     def __init__(self, run_cb):
         self._ctx_map = {}
         self._run_cb = run_cb
     
-    def zoomIn(self, nspace, hetero_graph=False):
+    def compile(self, nspace, hetero_graph=False):
         def wrapper(func):
             if not func.__name__ in self._ctx_map:
                 if not hetero_graph:
diff --git a/python/seastar/utils.py b/seastar/compiler/utils.py
similarity index 100%
rename from python/seastar/utils.py
rename to seastar/compiler/utils.py
diff --git a/python/seastar/val.py b/seastar/compiler/val.py
similarity index 100%
rename from python/seastar/val.py
rename to seastar/compiler/val.py
diff --git a/seastar/dataset/BitcoinOTC.py b/seastar/dataset/BitcoinOTC.py
new file mode 100644
index 0000000..6641cc0
--- /dev/null
+++ b/seastar/dataset/BitcoinOTC.py
@@ -0,0 +1,43 @@
+class BitcoinOtcDataloader():
+    def __init__(self):
+        self._dataset = pd.read_csv('../../dataset/bitcoin/soc-sign-bitcoinotc.csv',header=0, names=['source', 'target', 'rating','time'])
+        self._dataset.time = self._dataset.time / 10000000
+        self._dataset.time = self._dataset.time.astype(int)
+
+        lst = np.append(self._dataset.source.astype(int).unique(),self._dataset.target.astype(int).unique())
+        lst = np.sort(np.unique(lst))
+        
+        id_dict = {}
+        for i in range(len(lst)):
+            id_dict[lst[i]] = i
+        self.node_id_dict = id_dict
+            
+        self._dataset.source = self._dataset.source.apply(lambda x: self.node_id_dict[x])
+        self._dataset.target = self._dataset.target.apply(lambda x: self.node_id_dict[x])
+    
+    def _get_edges(self):
+        self._edges = []
+        time_lst = self._dataset.time.unique()
+        time_lst.sort()
+        for time in time_lst:
+            filtered_df = self._dataset.where(self._dataset.time == time).dropna()
+            self._edges.append(
+                np.array([filtered_df.source.astype(int).tolist(),filtered_df.target.astype(int).tolist()])
+            )
+
+    def _get_edge_weights(self):
+        self._edge_weights = []
+        for time in self._dataset.time.unique():
+            filtered_df = self._dataset.where(self._dataset.time == time).dropna()
+            self._edge_weights.append(
+                np.array(filtered_df.rating.tolist())
+            )
+
+    def _get_targets_and_features(self):
+        pass
+
+    def get_dataset(self):
+        self._get_edges()
+        self._get_edge_weights()
+        self._get_targets_and_features()
+        return self._edges, self._edge_weights, None, None
\ No newline at end of file
diff --git a/seastar/dataset/EnglandCOVID.py b/seastar/dataset/EnglandCOVID.py
new file mode 100644
index 0000000..2bc8d18
--- /dev/null
+++ b/seastar/dataset/EnglandCOVID.py
@@ -0,0 +1,127 @@
+import os
+import json
+import urllib.request
+import time
+
+import numpy as np
+
+from rich import inspect
+from rich.pretty import pprint
+from rich.progress import track
+from rich.console import Console
+
+console = Console()
+
+from rich.traceback import install
+install(show_locals=True)
+
+class EnglandCOVID:
+    def __init__(self, verbose: bool = False, lags: int = 8, split=0.75) -> None:
+        self.name = "EnglandCOVID"
+        self.lags = lags
+        self.split = split
+
+        self._graph_attr = {}
+        self._graph_updates = {}
+        self._max_num_nodes = 0
+
+        self._url_path = "https://raw.githubusercontent.com/benedekrozemberczki/pytorch_geometric_temporal/master/dataset/england_covid.json"
+        self._verbose = verbose
+
+        self._load_dataset()
+        self.total_timestamps = self._dataset["time_periods"]
+        self._get_edge_info()
+        self._get_targets_and_features()
+        self._presort_edge_weights()
+
+    def get_graph_data(self):
+        return self._graph_updates, self._max_num_nodes
+
+    def _load_dataset(self) -> None:
+        # loading the dataset by downloading them online
+        if self._verbose:
+            console.log(f"Downloading [cyan]{self.name}[/cyan] dataset")
+        self._dataset = json.loads(urllib.request.urlopen(self._url_path).read())
+
+    def _get_edge_info(self):
+        # getting the edge_list and edge_weights
+        self._edge_list = []
+        self._edge_weights = []
+
+        for time in range(self._dataset["time_periods"] - self.lags):
+            time_edge_list = []
+            time_edge_weights = []
+
+            for edge in self._dataset["edge_mapping"]["edge_index"][str(time)]:
+                time_edge_list.append((edge[0], edge[1]))
+
+            for weight in self._dataset["edge_mapping"]["edge_weight"][str(time)]:
+                time_edge_weights.append(weight)
+
+            self._edge_list.append(time_edge_list)
+            self._edge_weights.append(time_edge_weights)
+
+    def _get_targets_and_features(self):
+        stacked_target = np.array(self._dataset["y"])
+        standardized_target = (stacked_target - np.mean(stacked_target, axis=0)) / (
+            np.std(stacked_target, axis=0) + 10**-10
+        )
+
+        self._all_features = [
+            standardized_target[i : i + self.lags, :].T
+            for i in range(self._dataset["time_periods"] - self.lags)
+        ]
+        self._all_targets = [
+            standardized_target[i + self.lags, :].T
+            for i in range(self._dataset["time_periods"] - self.lags)
+        ]
+
+    def _presort_edge_weights(self):
+        """
+        Presorting edges according to (dest,src) since that is how eids are formed
+        allowing forward and backward kernel to access edge weights
+        """
+        final_edges_lst = []
+        final_edge_weights_lst = []
+
+        for i in range(len(self._edge_list)):
+            src_list = [edge[0] for edge in self._edge_list[i]]
+            dst_list = [edge[1] for edge in self._edge_list[i]]
+            weights = self._edge_weights[i]
+
+            edge_info_list = []
+            sorted_edges_lst = []
+            sorted_edge_weights_lst = []
+
+            for j in range(len(weights)):
+                edge_info = (src_list[j], dst_list[j], weights[j])
+                edge_info_list.append(edge_info)
+
+            # since it has to be sorted according to the reverse order
+            sorted_edge_info_list = sorted(
+                edge_info_list, key=lambda element: (element[1], element[0])
+            )
+
+            time_edge = []
+
+            for edge in sorted_edge_info_list:
+                time_edge.append((edge[0], edge[1]))
+                sorted_edge_weights_lst.append(edge[2])
+
+            final_edges_lst.append(time_edge)
+            final_edge_weights_lst.append(np.array(sorted_edge_weights_lst))
+
+        self._edge_list = final_edges_lst
+        self._edge_weights = final_edge_weights_lst
+
+    def get_edges(self):
+        return self._edge_list
+
+    def get_edge_weights(self):
+        return self._edge_weights
+
+    def get_all_features(self):
+        return self._all_features
+
+    def get_all_targets(self):
+        return self._all_targets
diff --git a/seastar/dataset/__init__.py b/seastar/dataset/__init__.py
new file mode 100644
index 0000000..7c40d61
--- /dev/null
+++ b/seastar/dataset/__init__.py
@@ -0,0 +1,1 @@
+'''Dataset loader provided by Seastar'''
\ No newline at end of file
diff --git a/seastar/dataset/cora.py b/seastar/dataset/cora.py
new file mode 100644
index 0000000..085fa52
--- /dev/null
+++ b/seastar/dataset/cora.py
@@ -0,0 +1,112 @@
+import os
+import json
+import urllib.request
+import time
+import random
+
+import numpy as np
+
+from rich import inspect
+from rich.pretty import pprint
+from rich.progress import track
+from rich.console import Console
+
+console = Console()
+
+class CoraDataset:
+    def __init__(self, verbose:bool = False, split=0.75) -> None:
+        self.name = "Cora"
+        self.num_nodes = 0
+        self.num_edges = 0
+        
+        self._train_split = split
+        self._test_split = 1-split
+
+        self._url_path = "https://raw.githubusercontent.com/bfGraph/Seastar-Datasets/main/cora.json"
+        self._verbose = verbose
+
+        self._load_dataset()
+        self._get_edge_info()
+        self._get_targets_and_features()
+        self._get_graph_attributes()
+        
+        self._train_mask = [0] * self.num_nodes
+        self._test_mask = [0] * self.num_nodes
+        
+        self._get_mask_info()
+
+    def _load_dataset(self) -> None:
+        
+        # loading the dataset by downloading them online
+        if self._verbose:
+            console.log(f'Downloading [cyan]{self.name}[/cyan] dataset')
+        self._dataset = json.loads(urllib.request.urlopen(self._url_path).read())
+
+
+    def _get_edge_info(self):
+        edges = np.array(self._dataset["edges"])
+        edge_list = []
+        for i in range(len(edges)):
+            edge = edges[i]
+            edge_list.append((edge[0], edge[1]))
+            
+        self._edge_list = edge_list
+
+    def _get_targets_and_features(self):
+        self._all_features = np.array(self._dataset["features"])
+        self._all_targets = np.array(self._dataset["labels"]).T
+
+    def _get_mask_info(self):
+        train_len = int(self.num_nodes * self._train_split)
+        
+        for i in range(0, train_len):
+            self._train_mask[i] = 1
+            
+        random.shuffle(self._train_mask)
+            
+        for i in range(len(self._train_mask)):
+            if self._train_mask[i] == 0:
+                self._test_mask[i] = 1
+            
+        self._train_mask = np.array(self._train_mask)
+        self._test_mask = np.array(self._test_mask)
+
+    def get_edges(self) -> np.ndarray:
+        return self._edge_list
+
+    def get_all_features(self) -> np.ndarray:
+        return self._all_features
+
+    def get_all_targets(self) -> np.ndarray:
+        return self._all_targets
+
+    def get_train_mask(self):
+        return self._train_mask
+    
+    def get_test_mask(self):
+        return self._test_mask
+
+    def get_train_features(self) -> np.ndarray:
+        train_range = int(len(self._all_features) * self.split)
+        return self._all_features[:train_range]
+
+    def get_train_targets(self) -> np.ndarray:
+        train_range = int(len(self._all_targets) * self.split)
+        return self._all_targets[:train_range]
+
+    def get_test_features(self) -> np.ndarray:
+        test_range = int(len(self._all_features) * self.split)
+        return self._all_features[test_range:]
+
+    def get_test_targets(self) -> np.ndarray:
+        test_range = int(len(self._all_targets) * self.split)
+        return self._all_targets[test_range:]
+
+    def _get_graph_attributes(self):
+        node_set = set()
+        for edge in self._edge_list:
+            node_set.add(edge[0])
+            node_set.add(edge[1])
+            
+        self.num_nodes = len(node_set)
+        self.num_edges = len(self._edge_list)
\ No newline at end of file
diff --git a/seastar/graph/SeastarGraph.py b/seastar/graph/SeastarGraph.py
new file mode 100644
index 0000000..5edd24e
--- /dev/null
+++ b/seastar/graph/SeastarGraph.py
@@ -0,0 +1,33 @@
+from abc import ABC, abstractmethod
+
+class SeastarGraph(ABC):
+    def __init__(self):
+        self.ndata = {}
+        
+        self._forward_graph = None
+        self._backward_graph = None
+        
+        self.fwd_row_offset_ptr = None
+        self.fwd_column_indices_ptr = None
+        self.fwd_eids_ptr = None
+        
+        self.bwd_row_offset_ptr = None
+        self.bwd_column_indices_ptr = None
+        self.bwd_eids_ptr = None
+    
+    @abstractmethod
+    def _get_graph_csr_ptrs(self):
+        pass
+    
+    @abstractmethod
+    def get_num_nodes(self):
+        pass
+    
+    @abstractmethod
+    def get_num_edges(self):
+        pass
+    
+    @property
+    @abstractmethod
+    def graph_type(self):
+        pass
\ No newline at end of file
diff --git a/seastar/graph/__init__.py b/seastar/graph/__init__.py
new file mode 100644
index 0000000..6a3e99e
--- /dev/null
+++ b/seastar/graph/__init__.py
@@ -0,0 +1,1 @@
+'''Graph Abstraction provided by Seastar'''
\ No newline at end of file
diff --git a/seastar/graph/dynamic/DynamicGraph.py b/seastar/graph/dynamic/DynamicGraph.py
new file mode 100644
index 0000000..04bdd17
--- /dev/null
+++ b/seastar/graph/dynamic/DynamicGraph.py
@@ -0,0 +1,132 @@
+# TODO: add this import as well
+from seastar.graph.SeastarGraph import SeastarGraph
+
+import copy
+
+from abc import ABC, abstractmethod
+
+class DynamicGraph(SeastarGraph):
+    def __init__(self, edge_list):
+        super().__init__()
+        self.graph_updates = {}
+        self.max_num_nodes = 0
+        self.graph_cache = {}
+        
+        # Indicates whether the graph is currently undergoing backprop
+        self._is_backprop_state = False
+        self.current_timestamp = 0
+        
+        self._preprocess_graph_structure(edge_list)
+        
+    def _get_graph_attr(self, edge_list):
+        graph_attr = {}
+        for time in range(len(edge_list)):
+            node_set = set()
+            edge_count = 0
+            for edge in edge_list[time]:
+                src = edge[0]
+                dst = edge[1]
+                edge_count += 1
+                node_set.add(src)
+                node_set.add(dst)
+            node_count = len(node_set)
+            graph_attr[str(time)] = (node_count, edge_count)
+        
+        return graph_attr
+        
+    def _preprocess_graph_structure(self, edge_list):
+        
+        graph_attr = self._get_graph_attr(edge_list)
+        
+        tmp_set = set()
+        for i in range(len(edge_list)):
+            tmp_set = set()
+            for j in range(len(edge_list[i])):
+                tmp_set.add(edge_list[i][j][0])
+                tmp_set.add(edge_list[i][j][1])
+        self.max_num_nodes = len(tmp_set)
+
+        edge_dict = {}
+        for i in range(len(edge_list)):
+            edge_set = set()
+            for j in range(len(edge_list[i])):
+                edge_set.add((edge_list[i][j][0], edge_list[i][j][1]))
+            edge_dict[str(i)] = edge_set
+
+        self.graph_updates = {}
+        self.graph_updates["0"] = {
+            "add": list(edge_dict["0"]),
+            "delete": [],
+            "num_nodes": graph_attr["0"][0],
+            "num_edges": graph_attr["0"][1],
+        }
+        for i in range(1, len(edge_list)):
+            self.graph_updates[str(i)] = {
+                "add": list(edge_dict[str(i)].difference(edge_dict[str(i - 1)])),
+                "delete": list(edge_dict[str(i - 1)].difference(edge_dict[str(i)])),
+                "num_nodes": graph_attr[str(i)][0],
+                "num_edges": graph_attr[str(i)][1],
+            }
+        
+    def get_graph(self, timestamp: int):
+
+        self._is_backprop_state = False
+        
+        if timestamp < self.current_timestamp:
+            raise Exception("⏰ Invalid timestamp during SeastarGraph.update_graph_forward()")
+
+        while self.current_timestamp < timestamp:
+            self._update_graph_forward()
+            self.current_timestamp += 1
+
+    def get_backward_graph(self, timestamp: int):
+
+        if not self._is_backprop_state:
+            self._init_reverse_graph()
+            self._is_backprop_state = True
+        
+        if timestamp > self.current_timestamp:
+            raise Exception("⏰ Invalid timestamp during SeastarGraph.update_graph_backward()")
+        
+        while self.current_timestamp > timestamp:
+            self._update_graph_backward()
+            self.current_timestamp -= 1
+    
+    def get_num_nodes(self):
+        return self.graph_updates[str(self.current_timestamp)]["num_nodes"]
+    
+    def get_num_edges(self):
+        return self.graph_updates[str(self.current_timestamp)]["num_edges"]
+    
+    # def _get_graph_csr_ptrs(self):
+    #     if not self._is_backprop_state:
+    #         csr_ptrs = self._forward_graph.get_csr_ptrs()
+    #     else:
+    #         csr_ptrs = self._backward_graph.get_csr_ptrs()
+        
+    #     self.row_offset_ptr = csr_ptrs[0]
+    #     self.column_indices_ptr = csr_ptrs[1]
+    #     self.eids_ptr = csr_ptrs[2]
+    
+    @abstractmethod
+    def in_degrees(self):
+        pass
+    
+    @abstractmethod
+    def out_degrees(self):
+        pass
+    
+    
+    @abstractmethod
+    def _update_graph_forward(self):
+        pass
+    
+    @abstractmethod
+    def _init_reverse_graph(self):
+        pass
+    
+    @abstractmethod
+    def _update_graph_backward(self):
+        pass
+    
+    
\ No newline at end of file
diff --git a/seastar/graph/dynamic/__init__.py b/seastar/graph/dynamic/__init__.py
new file mode 100644
index 0000000..e69de29
diff --git a/seastar/graph/dynamic/gpma/GPMAGraph.py b/seastar/graph/dynamic/gpma/GPMAGraph.py
new file mode 100644
index 0000000..4133437
--- /dev/null
+++ b/seastar/graph/dynamic/gpma/GPMAGraph.py
@@ -0,0 +1,121 @@
+import copy
+import numpy as np
+from rich import inspect
+
+from seastar.graph.dynamic.DynamicGraph import DynamicGraph
+from seastar.graph.dynamic.gpma.gpma import GPMA, init_gpma, print_gpma_info, edge_update_list, label_edges, copy_label_edges, build_reverse_gpma, get_csr_ptrs, get_graph_attr
+
+
+class GPMAGraph(DynamicGraph):
+    def __init__(self, edge_list):
+        super().__init__(edge_list)
+        
+        # forward and backward graphs for GPMA
+        self._forward_graph = GPMA()
+        self._backward_graph = GPMA()
+        
+        init_gpma(self._forward_graph, self.max_num_nodes)
+        init_gpma(self._backward_graph, self.max_num_nodes)
+        
+        # initialise the graph for the first time stamp
+        initial_graph_additions = self.graph_updates["0"]["add"]
+
+        edge_update_list(self._forward_graph, initial_graph_additions, is_reverse_edge=True)
+        label_edges(self._forward_graph)
+
+        self._get_graph_csr_ptrs()
+        
+    def graph_type(self):
+        return "gpma"
+        
+    def in_degrees(self):
+        return np.array(self._forward_graph.out_degree, dtype='int32')
+    
+    def out_degrees(self):
+        return np.array(self._forward_graph.in_degree, dtype='int32')
+    
+    def _update_reverse_graph_cache(self):
+        # saving reverse base graph in cache
+        self.graph_cache['reverse'] = copy.deepcopy(self._backward_graph)
+            
+    def _get_reverse_cached_graph(self):
+        return copy.deepcopy(self.graph_cache['reverse'])
+    
+    def _get_graph_csr_ptrs(self):
+        forward_csr_ptrs = get_csr_ptrs(self._forward_graph)
+        backward_csr_ptrs = get_csr_ptrs(self._backward_graph)
+        
+        self.fwd_row_offset_ptr = forward_csr_ptrs[0]
+        self.fwd_column_indices_ptr = forward_csr_ptrs[1]
+        self.fwd_eids_ptr = forward_csr_ptrs[2]
+        
+        self.bwd_row_offset_ptr = backward_csr_ptrs[0]
+        self.bwd_column_indices_ptr = backward_csr_ptrs[1]
+        self.bwd_eids_ptr = backward_csr_ptrs[2]
+    
+    #TODO: Right now this returns (max_num_nodes,num_edges) see if this is what is required
+    # def _get_graph_attributes(self):
+
+    #     if not self._is_backprop_state:
+    #         graph_attr = get_graph_attr(self._forward_graph)
+    #     else:
+    #         graph_attr = get_graph_attr(self._backward_graph)
+        
+    #     self.num_nodes = graph_attr[0]
+    #     self.num_edges = graph_attr[1]
+    
+    def _update_graph_forward(self):
+        
+        # if we went through the entire time-stamps
+        if str(self.current_timestamp + 1) not in self.graph_updates:
+            raise Exception("⏰ Invalid timestamp during SeastarGraph.update_graph_forward()")
+        
+        # getting the graph edge modifications in 
+        # the following form list[tuple(int, int)]
+        graph_additions = self.graph_updates[str(self.current_timestamp + 1)]["add"]
+        graph_deletions = self.graph_updates[str(self.current_timestamp + 1)]["delete"]
+
+        edge_update_list(self._forward_graph, graph_additions, is_reverse_edge=True)
+        edge_update_list(self._forward_graph, graph_deletions, is_delete=True, is_reverse_edge=True)
+
+        # TODO: UNCOMMENT LATER
+        label_edges(self._forward_graph)
+        self._get_graph_csr_ptrs()
+        # self._get_graph_attributes()  # NOTE:
+        
+    def _init_reverse_graph(self):
+        ''' Generates the reverse of the base graph'''
+
+        # checking if the reverse base graph exists in the cache
+        # we can load it from there instead of building it each time
+        if 'reverse' in self.graph_cache:
+            self._backward_graph = self._get_reverse_cached_graph()
+        else:
+            build_reverse_gpma(self._backward_graph, self._forward_graph)
+
+            # storing the reverse base graph in cache after building
+            # it for the first time
+            self._update_reverse_graph_cache()
+
+        self._get_graph_csr_ptrs()
+        # self._get_graph_attributes() # NOTE:
+        
+    def _update_graph_backward(self):
+        if self.current_timestamp < 0:
+            raise Exception("⏰ Invalid timestamp during SeastarGraph.update_graph_backward()")
+        
+        graph_additions = self.graph_updates[str(self.current_timestamp)]["delete"]
+        graph_deletions = self.graph_updates[str(self.current_timestamp)]["add"]
+
+        edge_update_list(self._backward_graph, graph_additions)
+        edge_update_list(self._backward_graph, graph_deletions, is_delete=True)
+
+        edge_update_list(self._forward_graph, graph_additions, is_reverse_edge=True)
+        edge_update_list(self._forward_graph, graph_deletions, is_delete=True, is_reverse_edge=True)
+
+        # TODO: UNCOMMENT LATER
+        label_edges(self._forward_graph)
+        copy_label_edges(self._backward_graph, self._forward_graph)
+
+        self._get_graph_csr_ptrs()
+        # self._get_graph_attributes()  # NOTE:
\ No newline at end of file
diff --git a/seastar/graph/dynamic/gpma/README.md b/seastar/graph/dynamic/gpma/README.md
new file mode 100644
index 0000000..25c311f
--- /dev/null
+++ b/seastar/graph/dynamic/gpma/README.md
@@ -0,0 +1,11 @@
+# GPMA 
+
+Seastar provides GPMA abstraction to represent dynamic graphs in GPU.
+
+## Building GPMA CUDA/C++ extension
+
+We use PyBind11 to provide Python bindings to the GPMA API written fully in C++/CUDA. You can run the following command to build the `gpma.so` file which can be imported into your python program.
+
+```
+/usr/local/cuda-11.7/bin/nvcc $(python3 -m pybind11 --includes) -w -shared -rdc=true --compiler-options '-fPIC' -o gpma.so gpma.cu
+```
\ No newline at end of file
diff --git a/seastar/graph/dynamic/gpma/__init__.py b/seastar/graph/dynamic/gpma/__init__.py
new file mode 100644
index 0000000..741218b
--- /dev/null
+++ b/seastar/graph/dynamic/gpma/__init__.py
@@ -0,0 +1,1 @@
+'''GPMA data structure to represent Dynamic Graphs in GPU'''
\ No newline at end of file
diff --git a/seastar/graph/dynamic/gpma/gpma.cu b/seastar/graph/dynamic/gpma/gpma.cu
new file mode 100644
index 0000000..77ff211
--- /dev/null
+++ b/seastar/graph/dynamic/gpma/gpma.cu
@@ -0,0 +1,1344 @@
+#include <thrust/device_vector.h>
+#include <thrust/host_vector.h>
+#include <thrust/remove.h>
+#include <thrust/sort.h>
+#include <cub/cub.cuh>
+
+#include <pybind11/pybind11.h>
+#include <pybind11/stl.h>
+
+#include <iostream>
+#include <iomanip>
+#include <string>
+#include <vector>
+#include <tuple>
+// #include <memory>
+#include "stdio.h"
+
+namespace py = pybind11;
+using namespace pybind11::literals;
+
+#define cErr(errcode)                             \
+    {                                             \
+        gpuAssert((errcode), __FILE__, __LINE__); \
+    }
+__inline__ __host__ __device__ void gpuAssert(cudaError_t code, const char *file, int line)
+{
+    if (code != cudaSuccess)
+    {
+        printf("GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
+    }
+}
+
+typedef unsigned long long KEY_TYPE;
+typedef unsigned int VALUE_TYPE;
+typedef unsigned int SIZE_TYPE;
+
+typedef thrust::device_vector<KEY_TYPE> DEV_VEC_KEY;
+typedef thrust::device_vector<VALUE_TYPE> DEV_VEC_VALUE;
+typedef thrust::device_vector<SIZE_TYPE> DEV_VEC_SIZE;
+
+typedef KEY_TYPE *KEY_PTR;
+typedef VALUE_TYPE *VALUE_PTR;
+
+#define RAW_PTR(x) thrust::raw_pointer_cast((x).data())
+
+const KEY_TYPE KEY_NONE = 0xFFFFFFFFFFFFFFFF;
+const KEY_TYPE KEY_MAX = 0xFFFFFFFFFFFFFFFE;
+const SIZE_TYPE SIZE_NONE = 0xFFFFFFFF;
+const VALUE_TYPE VALUE_NONE = 0;
+const KEY_TYPE COL_IDX_NONE = 0xFFFFFFFF;
+
+const SIZE_TYPE MAX_BLOCKS_NUM = 96 * 8;
+#define CALC_BLOCKS_NUM(ITEMS_PER_BLOCK, CALC_SIZE) min(MAX_BLOCKS_NUM, (CALC_SIZE - 1) / ITEMS_PER_BLOCK + 1)
+
+class GPMA
+{
+public:
+    // NOTE: Assuming that keys holds the column_indices
+    // and values is the column indices (subject to change)
+    DEV_VEC_KEY keys;
+    DEV_VEC_VALUE values;
+
+    SIZE_TYPE segment_length;
+    SIZE_TYPE tree_height;
+
+    double density_lower_thres_leaf = 0.08;
+    double density_lower_thres_root = 0.42;
+    double density_upper_thres_root = 0.84;
+    double density_upper_thres_leaf = 0.92;
+
+    thrust::host_vector<SIZE_TYPE> lower_element;
+    thrust::host_vector<SIZE_TYPE> upper_element;
+
+    // addition for csr
+    SIZE_TYPE row_num; // number of nodes
+    DEV_VEC_SIZE row_offset;
+
+    int edge_count = 0;
+    std::vector<unsigned int> in_degree;
+    std::vector<unsigned int> out_degree;
+
+    GPMA();
+
+    inline int get_size()
+    {
+        return keys.size();
+    }
+};
+
+GPMA::GPMA()
+{
+}
+
+__forceinline__ __host__ __device__
+    SIZE_TYPE
+    fls(SIZE_TYPE x)
+{
+    SIZE_TYPE r = 32;
+    if (!x)
+        return 0;
+    if (!(x & 0xffff0000u))
+        x <<= 16, r -= 16;
+    if (!(x & 0xff000000u))
+        x <<= 8, r -= 8;
+    if (!(x & 0xf0000000u))
+        x <<= 4, r -= 4;
+    if (!(x & 0xc0000000u))
+        x <<= 2, r -= 2;
+    if (!(x & 0x80000000u))
+        x <<= 1, r -= 1;
+    return r;
+}
+
+template <typename T>
+__global__ void memcpy_kernel(T *dest, const T *src, SIZE_TYPE size)
+{
+    SIZE_TYPE global_thread_id = blockDim.x * blockIdx.x + threadIdx.x;
+    SIZE_TYPE block_offset = gridDim.x * blockDim.x;
+    for (SIZE_TYPE i = global_thread_id; i < size; i += block_offset)
+    {
+        dest[i] = src[i];
+    }
+}
+
+template <typename T>
+__global__ void memset_kernel(T *data, T value, SIZE_TYPE size)
+{
+    SIZE_TYPE global_thread_id = blockDim.x * blockIdx.x + threadIdx.x;
+    SIZE_TYPE block_offset = gridDim.x * blockDim.x;
+    for (SIZE_TYPE i = global_thread_id; i < size; i += block_offset)
+    {
+        data[i] = value;
+    }
+}
+
+__host__ void recalculate_density(GPMA &gpma)
+{
+    gpma.lower_element.resize(gpma.tree_height + 1);
+    gpma.upper_element.resize(gpma.tree_height + 1);
+    cErr(cudaDeviceSynchronize());
+
+    SIZE_TYPE level_length = gpma.segment_length;
+    for (SIZE_TYPE i = 0; i <= gpma.tree_height; i++)
+    {
+        double density_lower = gpma.density_lower_thres_root + (gpma.density_lower_thres_leaf - gpma.density_lower_thres_root) * (gpma.tree_height - i) / gpma.tree_height;
+        double density_upper = gpma.density_upper_thres_root + (gpma.density_upper_thres_leaf - gpma.density_upper_thres_root) * (gpma.tree_height - i) / gpma.tree_height;
+
+        gpma.lower_element[i] = (SIZE_TYPE)ceil(density_lower * level_length);
+        gpma.upper_element[i] = (SIZE_TYPE)floor(density_upper * level_length);
+
+        // special trim for wrong threshold introduced by float-integer conversion
+        if (0 < i)
+        {
+            gpma.lower_element[i] = max(gpma.lower_element[i], 2 * gpma.lower_element[i - 1]);
+            gpma.upper_element[i] = min(gpma.upper_element[i], 2 * gpma.upper_element[i - 1]);
+        }
+        level_length <<= 1;
+    }
+}
+
+__device__ void cub_sort_key_value(KEY_TYPE *keys, VALUE_TYPE *values, SIZE_TYPE size, KEY_TYPE *tmp_keys,
+                                   VALUE_TYPE *tmp_values)
+{
+
+    void *d_temp_storage = NULL;
+    size_t temp_storage_bytes = 0;
+
+    cErr(cub::DeviceRadixSort::SortPairs(d_temp_storage, temp_storage_bytes, keys, tmp_keys, values, tmp_values, size));
+    cErr(cudaDeviceSynchronize());
+    cErr(cudaMalloc(&d_temp_storage, temp_storage_bytes));
+    cErr(cudaDeviceSynchronize());
+    cErr(cub::DeviceRadixSort::SortPairs(d_temp_storage, temp_storage_bytes, keys, tmp_keys, values, tmp_values, size));
+    cErr(cudaDeviceSynchronize());
+
+    SIZE_TYPE THREADS_NUM = 128;
+    SIZE_TYPE BLOCKS_NUM = CALC_BLOCKS_NUM(THREADS_NUM, size);
+    memcpy_kernel<KEY_TYPE><<<BLOCKS_NUM, THREADS_NUM>>>(keys, tmp_keys, size);
+    memcpy_kernel<VALUE_TYPE><<<BLOCKS_NUM, THREADS_NUM>>>(values, tmp_values, size);
+    cErr(cudaDeviceSynchronize());
+
+    cErr(cudaFree(d_temp_storage));
+}
+
+__device__ SIZE_TYPE handle_del_mod(KEY_TYPE *keys, VALUE_TYPE *values, SIZE_TYPE seg_length, KEY_TYPE key,
+                                    VALUE_TYPE value, SIZE_TYPE leaf)
+{
+
+    if (VALUE_NONE == value)
+        leaf = SIZE_NONE;
+    for (SIZE_TYPE i = 0; i < seg_length; i++)
+    {
+        if (keys[i] == key)
+        {
+            values[i] = value;
+            leaf = SIZE_NONE;
+            break;
+        }
+    }
+    return leaf;
+}
+
+__global__ void locate_leaf_kernel(KEY_TYPE *keys, VALUE_TYPE *values, SIZE_TYPE tree_size, SIZE_TYPE seg_length,
+                                   SIZE_TYPE tree_height, KEY_TYPE *update_keys, VALUE_TYPE *update_values, SIZE_TYPE update_size,
+                                   SIZE_TYPE *leaf, bool return_leaf_loc = true)
+{
+
+    SIZE_TYPE global_thread_id = blockDim.x * blockIdx.x + threadIdx.x;
+    SIZE_TYPE block_offset = gridDim.x * blockDim.x;
+    for (SIZE_TYPE i = global_thread_id; i < update_size; i += block_offset)
+    {
+        KEY_TYPE key = update_keys[i];
+        VALUE_TYPE value = update_values[i];
+
+        SIZE_TYPE prefix = 0;
+        SIZE_TYPE current_bit = seg_length << tree_height >> 1;
+
+        while (seg_length <= current_bit)
+        {
+            if (keys[prefix | current_bit] <= key)
+                prefix |= current_bit;
+            current_bit >>= 1;
+        }
+
+        prefix = handle_del_mod(keys + prefix, values + prefix, seg_length, key, value, prefix);
+
+        if (return_leaf_loc)
+            leaf[i] = prefix;
+    }
+}
+
+__host__ void locate_leaf_batch(KEY_TYPE *keys, VALUE_TYPE *values, SIZE_TYPE tree_size, SIZE_TYPE seg_length,
+                                SIZE_TYPE tree_height, KEY_TYPE *update_keys, VALUE_TYPE *update_values, SIZE_TYPE update_size,
+                                SIZE_TYPE *leaf, bool return_leaf_loc = true)
+{
+
+    SIZE_TYPE THREADS_NUM = 32;
+    SIZE_TYPE BLOCKS_NUM = CALC_BLOCKS_NUM(THREADS_NUM, update_size);
+
+    locate_leaf_kernel<<<BLOCKS_NUM, THREADS_NUM>>>(keys, values, tree_size, seg_length, tree_height, update_keys,
+                                                    update_values, update_size, leaf, return_leaf_loc);
+    cErr(cudaDeviceSynchronize());
+}
+
+template <SIZE_TYPE THREAD_PER_BLOCK, SIZE_TYPE ITEM_PER_THREAD>
+__device__ void block_compact_kernel(KEY_TYPE *keys, VALUE_TYPE *values, SIZE_TYPE &compacted_size)
+{
+    typedef cub::BlockScan<SIZE_TYPE, THREAD_PER_BLOCK> BlockScan;
+    SIZE_TYPE thread_id = threadIdx.x;
+
+    KEY_TYPE *block_keys = keys;
+    VALUE_TYPE *block_values = values;
+
+    KEY_TYPE thread_keys[ITEM_PER_THREAD];
+    VALUE_TYPE thread_values[ITEM_PER_THREAD];
+
+    SIZE_TYPE thread_offset = thread_id * ITEM_PER_THREAD;
+    for (SIZE_TYPE i = 0; i < ITEM_PER_THREAD; i++)
+    {
+        thread_keys[i] = block_keys[thread_offset + i];
+        thread_values[i] = block_values[thread_offset + i];
+        block_keys[thread_offset + i] = KEY_NONE;
+    }
+
+    __shared__ typename BlockScan::TempStorage temp_storage;
+    SIZE_TYPE thread_data[ITEM_PER_THREAD];
+    for (SIZE_TYPE i = 0; i < ITEM_PER_THREAD; i++)
+    {
+        thread_data[i] = (thread_keys[i] == KEY_NONE || thread_values[i] == VALUE_NONE) ? 0 : 1;
+    }
+    __syncthreads();
+
+    BlockScan(temp_storage).ExclusiveSum(thread_data, thread_data);
+    __syncthreads();
+
+    __shared__ SIZE_TYPE exscan[THREAD_PER_BLOCK * ITEM_PER_THREAD];
+    for (SIZE_TYPE i = 0; i < ITEM_PER_THREAD; i++)
+    {
+        exscan[i + thread_offset] = thread_data[i];
+    }
+    __syncthreads();
+
+    for (SIZE_TYPE i = 0; i < ITEM_PER_THREAD; i++)
+    {
+        if (thread_id == THREAD_PER_BLOCK - 1 && i == ITEM_PER_THREAD - 1)
+            continue;
+        if (exscan[thread_offset + i] != exscan[thread_offset + i + 1])
+        {
+            SIZE_TYPE loc = exscan[thread_offset + i];
+            block_keys[loc] = thread_keys[i];
+            block_values[loc] = thread_values[i];
+        }
+    }
+
+    // special logic for the last element
+    if (thread_id == THREAD_PER_BLOCK - 1)
+    {
+        SIZE_TYPE loc = exscan[THREAD_PER_BLOCK * ITEM_PER_THREAD - 1];
+        if (thread_keys[ITEM_PER_THREAD - 1] == KEY_NONE || thread_values[ITEM_PER_THREAD - 1] == VALUE_NONE)
+        {
+            compacted_size = loc;
+        }
+        else
+        {
+            compacted_size = loc + 1;
+            block_keys[loc] = thread_keys[ITEM_PER_THREAD - 1];
+            block_values[loc] = thread_values[ITEM_PER_THREAD - 1];
+        }
+    }
+}
+
+template <typename FIRST_TYPE, typename SECOND_TYPE>
+__device__ void block_pair_copy_kernel(FIRST_TYPE *dest_first, SECOND_TYPE *dest_second, FIRST_TYPE *src_first,
+                                       SECOND_TYPE *src_second, SIZE_TYPE size)
+{
+    for (SIZE_TYPE i = threadIdx.x; i < size; i += blockDim.x)
+    {
+        dest_first[i] = src_first[i];
+        dest_second[i] = src_second[i];
+    }
+}
+
+template <SIZE_TYPE THREAD_PER_BLOCK, SIZE_TYPE ITEM_PER_THREAD>
+__device__ void block_redispatch_kernel(KEY_TYPE *keys, VALUE_TYPE *values, SIZE_TYPE rebalance_width, SIZE_TYPE seg_length,
+                                        SIZE_TYPE merge_size, SIZE_TYPE *row_offset, SIZE_TYPE update_node)
+{
+
+    // step1: load KV in shared memory
+    __shared__ KEY_TYPE block_keys[THREAD_PER_BLOCK * ITEM_PER_THREAD];
+    __shared__ VALUE_TYPE block_values[THREAD_PER_BLOCK * ITEM_PER_THREAD];
+    block_pair_copy_kernel<KEY_TYPE, VALUE_TYPE>(block_keys, block_values, keys, values, rebalance_width);
+    __syncthreads();
+
+    // step2: sort by key with value on shared memory
+    typedef cub::BlockLoad<KEY_TYPE, THREAD_PER_BLOCK, ITEM_PER_THREAD, cub::BLOCK_LOAD_TRANSPOSE> BlockKeyLoadT;
+    typedef cub::BlockLoad<VALUE_TYPE, THREAD_PER_BLOCK, ITEM_PER_THREAD, cub::BLOCK_LOAD_TRANSPOSE> BlockValueLoadT;
+    typedef cub::BlockStore<KEY_TYPE, THREAD_PER_BLOCK, ITEM_PER_THREAD, cub::BLOCK_STORE_TRANSPOSE> BlockKeyStoreT;
+    typedef cub::BlockStore<VALUE_TYPE, THREAD_PER_BLOCK, ITEM_PER_THREAD, cub::BLOCK_STORE_TRANSPOSE> BlockValueStoreT;
+    typedef cub::BlockRadixSort<KEY_TYPE, THREAD_PER_BLOCK, ITEM_PER_THREAD, VALUE_TYPE> BlockRadixSortT;
+
+    __shared__ union
+    {
+        typename BlockKeyLoadT::TempStorage key_load;
+        typename BlockValueLoadT::TempStorage value_load;
+        typename BlockKeyStoreT::TempStorage key_store;
+        typename BlockValueStoreT::TempStorage value_store;
+        typename BlockRadixSortT::TempStorage sort;
+    } temp_storage;
+
+    KEY_TYPE thread_keys[ITEM_PER_THREAD];
+    VALUE_TYPE thread_values[ITEM_PER_THREAD];
+    BlockKeyLoadT(temp_storage.key_load).Load(block_keys, thread_keys);
+    BlockValueLoadT(temp_storage.value_load).Load(block_values, thread_values);
+    __syncthreads();
+
+    BlockRadixSortT(temp_storage.sort).Sort(thread_keys, thread_values);
+    __syncthreads();
+
+    BlockKeyStoreT(temp_storage.key_store).Store(block_keys, thread_keys);
+    BlockValueStoreT(temp_storage.value_store).Store(block_values, thread_values);
+    __syncthreads();
+
+    // step3: evenly re-dispatch KVs to leaf segments
+    KEY_TYPE frac = rebalance_width / seg_length;
+    KEY_TYPE deno = merge_size;
+    for (SIZE_TYPE i = threadIdx.x; i < merge_size; i += blockDim.x)
+    {
+        keys[i] = KEY_NONE;
+    }
+    __syncthreads();
+
+    for (SIZE_TYPE i = threadIdx.x; i < merge_size; i += blockDim.x)
+    {
+        SIZE_TYPE seg_idx = (SIZE_TYPE)(frac * i / deno);
+        SIZE_TYPE seg_lane = (SIZE_TYPE)(frac * i % deno / frac);
+        SIZE_TYPE proj_location = seg_idx * seg_length + seg_lane;
+
+        KEY_TYPE cur_key = block_keys[i];
+        VALUE_TYPE cur_value = block_values[i];
+        keys[proj_location] = cur_key;
+        values[proj_location] = cur_value;
+
+        // addition for csr
+        if ((cur_key & COL_IDX_NONE) == COL_IDX_NONE)
+        {
+            SIZE_TYPE cur_row = (SIZE_TYPE)(cur_key >> 32);
+            row_offset[cur_row + 1] = proj_location + update_node;
+        }
+    }
+}
+
+template <SIZE_TYPE THREAD_PER_BLOCK, SIZE_TYPE ITEM_PER_THREAD>
+__global__ void block_rebalancing_kernel(SIZE_TYPE seg_length, SIZE_TYPE level, KEY_TYPE *keys, VALUE_TYPE *values,
+                                         SIZE_TYPE *update_nodes, KEY_TYPE *update_keys, VALUE_TYPE *update_values, SIZE_TYPE *unique_update_nodes,
+                                         SIZE_TYPE *update_offset, SIZE_TYPE lower_bound, SIZE_TYPE upper_bound, SIZE_TYPE *row_offset)
+{
+
+    SIZE_TYPE update_id = blockIdx.x;
+    SIZE_TYPE update_node = unique_update_nodes[update_id];
+    KEY_TYPE *key = keys + update_node;
+    VALUE_TYPE *value = values + update_node;
+    SIZE_TYPE rebalance_width = seg_length << level;
+
+    // compact
+    __shared__ SIZE_TYPE compacted_size;
+    block_compact_kernel<THREAD_PER_BLOCK, ITEM_PER_THREAD>(key, value, compacted_size);
+    __syncthreads();
+
+    // judge whether fit the density threshold
+    SIZE_TYPE interval_a = update_offset[update_id];
+    SIZE_TYPE interval_b = update_offset[update_id + 1];
+    SIZE_TYPE interval_size = interval_b - interval_a;
+    SIZE_TYPE merge_size = compacted_size + interval_size;
+    __syncthreads();
+
+    if (lower_bound <= merge_size && merge_size <= upper_bound)
+    {
+        // move
+        block_pair_copy_kernel<KEY_TYPE, VALUE_TYPE>(key + compacted_size, value + compacted_size,
+                                                     update_keys + interval_a, update_values + interval_a, interval_size);
+        __syncthreads();
+
+        // set SIZE_NONE for executed update
+        for (SIZE_TYPE i = interval_a + threadIdx.x; i < interval_b; i += blockDim.x)
+        {
+            update_nodes[i] = SIZE_NONE;
+        }
+
+        // re-dispatch
+        block_redispatch_kernel<THREAD_PER_BLOCK, ITEM_PER_THREAD>(key, value, rebalance_width, seg_length,
+                                                                   merge_size, row_offset, update_node);
+    }
+}
+
+__global__ void label_key_whether_none_kernel(SIZE_TYPE *label, KEY_TYPE *keys, VALUE_TYPE *values, SIZE_TYPE size)
+{
+    SIZE_TYPE global_thread_id = blockDim.x * blockIdx.x + threadIdx.x;
+    SIZE_TYPE block_offset = gridDim.x * blockDim.x;
+    for (SIZE_TYPE i = global_thread_id; i < size; i += block_offset)
+    {
+        label[i] = (keys[i] == KEY_NONE || values[i] == VALUE_NONE) ? 0 : 1;
+    }
+}
+
+__global__ void copy_compacted_kv(SIZE_TYPE *exscan, KEY_TYPE *keys, VALUE_TYPE *values, SIZE_TYPE size, KEY_TYPE *tmp_keys,
+                                  VALUE_TYPE *tmp_values, SIZE_TYPE *compacted_size)
+{
+
+    SIZE_TYPE global_thread_id = blockDim.x * blockIdx.x + threadIdx.x;
+    SIZE_TYPE block_offset = gridDim.x * blockDim.x;
+    for (SIZE_TYPE i = global_thread_id; i < size; i += block_offset)
+    {
+        if (i == size - 1)
+            continue;
+        if (exscan[i] != exscan[i + 1])
+        {
+            SIZE_TYPE loc = exscan[i];
+            tmp_keys[loc] = keys[i];
+            tmp_values[loc] = values[i];
+        }
+    }
+
+    if (0 == global_thread_id)
+    {
+        SIZE_TYPE loc = exscan[size - 1];
+        if (keys[size - 1] == KEY_NONE || values[size - 1] == VALUE_NONE)
+        {
+            *compacted_size = loc;
+        }
+        else
+        {
+            *compacted_size = loc + 1;
+            tmp_keys[loc] = keys[size - 1];
+            tmp_values[loc] = values[size - 1];
+        }
+    }
+}
+
+__device__ void compact_kernel(SIZE_TYPE size, KEY_TYPE *keys, VALUE_TYPE *values, SIZE_TYPE *compacted_size,
+                               KEY_TYPE *tmp_keys, VALUE_TYPE *tmp_values, SIZE_TYPE *exscan, SIZE_TYPE *label)
+{
+
+    SIZE_TYPE THREADS_NUM = 32;
+    SIZE_TYPE BLOCKS_NUM = CALC_BLOCKS_NUM(THREADS_NUM, size);
+    label_key_whether_none_kernel<<<BLOCKS_NUM, THREADS_NUM>>>(label, keys, values, size);
+    cErr(cudaDeviceSynchronize());
+
+    // exscan
+    void *d_temp_storage = NULL;
+    size_t temp_storage_bytes = 0;
+    cErr(cub::DeviceScan::ExclusiveSum(d_temp_storage, temp_storage_bytes, label, exscan, size));
+    cErr(cudaDeviceSynchronize());
+    cErr(cudaMalloc(&d_temp_storage, temp_storage_bytes));
+    cErr(cudaDeviceSynchronize());
+    cErr(cub::DeviceScan::ExclusiveSum(d_temp_storage, temp_storage_bytes, label, exscan, size));
+    cErr(cudaDeviceSynchronize());
+    cErr(cudaFree(d_temp_storage));
+
+    // copy compacted kv to tmp, and set the original to none
+    copy_compacted_kv<<<BLOCKS_NUM, THREADS_NUM>>>(exscan, keys, values, size, tmp_keys, tmp_values, compacted_size);
+    cErr(cudaDeviceSynchronize());
+}
+
+__global__ void redispatch_kernel(KEY_TYPE *tmp_keys, VALUE_TYPE *tmp_values, KEY_TYPE *keys, VALUE_TYPE *values,
+                                  SIZE_TYPE update_width, SIZE_TYPE seg_length, SIZE_TYPE merge_size, SIZE_TYPE *row_offset,
+                                  SIZE_TYPE update_node)
+{
+
+    SIZE_TYPE global_thread_id = blockDim.x * blockIdx.x + threadIdx.x;
+    SIZE_TYPE block_offset = gridDim.x * blockDim.x;
+    KEY_TYPE frac = update_width / seg_length;
+    KEY_TYPE deno = merge_size;
+
+    for (SIZE_TYPE i = global_thread_id; i < merge_size; i += block_offset)
+    {
+        SIZE_TYPE seg_idx = (SIZE_TYPE)(frac * i / deno);
+        SIZE_TYPE seg_lane = (SIZE_TYPE)(frac * i % deno / frac);
+        SIZE_TYPE proj_location = seg_idx * seg_length + seg_lane;
+        KEY_TYPE cur_key = tmp_keys[i];
+        VALUE_TYPE cur_value = tmp_values[i];
+        keys[proj_location] = cur_key;
+        values[proj_location] = cur_value;
+
+        // addition for csr
+        if ((cur_key & COL_IDX_NONE) == COL_IDX_NONE)
+        {
+            SIZE_TYPE cur_row = (SIZE_TYPE)(cur_key >> 32);
+            row_offset[cur_row + 1] = proj_location + update_node;
+        }
+    }
+}
+
+__global__ void rebalancing_kernel(SIZE_TYPE unique_update_size, SIZE_TYPE seg_length, SIZE_TYPE level, KEY_TYPE *keys,
+                                   VALUE_TYPE *values, SIZE_TYPE *update_nodes, KEY_TYPE *update_keys, VALUE_TYPE *update_values,
+                                   SIZE_TYPE *unique_update_nodes, SIZE_TYPE *update_offset, SIZE_TYPE lower_bound, SIZE_TYPE upper_bound,
+                                   SIZE_TYPE *row_offset)
+{
+
+    SIZE_TYPE global_thread_id = blockDim.x * blockIdx.x + threadIdx.x;
+    SIZE_TYPE block_offset = gridDim.x * blockDim.x;
+    SIZE_TYPE update_width = seg_length << level;
+
+    SIZE_TYPE *compacted_size;
+    cErr(cudaMalloc(&compacted_size, sizeof(SIZE_TYPE)));
+    cErr(cudaDeviceSynchronize());
+
+    KEY_TYPE *tmp_keys;
+    VALUE_TYPE *tmp_values;
+    SIZE_TYPE *tmp_exscan;
+    SIZE_TYPE *tmp_label;
+
+    cErr(cudaMalloc(&tmp_keys, update_width * sizeof(KEY_TYPE)));
+    cErr(cudaMalloc(&tmp_values, update_width * sizeof(VALUE_TYPE)));
+    cErr(cudaMalloc(&tmp_exscan, update_width * sizeof(SIZE_TYPE)));
+    cErr(cudaMalloc(&tmp_label, update_width * sizeof(SIZE_TYPE)));
+    cErr(cudaDeviceSynchronize());
+
+    for (SIZE_TYPE i = global_thread_id; i < unique_update_size; i += block_offset)
+    {
+        SIZE_TYPE update_node = unique_update_nodes[i];
+        KEY_TYPE *key = keys + update_node;
+        VALUE_TYPE *value = values + update_node;
+
+        // compact
+        compact_kernel(update_width, key, value, compacted_size, tmp_keys, tmp_values, tmp_exscan, tmp_label);
+        cErr(cudaDeviceSynchronize());
+
+        // judge whether fit the density threshold
+        SIZE_TYPE interval_a = update_offset[i];
+        SIZE_TYPE interval_b = update_offset[i + 1];
+        SIZE_TYPE interval_size = interval_b - interval_a;
+        SIZE_TYPE merge_size = (*compacted_size) + interval_size;
+
+        if (lower_bound <= merge_size && merge_size <= upper_bound)
+        {
+            SIZE_TYPE THREADS_NUM = 32;
+            SIZE_TYPE BLOCKS_NUM;
+
+            // move
+            BLOCKS_NUM = CALC_BLOCKS_NUM(THREADS_NUM, interval_size);
+            memcpy_kernel<KEY_TYPE><<<BLOCKS_NUM, THREADS_NUM>>>(tmp_keys + (*compacted_size),
+                                                                 update_keys + interval_a, interval_size);
+            memcpy_kernel<VALUE_TYPE><<<BLOCKS_NUM, THREADS_NUM>>>(tmp_values + (*compacted_size),
+                                                                   update_values + interval_a, interval_size);
+            cErr(cudaDeviceSynchronize());
+
+            // set SIZE_NONE for executed updates
+            memset_kernel<SIZE_TYPE><<<BLOCKS_NUM, THREADS_NUM>>>(update_nodes + interval_a, SIZE_NONE, interval_size);
+            cErr(cudaDeviceSynchronize());
+
+            cub_sort_key_value(tmp_keys, tmp_values, merge_size, key, value);
+
+            // re-dispatch
+            BLOCKS_NUM = CALC_BLOCKS_NUM(THREADS_NUM, update_width);
+            memset_kernel<KEY_TYPE><<<BLOCKS_NUM, THREADS_NUM>>>(key, KEY_NONE, update_width);
+            cErr(cudaDeviceSynchronize());
+
+            BLOCKS_NUM = CALC_BLOCKS_NUM(THREADS_NUM, merge_size);
+            redispatch_kernel<<<BLOCKS_NUM, THREADS_NUM>>>(tmp_keys, tmp_values, key, value, update_width, seg_length,
+                                                           merge_size, row_offset, update_node);
+            cErr(cudaDeviceSynchronize());
+        }
+
+        cErr(cudaDeviceSynchronize());
+    }
+
+    cErr(cudaFree(compacted_size));
+    cErr(cudaFree(tmp_keys));
+    cErr(cudaFree(tmp_values));
+    cErr(cudaFree(tmp_exscan));
+    cErr(cudaFree(tmp_label));
+}
+
+__host__ void rebalance_batch(SIZE_TYPE level, SIZE_TYPE seg_length, KEY_TYPE *keys, VALUE_TYPE *values,
+                              SIZE_TYPE *update_nodes, KEY_TYPE *update_keys, VALUE_TYPE *update_values, SIZE_TYPE update_size,
+                              SIZE_TYPE *unique_update_nodes, SIZE_TYPE *update_offset, SIZE_TYPE unique_update_size,
+                              SIZE_TYPE lower_bound, SIZE_TYPE upper_bound, SIZE_TYPE *row_offset)
+{
+
+    SIZE_TYPE update_width = seg_length << level;
+
+    if (update_width <= 1024)
+    {
+        // func pointer for each template
+        void (*func_arr[10])(SIZE_TYPE, SIZE_TYPE, KEY_TYPE *, VALUE_TYPE *, SIZE_TYPE *, KEY_TYPE *, VALUE_TYPE *,
+                             SIZE_TYPE *, SIZE_TYPE *, SIZE_TYPE, SIZE_TYPE, SIZE_TYPE *);
+        func_arr[0] = block_rebalancing_kernel<2, 1>;
+        func_arr[1] = block_rebalancing_kernel<4, 1>;
+        func_arr[2] = block_rebalancing_kernel<8, 1>;
+        func_arr[3] = block_rebalancing_kernel<16, 1>;
+        func_arr[4] = block_rebalancing_kernel<32, 1>;
+        func_arr[5] = block_rebalancing_kernel<32, 2>;
+        func_arr[6] = block_rebalancing_kernel<32, 4>;
+        func_arr[7] = block_rebalancing_kernel<32, 8>;
+        func_arr[8] = block_rebalancing_kernel<32, 16>;
+        func_arr[9] = block_rebalancing_kernel<32, 32>;
+
+        // operate each tree node by cuda-block
+        SIZE_TYPE THREADS_NUM = update_width > 32 ? 32 : update_width;
+        SIZE_TYPE BLOCKS_NUM = unique_update_size;
+
+        func_arr[fls(update_width) - 2]<<<BLOCKS_NUM, THREADS_NUM>>>(seg_length, level, keys, values, update_nodes,
+                                                                     update_keys, update_values, unique_update_nodes, update_offset, lower_bound, upper_bound, row_offset);
+    }
+    else
+    {
+        // operate each tree node by cub-kernel (dynamic parallelsim)
+        SIZE_TYPE BLOCKS_NUM = min(2048, unique_update_size);
+        rebalancing_kernel<<<BLOCKS_NUM, 1>>>(unique_update_size, seg_length, level, keys, values, update_nodes,
+                                              update_keys, update_values, unique_update_nodes, update_offset, lower_bound, upper_bound, row_offset);
+    }
+    cErr(cudaDeviceSynchronize());
+}
+
+struct three_tuple_first_none
+{
+    typedef thrust::tuple<SIZE_TYPE, KEY_TYPE, VALUE_TYPE> Tuple;
+    __host__ __device__ bool operator()(const Tuple &a)
+    {
+        return SIZE_NONE == thrust::get<0>(a);
+    }
+};
+__host__ void compact_insertions(DEV_VEC_SIZE &update_nodes, DEV_VEC_KEY &update_keys, DEV_VEC_VALUE &update_values,
+                                 SIZE_TYPE &update_size)
+{
+
+    auto zip_begin = thrust::make_zip_iterator(
+        thrust::make_tuple(update_nodes.begin(), update_keys.begin(), update_values.begin()));
+    auto zip_end = thrust::remove_if(zip_begin, zip_begin + update_size, three_tuple_first_none());
+    cErr(cudaDeviceSynchronize());
+    update_size = zip_end - zip_begin;
+}
+
+__host__ SIZE_TYPE group_insertion_by_node(SIZE_TYPE *update_nodes, SIZE_TYPE update_size,
+                                           SIZE_TYPE *unique_update_nodes, SIZE_TYPE *update_offset)
+{
+
+    // step1: encode
+    void *d_temp_storage = NULL;
+    size_t temp_storage_bytes = 0;
+    SIZE_TYPE *tmp_offset;
+    cErr(cudaMalloc(&tmp_offset, sizeof(SIZE_TYPE) * update_size));
+
+    SIZE_TYPE *num_runs_out;
+    cErr(cudaMalloc(&num_runs_out, sizeof(SIZE_TYPE)));
+    cErr(cudaDeviceSynchronize());
+    cErr(cub::DeviceRunLengthEncode::Encode(d_temp_storage, temp_storage_bytes, update_nodes,
+                                            unique_update_nodes, tmp_offset, num_runs_out, update_size));
+    cErr(cudaDeviceSynchronize());
+    cErr(cudaMalloc(&d_temp_storage, temp_storage_bytes));
+    cErr(cudaDeviceSynchronize());
+    cErr(cub::DeviceRunLengthEncode::Encode(d_temp_storage, temp_storage_bytes, update_nodes,
+                                            unique_update_nodes, tmp_offset, num_runs_out, update_size));
+    cErr(cudaDeviceSynchronize());
+
+    SIZE_TYPE unique_node_size[1];
+    cErr(cudaMemcpy(unique_node_size, num_runs_out, sizeof(SIZE_TYPE), cudaMemcpyDeviceToHost));
+    cErr(cudaDeviceSynchronize());
+    cErr(cudaFree(num_runs_out));
+    cErr(cudaFree(d_temp_storage));
+
+    // step2: exclusive scan
+    d_temp_storage = NULL;
+    temp_storage_bytes = 0;
+    cErr(cub::DeviceScan::ExclusiveSum(d_temp_storage, temp_storage_bytes, tmp_offset,
+                                       update_offset, unique_node_size[0]));
+    cErr(cudaDeviceSynchronize());
+    cErr(cudaMalloc(&d_temp_storage, temp_storage_bytes));
+    cErr(cudaDeviceSynchronize());
+    cErr(cub::DeviceScan::ExclusiveSum(d_temp_storage, temp_storage_bytes, tmp_offset,
+                                       update_offset, unique_node_size[0]));
+    cErr(cudaDeviceSynchronize());
+    cErr(cudaFree(d_temp_storage));
+
+    cErr(cudaMemcpy(update_offset + unique_node_size[0], &update_size, sizeof(SIZE_TYPE), cudaMemcpyHostToDevice));
+    cErr(cudaDeviceSynchronize());
+    cErr(cudaFree(tmp_offset));
+
+    return unique_node_size[0];
+}
+
+__host__ void compress_insertions_by_node(DEV_VEC_SIZE &update_nodes, SIZE_TYPE update_size,
+                                          DEV_VEC_SIZE &unique_update_nodes, DEV_VEC_SIZE &update_offset, SIZE_TYPE &unique_node_size)
+{
+    unique_node_size = group_insertion_by_node(RAW_PTR(update_nodes), update_size, RAW_PTR(unique_update_nodes),
+                                               RAW_PTR(update_offset));
+    cErr(cudaDeviceSynchronize());
+}
+
+__global__ void up_level_kernel(SIZE_TYPE *update_nodes, SIZE_TYPE update_size, SIZE_TYPE update_width)
+{
+    SIZE_TYPE global_thread_id = blockDim.x * blockIdx.x + threadIdx.x;
+    SIZE_TYPE block_offset = gridDim.x * blockDim.x;
+
+    for (SIZE_TYPE i = global_thread_id; i < update_size; i += block_offset)
+    {
+        SIZE_TYPE node = update_nodes[i];
+        update_nodes[i] = node & ~update_width;
+    }
+}
+
+__host__ void up_level_batch(SIZE_TYPE *update_nodes, SIZE_TYPE update_size, SIZE_TYPE update_width)
+{
+    SIZE_TYPE THREADS_NUM = 32;
+    SIZE_TYPE BLOCKS_NUM = CALC_BLOCKS_NUM(THREADS_NUM, update_size);
+    up_level_kernel<<<BLOCKS_NUM, THREADS_NUM>>>(update_nodes, update_size, update_width);
+    cErr(cudaDeviceSynchronize());
+}
+
+struct kv_tuple_none
+{
+    typedef thrust::tuple<KEY_TYPE, VALUE_TYPE> Tuple;
+    __host__ __device__ bool operator()(const Tuple &a)
+    {
+        return KEY_NONE == thrust::get<0>(a) || VALUE_NONE == thrust::get<1>(a);
+    }
+};
+__host__ int resize_gpma(GPMA &gpma, DEV_VEC_KEY &update_keys, DEV_VEC_VALUE &update_values, SIZE_TYPE update_size)
+{
+    auto zip_begin = thrust::make_zip_iterator(thrust::make_tuple(gpma.keys.begin(), gpma.values.begin()));
+    auto zip_end = thrust::remove_if(zip_begin, zip_begin + gpma.keys.size(), kv_tuple_none());
+    cErr(cudaDeviceSynchronize());
+    SIZE_TYPE compacted_size = zip_end - zip_begin;
+    thrust::fill(gpma.keys.begin() + compacted_size, gpma.keys.end(), KEY_NONE);
+    cErr(cudaDeviceSynchronize());
+
+    SIZE_TYPE merge_size = compacted_size + update_size;
+    SIZE_TYPE original_tree_size = gpma.keys.size();
+
+    SIZE_TYPE tree_size = 4;
+    while (floor(gpma.density_upper_thres_root * tree_size) < merge_size)
+        tree_size <<= 1;
+    gpma.segment_length = 1 << (fls(fls(tree_size)) - 1);
+    gpma.tree_height = fls(tree_size / gpma.segment_length) - 1;
+
+    gpma.keys.resize(tree_size, KEY_NONE);
+    gpma.values.resize(tree_size);
+    cErr(cudaDeviceSynchronize());
+    recalculate_density(gpma);
+
+    return compacted_size;
+}
+
+__host__ void significant_insert(GPMA &gpma, DEV_VEC_KEY &update_keys, DEV_VEC_VALUE &update_values, int update_size)
+{
+    int valid_size = resize_gpma(gpma, update_keys, update_values, update_size);
+    thrust::copy(update_keys.begin(), update_keys.begin() + update_size, gpma.keys.begin() + valid_size);
+    thrust::copy(update_values.begin(), update_values.begin() + update_size, gpma.values.begin() + valid_size);
+
+    DEV_VEC_KEY tmp_update_keys(gpma.get_size());
+    DEV_VEC_VALUE tmp_update_values(gpma.get_size());
+    cErr(cudaDeviceSynchronize());
+
+    int merge_size = valid_size + update_size;
+    thrust::sort_by_key(gpma.keys.begin(), gpma.keys.begin() + merge_size, gpma.values.begin());
+    cErr(cudaDeviceSynchronize());
+
+    SIZE_TYPE THREADS_NUM = 32;
+    SIZE_TYPE BLOCKS_NUM;
+    BLOCKS_NUM = CALC_BLOCKS_NUM(THREADS_NUM, merge_size);
+    redispatch_kernel<<<BLOCKS_NUM, THREADS_NUM>>>(RAW_PTR(gpma.keys), RAW_PTR(gpma.values), RAW_PTR(tmp_update_keys),
+                                                   RAW_PTR(tmp_update_values), gpma.get_size(), gpma.segment_length, merge_size, RAW_PTR(gpma.row_offset), 0);
+    cErr(cudaDeviceSynchronize());
+
+    gpma.keys = tmp_update_keys;
+    gpma.values = tmp_update_values;
+    cErr(cudaDeviceSynchronize());
+}
+
+__host__ void update_gpma(GPMA &gpma, DEV_VEC_KEY &update_keys, DEV_VEC_VALUE &update_values)
+{
+    SIZE_TYPE ous = update_keys.size();
+
+    // step1: sort update keys with values
+    thrust::sort_by_key(update_keys.begin(), update_keys.end(), update_values.begin());
+    cErr(cudaDeviceSynchronize());
+
+    // step2: get leaf node of each update (execute del and mod)
+    DEV_VEC_SIZE update_nodes(update_keys.size());
+    cErr(cudaDeviceSynchronize());
+    locate_leaf_batch(RAW_PTR(gpma.keys), RAW_PTR(gpma.values), gpma.keys.size(), gpma.segment_length, gpma.tree_height,
+                      RAW_PTR(update_keys), RAW_PTR(update_values), update_keys.size(), RAW_PTR(update_nodes));
+    cErr(cudaDeviceSynchronize());
+
+    // step3: extract insertions
+    DEV_VEC_SIZE unique_update_nodes(update_keys.size());
+    DEV_VEC_SIZE update_offset(update_keys.size() + 1);
+    cErr(cudaDeviceSynchronize());
+    SIZE_TYPE update_size = update_nodes.size();
+    SIZE_TYPE unique_node_size = 0;
+    compact_insertions(update_nodes, update_keys, update_values, update_size);
+    compress_insertions_by_node(update_nodes, update_size, unique_update_nodes, update_offset, unique_node_size);
+    cErr(cudaDeviceSynchronize());
+
+    // step4: rebuild for significant update
+    int threshold = 5 * 1000 * 1000;
+    if (update_size >= threshold)
+    {
+        significant_insert(gpma, update_keys, update_values, update_size);
+        return;
+    }
+
+    // step5: rebalance each tree level
+    for (SIZE_TYPE level = 0; level <= gpma.tree_height && update_size; level++)
+    {
+        SIZE_TYPE lower_bound = gpma.lower_element[level];
+        SIZE_TYPE upper_bound = gpma.upper_element[level];
+
+        // re-balance
+        rebalance_batch(level, gpma.segment_length, RAW_PTR(gpma.keys), RAW_PTR(gpma.values), RAW_PTR(update_nodes),
+                        RAW_PTR(update_keys), RAW_PTR(update_values), update_size, RAW_PTR(unique_update_nodes),
+                        RAW_PTR(update_offset), unique_node_size, lower_bound, upper_bound, RAW_PTR(gpma.row_offset));
+
+        // compact
+        compact_insertions(update_nodes, update_keys, update_values, update_size);
+
+        // up level
+        up_level_batch(RAW_PTR(update_nodes), update_size, gpma.segment_length << level);
+
+        // re-compress
+        compress_insertions_by_node(update_nodes, update_size, unique_update_nodes, update_offset,
+                                    unique_node_size);
+    }
+
+    // step6: rebalance the root node if necessary
+    if (update_size > 0)
+    {
+        resize_gpma(gpma, update_keys, update_values, update_size);
+
+        SIZE_TYPE level = gpma.tree_height;
+        SIZE_TYPE lower_bound = gpma.lower_element[level];
+        SIZE_TYPE upper_bound = gpma.upper_element[level];
+
+        // re-balance
+        cErr(cudaDeviceSynchronize());
+        rebalance_batch(level, gpma.segment_length, RAW_PTR(gpma.keys), RAW_PTR(gpma.values), RAW_PTR(update_nodes),
+                        RAW_PTR(update_keys), RAW_PTR(update_values), update_size, RAW_PTR(unique_update_nodes),
+                        RAW_PTR(update_offset), unique_node_size, lower_bound, upper_bound, RAW_PTR(gpma.row_offset));
+    }
+
+    cErr(cudaDeviceSynchronize());
+}
+
+__host__ void build_gpma(GPMA &gpma)
+{
+    // we are creating the smallest possible GPMA tree
+    // height = 1
+    // levels = 2
+    // leaf segment length = 2
+    gpma.keys.resize(4, KEY_NONE);
+    gpma.values.resize(4);
+
+    cErr(cudaDeviceSynchronize());
+
+    gpma.segment_length = 2;
+    gpma.tree_height = 1;
+
+    // the minimal tree structure has 2 levels with 4 elements' space, and the leaf segment's length is 2
+    // put two MAX_KEY to keep minimal valid structure
+    gpma.keys[0] = gpma.keys[2] = KEY_MAX;
+    gpma.values[0] = gpma.values[2] = 1;
+
+    recalculate_density(gpma);
+}
+
+template <typename T>
+struct col_idx_none
+{
+    typedef T argument_type;
+    typedef T result_type;
+    __host__ __device__
+        T
+        operator()(const T &x) const
+    {
+        return (x << 32) + COL_IDX_NONE;
+    }
+};
+__host__ void init_gpma(GPMA &gpma, SIZE_TYPE row_num)
+{
+    // gpma     (GPMA)          : The GPMA object whose CSR arrays are to initialised
+    // row_num  (unsized int)   : Number of total nodes in the graph
+
+    // initialising the row_offset vector with all 0 value
+    gpma.row_num = row_num;
+    gpma.row_offset.resize(row_num + 1, 0);
+
+    // initialising in_degree and out_degree arrays
+    // with all zero values
+    gpma.in_degree.resize(row_num, 0);
+    gpma.out_degree.resize(row_num, 0);
+
+    // creates a device vector of size row_num
+    DEV_VEC_KEY row_wall(row_num);
+
+    // creates a device vector of size row_num all
+    // initialised to the value 1
+    DEV_VEC_VALUE tmp_value(row_num, 1);
+    cErr(cudaDeviceSynchronize());
+
+    thrust::tabulate(row_wall.begin(), row_wall.end(), col_idx_none<KEY_TYPE>());
+    build_gpma(gpma);
+    cErr(cudaDeviceSynchronize());
+    update_gpma(gpma, row_wall, tmp_value);
+    cErr(cudaDeviceSynchronize());
+}
+
+///////////////////////////////////////////////////////////////////////////////////////////////////
+
+void print_gpma_info(GPMA &gpma, int node)
+{
+    // Pretty prints the row_offset range and column_indices values
+    // for the given node
+    //
+    // row_offset:  Row offset vector
+    // col_indices: Column indices PMA
+    // node:        Node whose row_offset and column_indices is to be printed
+
+    // getting the (beg,end) range of the node
+
+    thrust::host_vector<SIZE_TYPE> row_offset = gpma.row_offset;
+    thrust::host_vector<KEY_TYPE> col_indices = gpma.keys;
+    thrust::host_vector<VALUE_TYPE> edge_values = gpma.values;
+
+    unsigned int beg = row_offset[node];
+    unsigned int end = row_offset[node + 1];
+    int print_count = 0;
+
+    std::cout << "\n🌟 Column Indices for node: " << node << "(" << beg << ", " << end << ")"
+              << "\n\n";
+    std::cout << std::setw(6);
+
+    for (int i = beg; i < end; ++i)
+    {
+        KEY_TYPE mask = (KEY_TYPE)node << 32;
+        unsigned int dst = (col_indices[i] - mask);
+        VALUE_TYPE val = edge_values[i];
+
+        if (dst != COL_IDX_NONE && val != VALUE_NONE)
+        {
+            std::cout << dst << "(" << val << ")" << std::setw(6);
+        }
+        else
+        {
+            std::cout << "-" << std::setw(6);
+        }
+        print_count += 1;
+
+        if (print_count > 10)
+        {
+            printf("\n");
+            print_count = 0;
+        }
+    }
+    printf("\n\n");
+}
+
+void load_data(const char *file_path, thrust::host_vector<int> &host_x, thrust::host_vector<int> &host_y,
+               int &node_size, int &edge_size)
+{
+
+    FILE *fp;
+    fp = fopen(file_path, "r");
+    if (not fp)
+    {
+        printf("Open graph file failed.\n");
+        exit(0);
+    }
+
+    // gets the number of nodes and edges from the first line
+    // [num_nodes] [num_edges]
+    fscanf(fp, "%d %d", &node_size, &edge_size);
+    printf("node_num: %d, edge_num: %d\n", node_size, edge_size);
+
+    // dynamic resizing to fit all edges
+    host_x.resize(edge_size);
+    host_y.resize(edge_size);
+
+    // loads the src and dst nodes into x_host
+    // and y_host respectively
+    for (int i = 0; i < edge_size; i++)
+    {
+        int x, y;
+
+        // [src_node] [dst_node]
+        fscanf(fp, "%d %d", &x, &y);
+        host_x[i] = x;
+        host_y[i] = y;
+    }
+
+    printf("Graph file is loaded.\n");
+    fclose(fp);
+}
+
+void load_graph(GPMA &gpma, const char *file_path)
+{
+    // sets the GPU Malloc Heap Size to 1GB
+    // we could change the limit accordingly
+    // cudaDeviceSetLimit(cudaLimitMallocHeapSize, 1024ll * 1024);
+
+    // controls the maximum nesting depth of a grid at which
+    // a thread can safely call cudaDeviceSynchronize()
+    // cudaDeviceSetLimit(cudaLimitDevRuntimeSyncDepth, 5);
+
+    // we are defining two host_vectors on the CPU
+    // host_x will contain the src nodes and
+    // host_y will contain the destination nodes of the edge
+    thrust::host_vector<int> host_x;
+    thrust::host_vector<int> host_y;
+
+    int node_size; // num of nodes
+    int edge_size; // num of edges
+
+    // loads the graph edges
+    load_data(file_path, host_x, host_y, node_size, edge_size);
+
+    // we are starting out with half the number of edges
+    // as the initial graph. After this is done, the other
+    // half of edges are added dynamically to the graph in
+    // batches of 100 edges
+    int half = edge_size / 2;
+
+    // h_base_key is a host_vector to maintain a unique ID
+    // for each edge. KEY_TYPE = unsigned long long.
+    // The unique ID is a 64-bit number, where the first 32 bits
+    // represent the source node and the last 32-bits the
+    // destination node.
+    thrust::host_vector<KEY_TYPE> h_base_keys(half);
+
+    // looping over half the edges to create the unique edge ID.
+    // It shifts the source node ID (64-bit after casting) by 32 bits
+    // to the left and the 32-bit destination node ID is added
+    // to the 32-bits on the left of source node ID
+    for (int i = 0; i < half; i++)
+        h_base_keys[i] = ((KEY_TYPE)host_x[i] << 32) + host_y[i];
+
+    // device_vector<unsigned long long>. The unique edge IDs are
+    // copied into the GPU device through a device_vector
+    DEV_VEC_KEY base_keys = h_base_keys;
+
+    // device_vector<double>. device_vector of size
+    // half = num_edges/2. Contains the value of the edges.
+    // Initialised to 1 and stored in GPU device
+    DEV_VEC_VALUE base_values(half, 1);
+
+    // the base keys and values are being transferred to
+    // GPU device. Host is waiting for that task to complete
+    // before continuing with the rest.
+    cudaDeviceSynchronize();
+
+    init_gpma(gpma, node_size); //
+    cudaDeviceSynchronize();
+
+    update_gpma(gpma, base_keys, base_values);
+    cudaDeviceSynchronize();
+
+    printf("Graph is updated.\n");
+}
+
+void edge_update_list(GPMA &gpma, std::vector<std::tuple<int, int>> edge_list, bool is_delete = false, bool is_reverse_edge = false)
+{
+    // NOTE:: Should we set these limits every single time?
+    // cudaDeviceSetLimit(cudaLimitMallocHeapSize, 1024ll * 1024);
+    // cudaDeviceSetLimit(cudaLimitDevRuntimeSyncDepth, 5);
+
+    thrust::host_vector<int> host_src;
+    thrust::host_vector<int> host_dst;
+
+    int edge_count = 0;
+
+    // iterating through a vector of tuples
+    // each tuple is of the form (src_node, dst_node)
+
+    for (auto &edge : edge_list)
+    {
+
+        int src = (is_reverse_edge == true) ? std::get<1>(edge) : std::get<0>(edge);
+        int dst = (is_reverse_edge == true) ? std::get<0>(edge) : std::get<1>(edge);
+
+        // updating the in-degree and out-degree of
+        // destination and source node respectively
+        if (is_delete)
+        {
+            gpma.in_degree[dst] -= 1;
+            gpma.out_degree[src] -= 1;
+        }
+        else
+        {
+            gpma.in_degree[dst] += 1;
+            gpma.out_degree[src] += 1;
+        }
+
+        host_src.push_back(src);
+        host_dst.push_back(dst);
+
+        ++edge_count;
+    }
+
+    gpma.edge_count = is_delete ? gpma.edge_count - edge_count : gpma.edge_count + edge_count;
+
+    thrust::host_vector<KEY_TYPE> h_base_keys(edge_count);
+
+    for (int i = 0; i < edge_count; i++)
+        h_base_keys[i] = ((KEY_TYPE)host_src[i] << 32) + host_dst[i];
+
+    DEV_VEC_KEY base_keys = h_base_keys;
+    cudaDeviceSynchronize();
+
+    DEV_VEC_VALUE base_values(edge_count);
+
+    if (is_delete)
+        thrust::fill(base_values.begin(), base_values.end(), VALUE_NONE);
+    else
+        thrust::fill(base_values.begin(), base_values.end(), 1);
+
+    cudaDeviceSynchronize();
+
+    update_gpma(gpma, base_keys, base_values);
+    cudaDeviceSynchronize();
+}
+
+void label_edges(GPMA &gpma)
+{
+    int edge_label_counter = 1;
+
+    thrust::host_vector<SIZE_TYPE> h_row_offset = gpma.row_offset;
+    thrust::host_vector<KEY_TYPE> h_keys = gpma.keys;
+    thrust::host_vector<VALUE_TYPE> h_values = gpma.values;
+
+    for (int node = 0; node < h_row_offset.size() - 1; ++node)
+    {
+        unsigned int beg = h_row_offset[node];
+        unsigned int end = h_row_offset[node + 1];
+
+        for (int i = beg; i < end; ++i)
+        {
+            KEY_TYPE mask = (KEY_TYPE)node << 32;
+            unsigned int dst = (h_keys[i] - mask);
+            if (dst != COL_IDX_NONE && h_values[i] != VALUE_NONE)
+            {
+                h_values[i] = edge_label_counter;
+                edge_label_counter += 1;
+            }
+        }
+    }
+
+    gpma.row_offset = h_row_offset;
+    gpma.keys = h_keys;
+    gpma.values = h_values;
+}
+
+void copy_label_edges(GPMA &gpma, GPMA &ref_gpma)
+{
+    int edge_counter = 0;
+
+    // creating the host vector version for keys
+    // and values for the new gpma
+    thrust::host_vector<KEY_TYPE> h_new_keys(ref_gpma.edge_count);
+    thrust::host_vector<VALUE_TYPE> h_new_values(ref_gpma.edge_count);
+
+    // creating the host vectors for the CSR arrays of the
+    // reference GPMA from it's device vectors
+    thrust::host_vector<SIZE_TYPE> h_ref_row_offset = ref_gpma.row_offset;
+    thrust::host_vector<KEY_TYPE> h_ref_keys = ref_gpma.keys;
+    thrust::host_vector<VALUE_TYPE> h_ref_values = ref_gpma.values;
+
+    DEV_VEC_KEY d_new_keys(ref_gpma.edge_count);
+    DEV_VEC_VALUE d_new_values(ref_gpma.edge_count);
+
+    for (int node = 0; node < h_ref_row_offset.size() - 1; ++node)
+    {
+        unsigned int beg = h_ref_row_offset[node];
+        unsigned int end = h_ref_row_offset[node + 1];
+
+        for (int i = beg; i < end; ++i)
+        {
+            KEY_TYPE mask = (KEY_TYPE)node << 32;
+            unsigned int dst = (h_ref_keys[i] - mask);
+            if (dst != COL_IDX_NONE && h_ref_values[i] != VALUE_NONE)
+            {
+                h_new_keys[edge_counter] = (h_ref_keys[i] << 32) + node;
+                h_new_values[edge_counter] = h_ref_values[i];
+                edge_counter += 1;
+            }
+        }
+    }
+
+    d_new_keys = h_new_keys;
+    d_new_values = h_new_values;
+
+    // NOTE: Verify if sorting is really required or not
+    thrust::sort_by_key(d_new_keys.begin(), d_new_keys.end(), d_new_values.begin());
+    cudaDeviceSynchronize();
+
+    locate_leaf_batch(RAW_PTR(gpma.keys), RAW_PTR(gpma.values), gpma.keys.size(), gpma.segment_length, gpma.tree_height,
+                      RAW_PTR(d_new_keys), RAW_PTR(d_new_values), d_new_keys.size(), NULL, false);
+
+    cudaDeviceSynchronize();
+}
+
+void build_reverse_gpma(GPMA &gpma, GPMA &ref_gpma)
+{
+    // This funtion expects an empty initialized GPMA
+    // as first param and the reference GPMA from which
+    // edges are to be added as the second param
+
+    int edge_counter = 0;
+
+    // creating host_vectors for the new GPMA
+    thrust::host_vector<KEY_TYPE> h_new_keys(ref_gpma.edge_count);
+    thrust::host_vector<VALUE_TYPE> h_new_values(ref_gpma.edge_count);
+
+    // creating the host vectors for the CSR arrays of the
+    // reference GPMA from it's device vectors
+    thrust::host_vector<SIZE_TYPE> h_ref_row_offset = ref_gpma.row_offset;
+    thrust::host_vector<KEY_TYPE> h_ref_keys = ref_gpma.keys;
+    thrust::host_vector<VALUE_TYPE> h_ref_values = ref_gpma.values;
+
+    DEV_VEC_KEY d_new_keys(ref_gpma.edge_count);
+    DEV_VEC_VALUE d_new_values(ref_gpma.edge_count);
+
+    for (int node = 0; node < h_ref_row_offset.size() - 1; ++node)
+    {
+        unsigned int beg = h_ref_row_offset[node];
+        unsigned int end = h_ref_row_offset[node + 1];
+
+        for (int i = beg; i < end; ++i)
+        {
+            KEY_TYPE mask = (KEY_TYPE)node << 32;
+            unsigned int dst = (h_ref_keys[i] - mask);
+            if (dst != COL_IDX_NONE && h_ref_values[i] != VALUE_NONE)
+            {
+                h_new_keys[edge_counter] = (h_ref_keys[i] << 32) + node;
+                h_new_values[edge_counter] = h_ref_values[i];
+                edge_counter += 1;
+            }
+        }
+    }
+
+    d_new_keys = h_new_keys;
+    d_new_values = h_new_values;
+
+    update_gpma(gpma, d_new_keys, d_new_values);
+    cudaDeviceSynchronize();
+}
+
+std::tuple<std::size_t, std::size_t, std::size_t> get_csr_ptrs(GPMA &gpma)
+{
+    std::tuple<std::size_t, std::size_t, std::size_t> t;
+    std::get<0>(t) = (std::size_t)RAW_PTR(gpma.row_offset);
+    std::get<1>(t) = (std::size_t)RAW_PTR(gpma.keys);
+    std::get<2>(t) = (std::size_t)RAW_PTR(gpma.values);
+    return t;
+}
+
+std::tuple<int, int> get_graph_attr(GPMA &gpma)
+{
+    std::tuple<int, int> t;
+
+    // Since row_offset has size num_nodes + 1
+    std::get<0>(t) = gpma.row_offset.size() - 1;
+    std::get<1>(t) = gpma.edge_count;
+    return t;
+}
+
+PYBIND11_MODULE(gpma, m)
+{
+    m.doc() = "CPython module for GPMA"; // optional module docstring
+
+    m.def("init_gpma", &init_gpma, "Initialises the CSR arrays using GPMA");
+    m.def("print_gpma_info", &print_gpma_info, "Prints row_offset and col_indices for a given node");
+    m.def("load_graph", &load_graph, "Loads a graph data into a GPMA");
+    m.def("edge_update_list", &edge_update_list, "Updates the GPMA by adding/deleting edges from the edge list", py::arg("gpma"), py::arg("edge_list"), py::arg("is_delete") = false, py::arg("is_reverse_edge") = false);
+    m.def("label_edges", &label_edges, "Creates edge labels for the current GPMA");
+    m.def("copy_label_edges", &copy_label_edges, "Label edges of a GPMA based on another GPMA");
+    m.def("build_reverse_gpma", &build_reverse_gpma, "Builds the reverse GPMA based on another GPMA");
+    m.def("get_csr_ptrs", &get_csr_ptrs, "Returns the pointers to row_offset, col_indices and edge_ids");
+    m.def("get_graph_attr", &get_graph_attr, "Returns the [num_nodes, num_edges]");
+
+    py::class_<GPMA>(m, "GPMA")
+        .def(py::init<>())
+        .def("get_size", &GPMA::get_size)
+        .def_readwrite("keys", &GPMA::keys)
+        .def_readwrite("values", &GPMA::values)
+        .def_readwrite("segment_length", &GPMA::segment_length)
+        .def_readwrite("tree_height", &GPMA::tree_height)
+        .def_readwrite("density_lower_thres_leaf", &GPMA::density_lower_thres_leaf)
+        .def_readwrite("density_lower_thres_root", &GPMA::density_lower_thres_root)
+        .def_readwrite("density_upper_thres_root", &GPMA::density_upper_thres_root)
+        .def_readwrite("density_upper_thres_leaf", &GPMA::density_upper_thres_leaf)
+        .def_readwrite("lower_element", &GPMA::lower_element)
+        .def_readwrite("upper_element", &GPMA::upper_element)
+        .def_readwrite("row_num", &GPMA::row_num)
+        .def_readwrite("row_offset", &GPMA::row_offset)
+        .def_readwrite("in_degree", &GPMA::in_degree)
+        .def_readwrite("out_degree", &GPMA::out_degree)
+        .def("__copy__", [](const GPMA &self)
+             { return GPMA(self); })
+        .def(
+            "__deepcopy__", [](const GPMA &self, py::dict)
+            { return GPMA(self); },
+            "memo"_a);
+}
+
+// Command used:
+// /usr/local/cuda-11.7/bin/nvcc $(python3 -m pybind11 --includes) -shared -rdc=true --compiler-options '-fPIC'  -D__CDPRT_SUPPRESS_SYNC_DEPRECATION_WARNING -o gpma.so gpma.cu
\ No newline at end of file
diff --git a/seastar/graph/dynamic/naive/NaiveGraph.py b/seastar/graph/dynamic/naive/NaiveGraph.py
new file mode 100644
index 0000000..d982e09
--- /dev/null
+++ b/seastar/graph/dynamic/naive/NaiveGraph.py
@@ -0,0 +1,77 @@
+import copy
+import numpy as np
+from rich import inspect
+
+from seastar.graph.dynamic.DynamicGraph import DynamicGraph
+from seastar.graph.static.csr import CSR
+from collections import deque
+
+class NaiveGraph(DynamicGraph):
+    def __init__(self, edge_list):
+        super().__init__(edge_list)
+        
+        self.graph_stack = deque()
+        self.edge_list = edge_list
+        
+        self._forward_graph = CSR(edge_list[0], self.graph_updates["0"]["num_nodes"], is_edge_reverse=True)
+        self._forward_graph.label_edges()
+
+        self._graph_stack_push(self._forward_graph)
+        self._get_graph_csr_ptrs()
+    
+    def _graph_stack_push(self, elem):
+        self.graph_stack.append(elem)
+    
+    def _graph_stack_pop(self):
+        elem = self.graph_stack[-1]
+        return elem
+        
+    def graph_type(self):
+        return "csr"
+        
+    def in_degrees(self):
+        return np.array(self._forward_graph.out_degrees, dtype='int32')
+    
+    def out_degrees(self):
+        return np.array(self._forward_graph.in_degrees, dtype='int32')
+    
+    def _get_graph_csr_ptrs(self):
+        fwd_csr_ptrs = self._forward_graph.get_csr_ptrs()
+        self.fwd_row_offset_ptr = fwd_csr_ptrs[0]
+        self.fwd_column_indices_ptr = fwd_csr_ptrs[1]
+        self.fwd_eids_ptr = fwd_csr_ptrs[2]
+        
+        if self._backward_graph is not None:
+            bwd_csr_ptrs = self._backward_graph.get_csr_ptrs()
+            self.bwd_row_offset_ptr = bwd_csr_ptrs[0]
+            self.bwd_column_indices_ptr = bwd_csr_ptrs[1]
+            self.bwd_eids_ptr = bwd_csr_ptrs[2]
+    
+    def _update_graph_forward(self):
+        ''' Updates the current base graph to the next timestamp
+        '''
+
+        if str(self.current_timestamp + 1) not in self.graph_updates:
+            raise Exception("⏰ Invalid timestamp during SeastarGraph.update_graph_forward()")
+
+        self._forward_graph = CSR(self.edge_list[self.current_timestamp], self.graph_updates[str(self.current_timestamp+1)]["num_nodes"], is_edge_reverse=True)
+        self._forward_graph.label_edges()
+        self._graph_stack_push(self._forward_graph)
+        self._get_graph_csr_ptrs()
+        
+    def _init_reverse_graph(self):
+        ''' Generates the reverse of the base graph'''
+
+        fwd_graph = self._graph_stack_pop()
+        self._backward_graph = CSR(self.edge_list[self.current_timestamp], self.graph_updates[str(self.current_timestamp)]["num_nodes"])
+        self._backward_graph.copy_label_edges(fwd_graph) 
+        self._get_graph_csr_ptrs()
+        
+    def _update_graph_backward(self):
+        if self.current_timestamp < 0:
+            raise Exception("⏰ Invalid timestamp during SeastarGraph.update_graph_backward()")
+        
+        fwd_graph = self._graph_stack_pop()
+        self._backward_graph = CSR(self.edge_list[self.current_timestamp - 1], self.graph_updates[str(self.current_timestamp - 1)]["num_nodes"])
+        self._backward_graph.copy_label_edges(fwd_graph) 
+        self._get_graph_csr_ptrs()
diff --git a/seastar/graph/dynamic/naive/__init__.py b/seastar/graph/dynamic/naive/__init__.py
new file mode 100644
index 0000000..e69de29
diff --git a/seastar/graph/dynamic/pcsr/PCSRGraph.py b/seastar/graph/dynamic/pcsr/PCSRGraph.py
new file mode 100644
index 0000000..1912794
--- /dev/null
+++ b/seastar/graph/dynamic/pcsr/PCSRGraph.py
@@ -0,0 +1,106 @@
+import copy
+import numpy as np
+from rich import inspect
+
+from seastar.graph.dynamic.DynamicGraph import DynamicGraph
+from seastar.graph.dynamic.pcsr.pcsr import PCSR, copy_label_edges, build_reverse_pcsr
+
+class PCSRGraph(DynamicGraph):
+    def __init__(self, edge_list):
+        super().__init__(edge_list)
+        
+        self._forward_graph = PCSR(self.max_num_nodes)
+        self._backward_graph = PCSR(self.max_num_nodes)
+        self._forward_graph.edge_update_list(self.graph_updates["0"]["add"],is_reverse_edge=True)
+
+        self._forward_graph.label_edges()
+        self._get_graph_csr_ptrs()
+        
+    def graph_type(self):
+        return "pcsr"
+        
+    def in_degrees(self):
+        return np.array([node.num_neighbors for node in self._forward_graph.nodes], dtype='int32')
+    
+    def out_degrees(self):
+        return np.array([node.in_degree for node in self._forward_graph.nodes], dtype='int32')
+    
+    def _update_reverse_graph_cache(self):
+        # saving reverse base graph in cache
+        self.graph_cache['reverse'] = copy.deepcopy(self._backward_graph)
+            
+    def _get_reverse_cached_graph(self):
+        return copy.deepcopy(self.graph_cache['reverse'])
+    
+    def _get_graph_csr_ptrs(self):
+        forward_csr_ptrs = self._forward_graph.get_csr_ptrs()
+        backward_csr_ptrs = self._backward_graph.get_csr_ptrs()
+        
+        self.fwd_row_offset_ptr = forward_csr_ptrs[0]
+        self.fwd_column_indices_ptr = forward_csr_ptrs[1]
+        self.fwd_eids_ptr = forward_csr_ptrs[2]
+        
+        self.bwd_row_offset_ptr = backward_csr_ptrs[0]
+        self.bwd_column_indices_ptr = backward_csr_ptrs[1]
+        self.bwd_eids_ptr = backward_csr_ptrs[2]
+
+    # def _get_graph_attributes(self):
+    #     if not self._is_backprop_state:
+    #         graph_attr = self._forward_graph.get_graph_attr()
+    #     else:
+    #         graph_attr = self._backward_graph.get_graph_attr()
+        
+    #     self.num_nodes = graph_attr[0]
+    #     self.num_edges = graph_attr[1]
+    
+    def _update_graph_forward(self):
+        ''' Updates the current base graph to the next timestamp
+        '''
+
+        if str(self.current_timestamp + 1) not in self.graph_updates:
+            raise Exception("⏰ Invalid timestamp during SeastarGraph.update_graph_forward()")
+        
+        graph_additions = self.graph_updates[str(self.current_timestamp + 1)]["add"]
+        graph_deletions = self.graph_updates[str(self.current_timestamp + 1)]["delete"]
+
+        self._forward_graph.edge_update_list(graph_additions,is_reverse_edge=True)
+        self._forward_graph.edge_update_list(graph_deletions,is_delete=True,is_reverse_edge=True)
+
+        self._forward_graph.label_edges()
+        self._get_graph_csr_ptrs()
+        # self._get_graph_attributes()
+        
+    def _init_reverse_graph(self):
+        ''' Generates the reverse of the base graph'''
+
+        # checking if the reverse base graph exists in the cache
+        # we can load it from there instead of building it each time
+        if 'reverse' in self.graph_cache:
+            self._backward_graph = self._get_reverse_cached_graph()
+        else:
+            build_reverse_pcsr(self._backward_graph, self._forward_graph)
+
+            # storing the reverse base graph in cache after building
+            # it for the first time
+            self._update_reverse_graph_cache()
+
+        self._get_graph_csr_ptrs()
+        # self._get_graph_attributes()
+        
+    def _update_graph_backward(self):
+        if self.current_timestamp < 0:
+            raise Exception("⏰ Invalid timestamp during SeastarGraph.update_graph_backward()")
+        
+        graph_additions = self.graph_updates[str(self.current_timestamp)]["delete"]
+        graph_deletions = self.graph_updates[str(self.current_timestamp)]["add"]
+
+        self._backward_graph.edge_update_list(graph_additions)
+        self._forward_graph.edge_update_list(graph_additions,is_reverse_edge=True)
+        
+        self._backward_graph.edge_update_list(graph_deletions,is_delete=True)
+        self._forward_graph.edge_update_list(graph_deletions, is_delete=True, is_reverse_edge=True)
+
+        self._forward_graph.label_edges()
+        copy_label_edges(self._backward_graph, self._forward_graph)
+        self._get_graph_csr_ptrs()
+        # self._get_graph_attributes()
diff --git a/seastar/graph/dynamic/pcsr/README.md b/seastar/graph/dynamic/pcsr/README.md
new file mode 100644
index 0000000..986fdb6
--- /dev/null
+++ b/seastar/graph/dynamic/pcsr/README.md
@@ -0,0 +1,9 @@
+# PCSR Python implementation with C++
+
+## Building the shared object
+
+Run the following command while in this same directory to build the `pcsr.so` file.
+
+```
+/usr/local/cuda-11.7/bin/nvcc $(python3 -m pybind11 --includes) -w -shared -rdc=true --compiler-options '-fPIC' -o pcsr.so pcsr.cu
+```
\ No newline at end of file
diff --git a/seastar/graph/dynamic/pcsr/__init__.py b/seastar/graph/dynamic/pcsr/__init__.py
new file mode 100644
index 0000000..e69de29
diff --git a/seastar/graph/dynamic/pcsr/pcsr.cu b/seastar/graph/dynamic/pcsr/pcsr.cu
new file mode 100644
index 0000000..1f7e0c0
--- /dev/null
+++ b/seastar/graph/dynamic/pcsr/pcsr.cu
@@ -0,0 +1,978 @@
+#include <pybind11/pybind11.h>
+#include <pybind11/stl.h>
+
+#include <algorithm>
+#include <queue>
+#include <stdint.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <tuple>
+#include <vector>
+#include <string>
+#include <cstring>
+#include <iostream>
+
+#include <thrust/device_vector.h>
+#include <thrust/host_vector.h>
+#include <thrust/remove.h>
+#include <thrust/sort.h>
+#include <cub/cub.cuh>
+
+namespace py = pybind11;
+using namespace pybind11::literals;
+using namespace std;
+
+typedef thrust::device_vector<int> DEV_VEC;
+#define RAW_PTR(x) thrust::raw_pointer_cast((x).data())
+
+////////////////////////////////////////////////////////////////////////////
+
+typedef struct _node
+{
+    // beginning and end of the associated region in the edge list
+    uint32_t beginning;     // deleted = max int
+    uint32_t end;           // end pointer is exclusive
+    uint32_t num_neighbors; // number of edges with this node as source
+    uint32_t in_degree;     // in-degree of a node -  number of edges going into the node
+
+    _node(int beg = 0, int _end = 0, int num_neigh = 0, int in_deg = 0)
+    {
+        beginning = beg;
+        end = _end;
+        num_neighbors = num_neigh;
+        in_degree = in_deg;
+    }
+} node_t;
+
+typedef struct _edge
+{
+    uint32_t dest;  // destination of this edge in the graph, MAX_INT if this is a sentinel
+    uint32_t value; // edge value of zero means it a null since we don't store 0 edges
+
+    _edge(int _dest = 0, int _value = 0)
+    {
+        dest = _dest;
+        value = _value;
+    }
+} edge_t;
+
+typedef struct edge_list
+{
+    int N;
+    int H;
+    int logN;
+    vector<edge_t> items;
+
+    edge_list()
+    {
+        N = 0;
+        H = 0;
+        logN = 0;
+
+        vector<edge_t> temp(0, 0);
+        items = temp;
+    }
+} edge_list_t;
+
+typedef struct _pair_int
+{
+    int x; // length in array
+    int y; // depth
+
+    _pair_int(int _x = 0, int _y = 0)
+    {
+        x = _x;
+        y = _y;
+    }
+} pair_int;
+
+typedef struct _pair_double
+{
+    double x;
+    double y;
+
+    _pair_double(double _x = 0, double _y = 0)
+    {
+        x = _x;
+        y = _y;
+    }
+} pair_double;
+
+////////////////////////////////////////////////////////////////////////////
+
+static inline int bsf_word(int word)
+{
+    int result;
+    __asm__ volatile("bsf %1, %0"
+                     : "=r"(result)
+                     : "r"(word));
+    return result;
+}
+
+static inline int bsr_word(int word)
+{
+    int result;
+    __asm__ volatile("bsr %1, %0"
+                     : "=r"(result)
+                     : "r"(word));
+    return result;
+}
+
+// given index, return the starting index of the leaf it is in
+int find_leaf(edge_list_t *list, int index)
+{
+    return (index / list->logN) * list->logN;
+}
+
+bool is_null(edge_t e) { return e.value == 0; }
+
+bool is_sentinel(edge_t e)
+{
+    return e.dest == UINT32_MAX || e.value == UINT32_MAX;
+}
+
+uint32_t binary_search(edge_list_t *list, edge_t *elem, uint32_t start,
+                       uint32_t end)
+{
+    while (start + 1 < end)
+    {
+        uint32_t mid = (start + end) / 2;
+
+        edge_t item = list->items[mid];
+        uint32_t change = 1;
+        uint32_t check = mid;
+
+        bool flag = true;
+        while (is_null(item) && flag)
+        {
+            flag = false;
+            check = mid + change;
+            if (check < end)
+            {
+                flag = true;
+                if (check <= end)
+                {
+                    item = list->items[check];
+                    if (!is_null(item))
+                    {
+                        break;
+                    }
+                    else if (check == end)
+                    {
+                        break;
+                    }
+                }
+            }
+            check = mid - change;
+            if (check >= start)
+            {
+                flag = true;
+                item = list->items[check];
+            }
+            change++;
+        }
+
+        if (is_null(item) || start == check || end == check)
+        {
+            if (!is_null(item) && start == check && elem->dest <= item.dest)
+            {
+                return check;
+            }
+            return mid;
+        }
+
+        // if we found it, return
+        if (elem->dest == item.dest)
+        {
+            return check;
+        }
+        else if (elem->dest < item.dest)
+        {
+            end =
+                check; // if the searched for item is less than current item, set end
+        }
+        else
+        {
+            start = check;
+            // otherwise, searched for item is more than current and we set start
+        }
+    }
+    if (end < start)
+    {
+        start = end;
+    }
+    // handling the case where there is one element left
+    // if you are leq, return start (index where elt is)
+    // otherwise, return end (no element greater than you in the range)
+    // printf("start = %d, end = %d, n = %d\n", start,end, list->N);
+    if (elem->dest <= list->items[start].dest && !is_null(list->items[start]))
+    {
+        return start;
+    }
+    return end;
+}
+
+// get density of a node
+double get_density(edge_list_t *list, int index, int len)
+{
+    int full = 0;
+    for (int i = index; i < index + len; i++)
+    {
+        full += (!is_null(list->items[i]));
+    }
+    double full_d = (double)full;
+    return full_d / len;
+}
+
+int find_node(int index, int len) { return (index / len) * len; }
+
+pair_double density_bound(edge_list_t *list, int depth)
+{
+    pair_double pair;
+
+    // between 1/4 and 1/2
+    // pair.x = 1.0/2.0 - (( .25*depth)/list->H);
+    // between 1/8 and 1/4
+    pair.x = 1.0 / 4.0 - ((.125 * depth) / list->H);
+    pair.y = 3.0 / 4.0 + ((.25 * depth) / list->H);
+    return pair;
+}
+
+bool edge_equals(edge_t e1, edge_t e2)
+{
+    return e1.dest == e2.dest && e1.value == e2.value;
+}
+
+uint32_t find_elem_pointer(edge_list_t *list, uint32_t index, edge_t elem)
+{
+    edge_t item = list->items[index];
+    while (!edge_equals(item, elem))
+    {
+        item = list->items[++index];
+    }
+    return index;
+}
+
+////////////////////////////////////////////////////////////////////////////
+
+class PCSR
+{
+public:
+    // data members
+    std::vector<node_t> nodes;
+    edge_list_t edges;
+    uint32_t edge_count = 0;
+
+    // member functions
+    PCSR(uint32_t init_n);
+    void init_graph(std::string graph_file_path);
+
+    void add_node();
+    void add_edge(uint32_t src, uint32_t dest, uint32_t value);
+    void add_edge_update(uint32_t src, uint32_t dest, uint32_t value);
+    void edge_update_list(std::vector<std::tuple<uint32_t, uint32_t>> edge_list, bool is_delete, bool is_reverse_edge);
+    void delete_edge(uint32_t src, uint32_t dest);
+    uint64_t get_n();
+    vector<tuple<uint32_t, uint32_t, uint32_t>> get_edges();
+    void print_graph();
+    void print_array();
+    std::tuple<std::size_t, std::size_t, std::size_t> get_csr_ptrs();
+    void label_edges();
+    uint32_t find_edge_id(uint32_t src, uint32_t dest);
+    std::tuple<uint32_t, uint32_t> get_graph_attr();
+
+    uint32_t insert(uint32_t index, edge_t elem, uint32_t src);
+    void double_list();
+    int slide_right(int index);
+    void slide_left(int index);
+    void fix_sentinel(int32_t node_index, int in);
+    void redistribute(int index, int len);
+};
+
+////////////////////////////////////////////////////////////////////////////
+
+// removed default argument of 0
+PCSR::PCSR(uint32_t init_n)
+{
+    if (init_n != 0)
+    {
+        edges.N = 2 << bsr_word(init_n);
+        edges.logN = (1 << bsr_word(bsr_word(edges.N) + 1));
+        edges.H = bsr_word(edges.N / edges.logN);
+
+        for (int i = 0; i < edges.N; i++)
+        {
+            edge_t new_edge(0, 0);
+            edges.items.push_back(new_edge);
+        }
+
+        for (int i = 0; i < init_n; i++)
+        {
+            add_node();
+        }
+    }
+}
+
+void PCSR::init_graph(std::string graph_file_path)
+{
+
+    FILE *fp;
+    fp = fopen(graph_file_path.c_str(), "r");
+    if (!fp)
+    {
+        printf("Open file failed.\n");
+        exit(-1);
+    }
+
+    // loading the number of nodes and edges present in
+    // the graph into the variables node_size and edge_size respectively
+    int node_size, edge_size;
+    fscanf(fp, "%d %d", &node_size, &edge_size);
+
+    edges.N = 2 << bsr_word(node_size);
+    edges.logN = (1 << bsr_word(bsr_word(edges.N) + 1));
+    edges.H = bsr_word(edges.N / edges.logN);
+
+    for (int i = 0; i < edges.N; i++)
+    {
+        edge_t new_edge(0, 0);
+        edges.items.push_back(new_edge);
+    }
+
+    for (int i = 0; i < node_size; i++)
+    {
+        add_node();
+    }
+
+    for (int i = 0; i < edge_size; i++)
+    {
+        int source, dest, edge_value;
+        fscanf(fp, "%d %d %d", &source, &dest, &edge_value);
+        add_edge(source, dest, edge_value);
+    }
+}
+
+// add a node to the graph
+void PCSR::add_node()
+{
+    node_t node;
+    int len = nodes.size();
+
+    edge_t sentinel;
+    sentinel.dest = UINT32_MAX; // placeholder
+    sentinel.value = len;       // back pointer
+
+    if (len > 0)
+    {
+        node.beginning = nodes[len - 1].end;
+        node.end = node.beginning + 1;
+    }
+    else
+    {
+        node.beginning = 0;
+        node.end = 1;
+        sentinel.value = UINT32_MAX;
+    }
+    node.num_neighbors = 0;
+
+    nodes.push_back(node);
+    insert(node.beginning, sentinel, nodes.size() - 1);
+}
+
+uint32_t PCSR::insert(uint32_t index, edge_t elem, uint32_t src)
+{
+    int node_index = find_leaf(&edges, index);
+    int level = edges.H;
+    int len = edges.logN;
+
+    // always deposit on the left
+    if (is_null(edges.items[index]))
+    {
+        edges.items[index].value = elem.value;
+        edges.items[index].dest = elem.dest;
+    }
+    else
+    {
+        // if the edge already exists in the graph, update its value
+        // do not make another edge
+        // return index of the edge that already exists
+        if (!is_sentinel(elem) && edges.items[index].dest == elem.dest)
+        {
+            edges.items[index].value = elem.value;
+            return index;
+        }
+        if (index == edges.N - 1)
+        {
+            // when adding to the end double then add edge
+            double_list();
+            node_t node = nodes[src];
+            uint32_t loc_to_add =
+                binary_search(&edges, &elem, node.beginning + 1, node.end);
+            return insert(loc_to_add, elem, src);
+        }
+        else
+        {
+            if (slide_right(index) == -1)
+            {
+                index -= 1;
+                slide_left(index);
+            }
+        }
+        edges.items[index].value = elem.value;
+        edges.items[index].dest = elem.dest;
+    }
+
+    double density = get_density(&edges, node_index, len);
+
+    // spill over into next level up, node is completely full.
+    if (density == 1)
+    {
+        node_index = find_node(node_index, len * 2);
+        redistribute(node_index, len * 2);
+    }
+    else
+    {
+        // makes the last slot in a section empty so you can always slide right
+        redistribute(node_index, len);
+    }
+
+    // get density of the leaf you are in
+    pair_double density_b = density_bound(&edges, level);
+    density = get_density(&edges, node_index, len);
+
+    // while density too high, go up the implicit tree
+    // go up to the biggest node above the density bound
+    while (density >= density_b.y)
+    {
+        len *= 2;
+        if (len <= edges.N)
+        {
+            level--;
+            node_index = find_node(node_index, len);
+            density_b = density_bound(&edges, level);
+            density = get_density(&edges, node_index, len);
+        }
+        else
+        {
+            // if you reach the root, double the list
+            double_list();
+
+            // search from the beginning because list was doubled
+            return find_elem_pointer(&edges, 0, elem);
+        }
+    }
+    redistribute(node_index, len);
+
+    return find_elem_pointer(&edges, node_index, elem);
+}
+
+void PCSR::double_list()
+{
+    edges.N *= 2;
+    edges.logN = (1 << bsr_word(bsr_word(edges.N) + 1));
+    edges.H = bsr_word(edges.N / edges.logN);
+
+    for (int i = edges.N / 2; i < edges.N; i++)
+    {
+        edge_t new_edge(0, 0);
+        edges.items.push_back(new_edge);
+    }
+
+    redistribute(0, edges.N);
+}
+
+int PCSR::slide_right(int index)
+{
+    int rval = 0;
+    edge_t el = edges.items[index];
+    edges.items[index].dest = 0;
+    edges.items[index].value = 0;
+    index++;
+    while (index < edges.N && !is_null(edges.items[index]))
+    {
+        edge_t temp = edges.items[index];
+        edges.items[index] = el;
+        if (!is_null(el) && is_sentinel(el))
+        {
+            // fixing pointer of node that goes to this sentinel
+            uint32_t node_index = el.value;
+            if (node_index == UINT32_MAX)
+            {
+                node_index = 0;
+            }
+            fix_sentinel(node_index, index);
+        }
+        el = temp;
+        index++;
+    }
+    if (!is_null(el) && is_sentinel(el))
+    {
+        // fixing pointer of node that goes to this sentinel
+        uint32_t node_index = el.value;
+        if (node_index == UINT32_MAX)
+        {
+            node_index = 0;
+        }
+        fix_sentinel(node_index, index);
+    }
+    // TODO There might be an issue with this going of the end sometimes
+    if (index == edges.N)
+    {
+        index--;
+        slide_left(index);
+        rval = -1;
+        printf("slide off the end on the right, should be rare\n");
+    }
+    edges.items[index] = el;
+    return rval;
+}
+
+void PCSR::slide_left(int index)
+{
+    edge_t el = edges.items[index];
+    edges.items[index].dest = 0;
+    edges.items[index].value = 0;
+
+    index--;
+    while (index >= 0 && !is_null(edges.items[index]))
+    {
+        edge_t temp = edges.items[index];
+        edges.items[index] = el;
+        if (!is_null(el) && is_sentinel(el))
+        {
+            // fixing pointer of node that goes to this sentinel
+            uint32_t node_index = el.value;
+            if (node_index == UINT32_MAX)
+            {
+                node_index = 0;
+            }
+
+            fix_sentinel(node_index, index);
+        }
+        el = temp;
+        index--;
+    }
+
+    if (index == -1)
+    {
+        double_list();
+
+        slide_right(0);
+        index = 0;
+    }
+    if (!is_null(el) && is_sentinel(el))
+    {
+        // fixing pointer of node that goes to this sentinel
+        uint32_t node_index = el.value;
+        if (node_index == UINT32_MAX)
+        {
+            node_index = 0;
+        }
+        fix_sentinel(node_index, index);
+    }
+
+    edges.items[index] = el;
+}
+
+void PCSR::fix_sentinel(int32_t node_index, int in)
+{
+    nodes[node_index].beginning = in;
+    if (node_index > 0)
+    {
+        nodes[node_index - 1].end = in;
+    }
+    if (node_index == nodes.size() - 1)
+    {
+        nodes[node_index].end = edges.N - 1;
+    }
+}
+
+void PCSR::redistribute(int index, int len)
+{
+    // printf("REDISTRIBUTE: \n");
+    // print_array();
+    // std::vector<edge_t> space(len); //
+    edge_t new_edge;
+    vector<edge_t> space(len, new_edge);
+
+    int j = 0;
+
+    // move all items in ofm in the range into
+    // a temp array
+    for (int i = index; i < index + len; i++)
+    {
+        space[j] = edges.items[i];
+        // counting non-null edges
+        j += (!is_null(edges.items[i]));
+        // setting section to null
+        edges.items[i].value = 0;
+        edges.items[i].dest = 0;
+    }
+
+    // evenly redistribute for a uniform density
+    double index_d = index;
+    double step = ((double)len) / j;
+    for (int i = 0; i < j; i++)
+    {
+        int in = index_d;
+
+        edges.items[in] = space[i];
+        if (is_sentinel(space[i]))
+        {
+            // fixing pointer of node that goes to this sentinel
+            uint32_t node_index = space[i].value;
+            if (node_index == UINT32_MAX)
+            {
+                node_index = 0;
+            }
+            fix_sentinel(node_index, in);
+        }
+        index_d += step;
+    }
+    //   free(space);
+}
+
+void PCSR::add_edge(uint32_t src, uint32_t dest, uint32_t value)
+{
+    if (value != 0)
+    {
+        node_t node = nodes[src];
+        nodes[src].num_neighbors++;
+        nodes[dest].in_degree++;
+
+        edge_t e;
+        e.dest = dest;
+        e.value = value;
+
+        uint32_t loc_to_add =
+            binary_search(&edges, &e, node.beginning + 1, node.end);
+        insert(loc_to_add, e, src);
+        ++edge_count;
+    }
+}
+
+void PCSR::add_edge_update(uint32_t src, uint32_t dest, uint32_t value)
+{
+    if (value != 0)
+    {
+        node_t node = nodes[src];
+
+        edge_t e;
+        e.dest = dest;
+        e.value = value;
+
+        uint32_t loc_to_add =
+            binary_search(&edges, &e, node.beginning + 1, node.end);
+        if (edges.items[loc_to_add].dest == dest)
+        {
+            edges.items[loc_to_add].value = value;
+            return;
+        }
+        nodes[src].num_neighbors++;
+        nodes[dest].in_degree++;
+        insert(loc_to_add, e, src);
+        ++edge_count;
+    }
+}
+
+void PCSR::delete_edge(uint32_t src, uint32_t dest)
+{
+    edge_t e;
+    e.value = 0;
+    e.dest = dest;
+    uint32_t loc = binary_search(&edges, &e, nodes[src].beginning + 1, nodes[src].end);
+
+    if (!is_null(edges.items[loc]) && edges.items[loc].dest == dest)
+    {
+        edges.items[loc].value = 0;
+        nodes[src].num_neighbors -= 1;
+        nodes[dest].in_degree -= 1;
+    }
+    --edge_count;
+}
+
+uint64_t PCSR::get_n()
+{
+    return nodes.size();
+}
+
+vector<tuple<uint32_t, uint32_t, uint32_t>> PCSR::get_edges()
+{
+    uint64_t n = get_n();
+    vector<tuple<uint32_t, uint32_t, uint32_t>> output;
+
+    for (int i = 0; i < n; i++)
+    {
+        uint32_t start = nodes[i].beginning;
+        uint32_t end = nodes[i].end;
+        for (int j = start + 1; j < end; j++)
+        {
+            if (!is_null(edges.items[j]))
+            {
+                output.push_back(
+                    make_tuple(i, edges.items[j].dest, edges.items[j].value));
+            }
+        }
+    }
+    return output;
+}
+
+void PCSR::print_graph()
+{
+    int num_vertices = nodes.size();
+
+    // printing the graph matrix column indices
+    for (int i = 0; i < num_vertices; ++i)
+        printf("   %d", i);
+
+    printf("\n");
+
+    for (int i = 0; i < num_vertices; i++)
+    {
+        // +1 to avoid sentinel
+
+        // printing the graph matrix row indices
+        printf("%d ", i);
+        int matrix_index = 0;
+
+        for (uint32_t j = nodes[i].beginning + 1; j < nodes[i].end; j++)
+        {
+            if (!is_null(edges.items[j]))
+            {
+                while (matrix_index < edges.items[j].dest)
+                {
+                    printf("    ");
+                    matrix_index++;
+                }
+                // printf("%03d ", edges.items[j].value);
+                printf(" %d  ", edges.items[j].value);
+                matrix_index++;
+            }
+        }
+        for (uint32_t j = matrix_index; j < num_vertices; j++)
+        {
+            printf("    ");
+        }
+        printf("\n");
+    }
+}
+
+void PCSR::print_array()
+{
+    for (int i = 0; i < edges.N; i++)
+    {
+        if (is_null(edges.items[i]))
+        {
+            printf("%d-x ", i);
+        }
+        else if (is_sentinel(edges.items[i]))
+        {
+            uint32_t value = edges.items[i].value;
+            if (value == UINT32_MAX)
+            {
+                value = 0;
+            }
+            printf("\n%d-s(%u):(%d, %d) ", i, value, nodes[value].beginning,
+                   nodes[value].end);
+        }
+        else
+        {
+            printf("%d-(%d, %u) ", i, edges.items[i].dest, edges.items[i].value);
+        }
+    }
+    printf("\n\n");
+}
+
+std::tuple<std::size_t, std::size_t, std::size_t> PCSR::get_csr_ptrs()
+{
+
+    // we are building a compressed CSR arrays without
+    // the -1 values which indicates an empty edge
+
+    thrust::host_vector<int> row_offset;
+    thrust::host_vector<int> column_indices;
+    thrust::host_vector<int> eids;
+
+    int row_offset_size = nodes.size() + 1;
+    int column_indices_size = edges.items.size();
+
+    // first element of row offset is always zero
+    row_offset.push_back(0);
+
+    for (int i = 0; i < row_offset_size - 1; ++i)
+    {
+        row_offset.push_back(nodes[i].num_neighbors + row_offset[i]);
+    }
+
+    for (int i = 0; i < column_indices_size; ++i)
+    {
+        if (!is_sentinel(edges.items[i]) && !is_null(edges.items[i]))
+        {
+            column_indices.push_back(edges.items[i].dest);
+            eids.push_back(edges.items[i].value);
+        }
+    }
+
+    DEV_VEC row_offset_device = row_offset;
+    DEV_VEC column_indices_device = column_indices;
+    DEV_VEC eids_device = eids;
+
+    std::tuple<std::size_t, std::size_t, std::size_t> t;
+    std::get<0>(t) = (std::size_t)RAW_PTR(row_offset_device);
+    std::get<1>(t) = (std::size_t)RAW_PTR(column_indices_device);
+    std::get<2>(t) = (std::size_t)RAW_PTR(eids_device);
+    return t;
+}
+
+void PCSR::label_edges()
+{
+
+    uint32_t column_indices_size = edges.items.size();
+    uint32_t counter = 1;
+    for (int i = 0; i < column_indices_size; ++i)
+    {
+        if (!is_sentinel(edges.items[i]) && !is_null(edges.items[i]))
+        {
+            edges.items[i].value = counter;
+            ++counter;
+        }
+    }
+}
+
+uint32_t PCSR::find_edge_id(uint32_t src, uint32_t dest)
+{
+    edge_t e;
+    e.dest = dest;
+    // setting a random number
+    e.value = 0;
+
+    uint32_t loc = binary_search(&edges, &e, nodes[src].beginning + 1, nodes[src].end);
+    return edges.items[loc].value;
+}
+
+std::tuple<uint32_t, uint32_t> PCSR::get_graph_attr()
+{
+    std::tuple<uint32_t, uint32_t> t;
+    std::get<0>(t) = nodes.size();
+    std::get<1>(t) = edge_count;
+    return t;
+}
+
+void copy_label_edges(PCSR &pcsr, PCSR &ref_pcsr)
+{
+    uint64_t n = pcsr.get_n();
+
+    for (int i = 0; i < n; i++)
+    {
+        uint32_t start = pcsr.nodes[i].beginning;
+        uint32_t end = pcsr.nodes[i].end;
+        for (int j = start + 1; j < end; j++)
+        {
+            if (!is_null(pcsr.edges.items[j]))
+            {
+                // Searching for the edge_id of reverse edge from reference_pcsr
+                pcsr.edges.items[j].value = ref_pcsr.find_edge_id(pcsr.edges.items[j].dest, i);
+            }
+        }
+    }
+}
+
+void PCSR::edge_update_list(std::vector<std::tuple<uint32_t, uint32_t>> edge_list, bool is_delete = false, bool is_reverse_edge = false)
+{
+    for (auto &edge : edge_list)
+    {
+        uint32_t src = (is_reverse_edge == true) ? std::get<1>(edge) : std::get<0>(edge);
+        uint32_t dst = (is_reverse_edge == true) ? std::get<0>(edge) : std::get<1>(edge);
+
+        if (is_delete)
+            delete_edge(src, dst);
+        else
+            add_edge(src, dst, 1);
+    }
+}
+
+void build_reverse_pcsr(PCSR &pcsr, PCSR &ref_pcsr)
+{
+    vector<tuple<uint32_t, uint32_t, uint32_t>> edge_list = ref_pcsr.get_edges();
+    for (auto &edge : edge_list)
+    {
+        // since reverse graph the src and dst are swapped
+        uint32_t src = std::get<1>(edge);
+        uint32_t dst = std::get<0>(edge);
+        uint32_t val = std::get<2>(edge);
+        pcsr.add_edge(src, dst, val);
+    }
+}
+
+////////////////////////////////////////////////////////////////////////////
+
+// PCSR Python Module
+
+PYBIND11_MODULE(pcsr, m)
+{
+    m.doc() = "PCSR Dynamic Graph Representation";
+    m.def("copy_label_edges", &copy_label_edges, "Label edges of a PCSR based on another PCSR");
+    m.def("build_reverse_pcsr", &build_reverse_pcsr, "Builds the reverse PCSR based on another PCSR");
+
+    py::class_<node_t>(m, "Node")
+        .def(py::init<int, int, int>(), py::arg("beg") = 0, py::arg("_end") = 0, py::arg("num_neigh") = 0)
+        .def_readwrite("beginning", &_node::beginning)
+        .def_readwrite("end", &_node::end)
+        .def_readwrite("num_neighbors", &_node::num_neighbors)
+        .def_readwrite("in_degree", &_node::in_degree);
+
+    py::class_<edge_t>(m, "Edge")
+        .def(py::init<int, int>(), py::arg("_dest") = 0, py::arg("_value") = 0)
+        .def_readwrite("dest", &edge_t::dest)
+        .def_readwrite("value", &edge_t::value);
+
+    py::class_<edge_list_t>(m, "EdgeList")
+        .def(py::init<>())
+        .def_readwrite("N", &edge_list_t::N)
+        .def_readwrite("H", &edge_list_t::H)
+        .def_readwrite("logN", &edge_list_t::logN)
+        .def_readwrite("items", &edge_list_t::items);
+
+    py::class_<pair_int>(m, "PairInt")
+        .def(py::init<int, int>())
+        .def_readwrite("x", &pair_int::x)
+        .def_readwrite("y", &pair_int::y);
+
+    py::class_<pair_double>(m, "PairDouble")
+        .def(py::init<double, double>())
+        .def_readwrite("x", &pair_double::x)
+        .def_readwrite("y", &pair_double::y);
+
+    py::class_<PCSR>(m, "PCSR")
+        .def(py::init<int>(), py::arg("init_n") = 0)
+        .def("init_graph", &PCSR::init_graph)
+        .def("add_node", &PCSR::add_node)
+        .def("insert", &PCSR::insert)
+        .def("double_list", &PCSR::double_list)
+        .def("slide_right", &PCSR::slide_right)
+        .def("slide_left", &PCSR::slide_left)
+        .def("fix_sentinel", &PCSR::fix_sentinel)
+        .def("redistribute", &PCSR::redistribute)
+        .def("print_graph", &PCSR::print_graph)
+        .def("add_edge", &PCSR::add_edge)
+        .def("add_edge_update", &PCSR::add_edge_update)
+        .def("edge_update_list", &PCSR::edge_update_list, py::arg("edge_list"), py::arg("is_delete") = false, py::arg("is_reverse_edge") = false)
+        .def("label_edges", &PCSR::label_edges, "Creates edge labels for the current GPMA")
+        .def("delete_edge", &PCSR::delete_edge)
+        .def("get_n", &PCSR::get_n)
+        .def("get_edges", &PCSR::get_edges)
+        .def("print_array", &PCSR::print_array)
+        .def("get_csr_ptrs", &PCSR::get_csr_ptrs)
+        .def("get_graph_attr", &PCSR::get_graph_attr)
+        .def("find_edge_id", &PCSR::find_edge_id)
+        .def_readwrite("nodes", &PCSR::nodes)
+        .def_readwrite("edges", &PCSR::edges)
+        .def("__copy__", [](const PCSR &self)
+             { return PCSR(self); })
+        .def(
+            "__deepcopy__", [](const PCSR &self, py::dict)
+            { return PCSR(self); },
+            "memo"_a);
+}
+
+// /usr/local/cuda-11.7/bin/nvcc $(python3 -m pybind11 --includes) -shared -rdc=true --compiler-options '-fPIC'  -D__CDPRT_SUPPRESS_SYNC_DEPRECATION_WARNING -o pcsr.so pcsr.cu
\ No newline at end of file
diff --git a/seastar/graph/static/README.md b/seastar/graph/static/README.md
new file mode 100644
index 0000000..e52da4b
--- /dev/null
+++ b/seastar/graph/static/README.md
@@ -0,0 +1,7 @@
+# CSR Implementation using Python/C++
+
+## Building the shared file
+
+```
+/usr/local/cuda-11.7/bin/nvcc $(python3 -m pybind11 --includes) -shared -rdc=true --compiler-options '-fPIC' -D__CDPRT_SUPPRESS_SYNC_DEPRECATION_WARNING -o csr.so csr.cu
+```
\ No newline at end of file
diff --git a/seastar/graph/static/StaticGraph.py b/seastar/graph/static/StaticGraph.py
new file mode 100644
index 0000000..8b5a5af
--- /dev/null
+++ b/seastar/graph/static/StaticGraph.py
@@ -0,0 +1,61 @@
+from abc import ABC, abstractmethod
+import copy
+
+import numpy as np
+
+from seastar.graph.SeastarGraph import SeastarGraph
+
+
+from seastar.graph.static.csr import CSR
+
+class StaticGraph(SeastarGraph):
+    def __init__(self, edge_list):    
+        super().__init__()
+        self._num_nodes = 0
+        self._num_edges = 0
+        
+        self._get_graph_attr(edge_list)
+        
+        self._forward_graph = CSR(edge_list, self._num_nodes, is_edge_reverse=True)
+        self._backward_graph = CSR(edge_list, self._num_nodes)
+        self._forward_graph.label_edges()
+        self._backward_graph.copy_label_edges(self._forward_graph)
+        
+        self._get_graph_csr_ptrs()
+        
+    def _get_graph_attr(self, edge_list):
+        node_set = set()
+        for edge in edge_list:
+            src, dst = edge[0], edge[1]
+            node_set.add(src)
+            node_set.add(dst)
+        
+        self._num_nodes = len(node_set)
+        self._num_edges = len(edge_list)
+            
+        
+    def _get_graph_csr_ptrs(self):
+        fwd_csr_ptrs = self._forward_graph.get_csr_ptrs()
+        self.fwd_row_offset_ptr = fwd_csr_ptrs[0]
+        self.fwd_column_indices_ptr = fwd_csr_ptrs[1]
+        self.fwd_eids_ptr = fwd_csr_ptrs[2]
+        
+        bwd_csr_ptrs = self._backward_graph.get_csr_ptrs()
+        self.bwd_row_offset_ptr = bwd_csr_ptrs[0]
+        self.bwd_column_indices_ptr = bwd_csr_ptrs[1]
+        self.bwd_eids_ptr = bwd_csr_ptrs[2]
+        
+    def get_num_nodes(self):
+        return self._num_nodes
+    
+    def get_num_edges(self):
+        return self._num_edges
+        
+    def graph_type(self):
+        return "csr"
+    
+    def in_degrees(self):
+        return np.array(self._forward_graph.out_degrees, dtype='int32')
+    
+    def out_degrees(self):
+        return np.array(self._forward_graph.in_degrees, dtype='int32')
\ No newline at end of file
diff --git a/seastar/graph/static/__init__.py b/seastar/graph/static/__init__.py
new file mode 100644
index 0000000..e69de29
diff --git a/seastar/graph/static/csr.cu b/seastar/graph/static/csr.cu
new file mode 100644
index 0000000..498f845
--- /dev/null
+++ b/seastar/graph/static/csr.cu
@@ -0,0 +1,277 @@
+#include "stdio.h"
+#include <vector>
+#include <tuple>
+
+#include <thrust/device_vector.h>
+#include <thrust/host_vector.h>
+#include <thrust/remove.h>
+#include <thrust/sort.h>
+#include <cub/cub.cuh>
+
+#include <pybind11/pybind11.h>
+#include <pybind11/stl.h>
+
+namespace py = pybind11;
+using namespace pybind11::literals;
+
+#define RAW_PTR(x) thrust::raw_pointer_cast((x).data())
+
+class CSR
+{
+public:
+    thrust::device_vector<int> row_offset;
+    thrust::device_vector<int> column_indices;
+    thrust::device_vector<int> eids;
+
+    std::vector<int> in_degrees;
+    std::vector<int> out_degrees;
+
+    CSR(std::vector<std::tuple<int, int>> edge_list, int num_nodes, bool is_edge_reverse);
+    std::tuple<std::size_t, std::size_t, std::size_t> get_csr_ptrs();
+    void label_edges();
+    int find_edge_id(int src, int dst);
+    void copy_label_edges(CSR new_csr);
+    void print_csr_arrays();
+    void print_graph();
+};
+
+bool sort_by_sec(const std::tuple<int, int> &a,
+                 const std::tuple<int, int> &b)
+{
+    return std::tie(std::get<1>(a), std::get<0>(a)) <
+           std::tie(std::get<1>(b), std::get<0>(b));
+}
+
+CSR::CSR(std::vector<std::tuple<int, int>> edge_list, int num_nodes, bool is_edge_reverse = false)
+{
+    if (is_edge_reverse)
+        sort(edge_list.begin(), edge_list.end(), sort_by_sec);
+    else
+        sort(edge_list.begin(), edge_list.end());
+
+    thrust::host_vector<int> h_row_offset;
+    thrust::host_vector<int> h_column_indices;
+
+    // initialising row_offset values all to -1
+    h_row_offset.resize(num_nodes + 1);
+    thrust::fill(h_row_offset.begin(), h_row_offset.end(), -1);
+    h_row_offset[0] = 0;
+
+    in_degrees.resize(num_nodes, 0);
+    out_degrees.resize(num_nodes, 0);
+
+    int current_src;
+    int beg = 0;
+    int end = 0;
+
+    // iterating through the edge_list
+    for (auto &edge : edge_list)
+    {
+        int src = is_edge_reverse ? std::get<1>(edge) : std::get<0>(edge);
+        int dst = is_edge_reverse ? std::get<0>(edge) : std::get<1>(edge);
+
+        // first edge
+        if (beg == 0 && end == 0)
+            current_src = src;
+
+        // source node we are dealing with changes
+        if (current_src != src)
+        {
+            // update row_offset
+            h_row_offset[current_src] = beg;
+            h_row_offset[current_src + 1] = end;
+
+            current_src = src;
+            beg = end;
+        }
+
+        // adding the dst node to the column indices
+        // and incrementing the end range by 1
+        h_column_indices.push_back(dst);
+        end += 1;
+
+        // updating the degree arrays
+        out_degrees[src] += 1;
+        in_degrees[dst] += 1;
+    }
+
+    h_row_offset[current_src + 1] = end;
+
+    // removing the -1
+    int curr_val = h_row_offset[0];
+    for (int i = 1; i < h_row_offset.size(); ++i)
+    {
+        if (h_row_offset[i] != curr_val && h_row_offset[i] != -1)
+            curr_val = h_row_offset[i];
+
+        if (h_row_offset[i] == -1)
+            h_row_offset[i] = curr_val;
+    }
+
+    row_offset = h_row_offset;
+    column_indices = h_column_indices;
+}
+
+std::tuple<std::size_t, std::size_t, std::size_t> CSR::get_csr_ptrs()
+{
+    std::tuple<std::size_t, std::size_t, std::size_t> t;
+    std::get<0>(t) = (std::size_t)RAW_PTR(row_offset);
+    std::get<1>(t) = (std::size_t)RAW_PTR(column_indices);
+    std::get<2>(t) = (std::size_t)RAW_PTR(eids);
+    return t;
+}
+
+int binary_search_idx(thrust::host_vector<int> vec, int start, int end, int elem)
+{
+    auto lower = std::lower_bound(vec.begin() + start, vec.begin() + end, elem);
+    const bool found = lower != vec.begin() + end && *lower == elem;
+
+    if (found)
+        return std::distance(vec.begin(), lower);
+    else
+        return -1;
+}
+
+int CSR::find_edge_id(int src, int dst)
+{
+    thrust::host_vector<int> h_row_offset = row_offset;
+    thrust::host_vector<int> h_column_indices = column_indices;
+    thrust::host_vector<int> h_eids = eids;
+
+    int beg = h_row_offset[src];
+    int end = h_row_offset[src + 1];
+
+    int col_idx = binary_search_idx(h_column_indices, beg, end, dst);
+
+    if (col_idx != -1)
+        return h_eids[col_idx];
+    else
+        return -1;
+}
+
+void CSR::label_edges()
+{
+    eids.resize(column_indices.size());
+    thrust::sequence(eids.begin(), eids.end());
+}
+
+void CSR::copy_label_edges(CSR ref_csr)
+{
+
+    thrust::host_vector<int> h_row_offset = row_offset;
+    thrust::host_vector<int> h_column_indices = column_indices;
+    thrust::host_vector<int> h_eids(h_column_indices.size());
+
+    int num_nodes = h_row_offset.size() - 1;
+    for (int src = 0; src < num_nodes; ++src)
+    {
+        int beg = h_row_offset[src];
+        int end = h_row_offset[src + 1];
+
+        for (int dst_idx = beg; dst_idx < end; ++dst_idx)
+        {
+            int dst = h_column_indices[dst_idx];
+            int id = ref_csr.find_edge_id(dst, src);
+            h_eids[dst_idx] = id;
+        }
+    }
+
+    eids = h_eids;
+}
+
+void CSR::print_csr_arrays()
+{
+    thrust::host_vector<int> h_row = row_offset;
+    thrust::host_vector<int> h_col = column_indices;
+    thrust::host_vector<int> h_eids = eids;
+
+    cudaDeviceSynchronize();
+
+    printf("\nRow offsets:\n");
+    for (int i = 0; i < h_row.size(); ++i)
+        printf("%d ", h_row[i]);
+
+    printf("\n");
+
+    printf("\nOut-degrees:\n");
+    for (int i = 0; i < out_degrees.size(); ++i)
+        printf("%d ", out_degrees[i]);
+
+    printf("\nIn-degrees:\n");
+    for (int i = 0; i < in_degrees.size(); ++i)
+        printf("%d ", in_degrees[i]);
+
+    printf("\nColumn Indices:\n");
+    for (int i = 0; i < h_col.size(); ++i)
+        printf("%d ", h_col[i]);
+
+    printf("\nEids:\n");
+    for (int i = 0; i < h_eids.size(); ++i)
+        printf("%d ", h_eids[i]);
+
+    printf("\n");
+}
+
+void CSR::print_graph()
+{
+    thrust::host_vector<int> h_row_offset = row_offset;
+    thrust::host_vector<int> h_column_indices = column_indices;
+    thrust::host_vector<int> h_eids = eids;
+
+    int num_vertices = h_row_offset.size() - 1;
+
+    // printing the graph matrix column indices
+    for (int dst = 0; dst < num_vertices; ++dst)
+        std::cout << "   " << dst << std::flush;
+
+    printf("\n");
+
+    for (int src = 0; src < num_vertices; ++src)
+    {
+        std::cout << src << " " << std::flush;
+        int matrix_index = 0;
+
+        int beg = h_row_offset[src];
+        int end = h_row_offset[src + 1];
+
+        for (int dst_idx = beg; dst_idx < end; ++dst_idx)
+        {
+            while (matrix_index < h_column_indices[dst_idx])
+            {
+                std::cout << "    " << std::flush;
+                matrix_index++;
+            }
+            std::cout << " " << h_eids[dst_idx] << "  " << std::flush;
+            matrix_index++;
+        }
+
+        for (int j = matrix_index; j < num_vertices; ++j)
+            std::cout << "    " << std::flush;
+
+        printf("\n");
+    }
+}
+
+PYBIND11_MODULE(csr, m)
+{
+    m.doc() = "CPython module for CSR"; // optional module docstring
+
+    py::class_<CSR>(m, "CSR")
+        .def(py::init<std::vector<std::tuple<int, int>>, int, bool>(), py::arg("edge_list"), py::arg("num_nodes"), py::arg("is_edge_reverse") = false)
+        .def("label_edges", &CSR::label_edges)
+        .def("copy_label_edges", &CSR::copy_label_edges)
+        .def("get_csr_ptrs", &CSR::get_csr_ptrs)
+        .def("print_csr_arrays", &CSR::print_csr_arrays)
+        .def("print_graph", &CSR::print_graph)
+        .def_readwrite("row_offset", &CSR::row_offset)
+        .def_readwrite("column_indices", &CSR::column_indices)
+        .def_readwrite("eids", &CSR::eids)
+        .def_readwrite("out_degrees", &CSR::out_degrees)
+        .def_readwrite("in_degrees", &CSR::in_degrees)
+        .def("__copy__", [](const CSR &self)
+             { return CSR(self); })
+        .def(
+            "__deepcopy__", [](const CSR &self, py::dict)
+            { return CSR(self); },
+            "memo"_a);
+}
\ No newline at end of file
diff --git a/seastar/nn/README.md b/seastar/nn/README.md
new file mode 100644
index 0000000..d3d8d59
--- /dev/null
+++ b/seastar/nn/README.md
@@ -0,0 +1,3 @@
+# Seastar Neural Networks
+
+Module that contains already built Neural Networks written using Seastar
\ No newline at end of file
diff --git a/seastar/nn/__init__.py b/seastar/nn/__init__.py
new file mode 100644
index 0000000..ca53a6e
--- /dev/null
+++ b/seastar/nn/__init__.py
@@ -0,0 +1,1 @@
+'''State of the art Graph Neural Networks written using Seastar'''
\ No newline at end of file
diff --git a/python/setup.py b/setup.py
similarity index 57%
rename from python/setup.py
rename to setup.py
index 2730428..a204764 100644
--- a/python/setup.py
+++ b/setup.py
@@ -2,12 +2,15 @@
 
 setuptools.setup(
     name='seastar',
-    version='0.1',
-    author='ydwu',
-    author_email='ydwu2014@gmail.com',
-    description='Easy Fast Graph learning',
-    url='https://github.com/ydwu4/seastar-paper-version',
+    version='1.0.0',
+    author='Joel Mathew Cherian, Nithin Puthalath Manoj',
+    author_email='joelmathewcherian@gmail.com, nithinp.manoj@gmail.com',
+    description='Vertex-Centric Approach to building Graph Neural Networks',
+    url='https://github.com/bfGraph/Seastar',
     packages=setuptools.find_packages(),
+    package_data={
+        "": ["*.so","*.jinja"]
+    },
     classifiers=[
         "Programming Language :: Python :: 3",
         "License :: OSI Approved :: MIT License",


Test patch:
diff --git a/python/dynamic_graph/PCSR/datasets/test_graph.txt b/python/dynamic_graph/PCSR/datasets/test_graph.txt
deleted file mode 100644
index d9a4816..0000000
--- a/python/dynamic_graph/PCSR/datasets/test_graph.txt
+++ /dev/null
@@ -1,7 +0,0 @@
-3 6
-0 0 1
-0 2 2
-2 2 3
-2 0 4
-2 1 5
-1 2 6
\ No newline at end of file