# -*- mode: python -*-
from gg.ast import *
from gg.lib.graph import Graph
from gg.lib.wl import Worklist
from gg.ast.params import GraphParam
import cgen
G = Graph("graph")
EL = Worklist() # EdgeList
ast = Module([
CBlock(["typedef int edge_data_type", "typedef int node_data_type"]),
Kernel(
"ebgc_init",
[
G.param(), ("int *", "edge_src"), ("int *", "VForbidden"),
("int *", "CS"), ("int *", "TVForbidden")
],
[
ForAll(
"node",
G.nodes(),
[
CBlock([
"graph.node_data[node] = 0", # 0 represents not colored
"VForbidden[node] = 0",
"TVForbidden[node] = 0",
"CS[node] = 0"
]),
from gg.ast import *
from gg.lib.graph import Graph
from gg.lib.wl import Worklist
from gg.ast.params import GraphParam
import cgen
G = Graph("graph")
WL = Worklist()
ast = Module([
CBlock([cgen.Include("kernels/reduce.cuh", system=False)], parse=False),
CBlock([cgen.Include("gen_cuda.cuh", system=False)], parse=False),
Kernel("InitializeGraph", [
G.param(), ('unsigned int', '__begin'), ('unsigned int', '__end'),
('uint32_t *', 'p_comp_current'), ('uint32_t *', 'p_comp_old')
], [
ForAll("src", G.nodes("__begin", "__end"), [
CDecl([("bool", "pop", " = src < __end")]),
If("pop", [
CBlock(["p_comp_current[src] = graph.node_data[src]"]),
CBlock(["p_comp_old[src] = graph.node_data[src]"]),
]),
]),
]),
Kernel("FirstItr_ConnectedComp", [
G.param(), ('unsigned int', '__begin'), ('unsigned int', '__end'),
('uint32_t *', 'p_comp_current'), ('uint32_t *', 'p_comp_old'),
('DynamicBitset&', 'bitset_comp_current')
], [
ForAll("src", G.nodes("__begin", "__end"), [
CDecl([("bool", "pop", " = src < __end")]),
If("pop", [
CBlock(["p_comp_old[src] = p_comp_current[src]"]),
from gg.ast import *
from gg.lib.graph import Graph
from gg.lib.wl import Worklist
from gg.ast.params import GraphParam
import cgen
G = Graph("graph")
WL = Worklist()
ast = Module([
CBlock([cgen.Include("kernels/reduce.cuh", system = False)], parse = False),
CBlock([cgen.Include("gen_cuda.cuh", system = False)], parse = False),
Kernel("ResetGraph", [G.param(), ('unsigned int', '__begin'), ('unsigned int', '__end'), ('const float ', 'local_alpha'), ('float *', 'p_delta'), ('uint32_t *', 'p_nout'), ('float *', 'p_residual'), ('float *', 'p_value')],
[
ForAll("src", G.nodes("__begin", "__end"),
[
CDecl([("bool", "pop", " = src < __end")]),
If("pop", [
CBlock(["p_value[src] = 0"]),
CBlock(["p_nout[src] = 0"]),
CBlock(["p_delta[src] = 0"]),
CBlock(["p_residual[src] = local_alpha"]),
]),
]),
]),
Kernel("InitializeGraph", [G.param(), ('unsigned int', '__begin'), ('unsigned int', '__end'), ('uint32_t *', 'p_nout'), ('DynamicBitset&', 'bitset_nout')],
[
ForAll("src", G.nodes("__begin", "__end"),
[
CDecl([("bool", "pop", " = src < __end")]),
If("pop", [
]),
UniformConditional(If("!pop", [CBlock("continue")]), uniform_only = False, _only_if_np = True),
# -*- mode: python -*-
from gg.ast import *
from gg.lib.graph import Graph
from gg.lib.wl import Worklist
from gg.ast.params import GraphParam
import cgen
G = Graph("graph")
WL = Worklist()
ast = Module([
CBlock(["typedef int edge_data_type", "typedef int node_data_type"]),
Kernel("ipgc_init", [G.param(), ("int *", "max_degree")], [
ForAll("node", G.nodes(), [
CBlock([
"graph.node_data[node] = 0",
"atomicMax(max_degree, graph.getOutDegree(node))"
]),
WL.push("node")
])
]), #initializing all nodes to color 0
Kernel(
"assignColors",
[G.param(), ("bool *", "forbidden"), ("int *", "max_degree")],
[
ForAll(
"wlnode",
WL.items(),
[
from gg.ast import *
from gg.lib.graph import Graph
from gg.lib.wl import Worklist
from gg.ast.params import GraphParam
import cgen
G = Graph("graph")
WL = Worklist()
ast = Module([
CBlock([cgen.Include("kernels/reduce.cuh", system = False)], parse = False),
CBlock([cgen.Include("cc_pull_cuda.cuh", system = False)], parse = False),
Kernel("InitializeGraph", [G.param(), ('unsigned int', '__begin'), ('unsigned int', '__end'), ('uint32_t *', 'p_comp_current')],
[
ForAll("src", G.nodes("__begin", "__end"),
[
CDecl([("bool", "pop", " = src < __end")]),
If("pop", [
CBlock(["p_comp_current[src] = graph.node_data[src]"]),
]),
]),
]),
Kernel("ConnectedComp", [G.param(), ('unsigned int', '__begin'), ('unsigned int', '__end'), ('uint32_t *', 'p_comp_current'), ('DynamicBitset&', 'bitset_comp_current'), ('HGAccumulator<unsigned int>', 'active_vertices')],
[
CDecl([("__shared__ cub::BlockReduce<unsigned int, TB_SIZE>::TempStorage", "active_vertices_ts", "")]),
CBlock(["active_vertices.thread_entry()"]),
ForAll("src", G.nodes("__begin", "__end"),
[
CDecl([("bool", "pop", " = src < __end")]),
If("pop", [
]),
UniformConditional(If("!pop", [CBlock("continue")]), uniform_only = False, _only_if_np = True),
ClosureHint(
# -*- mode: python -*-
from gg.ast import *
from gg.lib.graph import Graph
from gg.lib.wl import Worklist
from gg.ast.params import GraphParam
import cgen
G = Graph("graph")
WL = Worklist()
ast = Module([
CBlock(["typedef int edge_data_type", "typedef int node_data_type"]),
CBlock([cgen.Define("MAXFORBID", "32")]),
Kernel("ipgc_init", [G.param()], [
ForAll("node", G.nodes(),
[CBlock(["graph.node_data[node] = 0"]),
WL.push("node")])
]), #initializing all nodes to color 0
Kernel(
"assignColors",
[G.param()],
[
ForAll(
"node",
G.nodes(),
[
CDecl([("int", "FORBIDDEN", ""), ("int", "offset", "=0"),
("bool", "colored", "=false")]),
If(
from gg.ast import *
from gg.lib.graph import Graph
from gg.lib.wl import Worklist
from gg.ast.params import GraphParam
import cgen
G = Graph("graph")
WL = Worklist()
ast = Module([
CBlock([cgen.Include("kernels/reduce.cuh", system = False)], parse = False),
CBlock([cgen.Include("tc_cuda.cuh", system = False)], parse = False),
CBlock([cgen.Include("kernels/segmentedsort.cuh", system = False)], parse = False),
CBlock([cgen.Include("moderngpu.cuh", system = False)], parse = False),
CBlock([cgen.Include("util/mgpucontext.h", system = False)], parse = False),
CBlock([cgen.Include("cuda_profiler_api.h", system = True)], parse = False),
CBlock('mgpu::ContextPtr mgc', parse = False),
Kernel("intersect", [G.param(), ('index_type', 'u'), ('index_type', 'v')],
[
CDecl([('index_type', 'u_start', '= graph.getFirstEdge(u)'),
('index_type', 'u_end', '= u_start + graph.getOutDegree(u)'),
('index_type', 'v_start', '= graph.getFirstEdge(v)'),
('index_type', 'v_end', '= v_start + graph.getOutDegree(v)'),
('int', 'count', '= 0'),
('index_type', 'u_it', '= u_start'),
('index_type', 'v_it', '= v_start'),
('index_type', 'a', ''),
('index_type', 'b', ''),
]),
While('u_it < u_end && v_it < v_end',
[
from gg.ast import *
from gg.lib.graph import Graph
from gg.lib.wl import Worklist
from gg.ast.params import GraphParam
import cgen
G = Graph("graph")
WL = Worklist()
ast = Module([
CBlock([cgen.Include("kernels/reduce.cuh", system = False)], parse = False),
CBlock([cgen.Include("sssp_pull_cuda.cuh", system = False)], parse = False),
Kernel("InitializeGraph", [G.param(), ('unsigned int', '__begin'), ('unsigned int', '__end'), ('const uint32_t ', 'local_infinity'), ('unsigned long long', 'local_src_node'), ('uint32_t *', 'p_dist_current')],
[
ForAll("src", G.nodes("__begin", "__end"),
[
CDecl([("bool", "pop", " = src < __end")]),
If("pop", [
CBlock(["p_dist_current[src] = (graph.node_data[src] == local_src_node) ? 0 : local_infinity"]),
]),
]),
]),
Kernel("SSSP", [G.param(), ('unsigned int', '__begin'), ('unsigned int', '__end'), ('uint32_t *', 'p_dist_current'), ('DynamicBitset&', 'bitset_dist_current'), ('HGAccumulator<unsigned int>', 'active_vertices')],
[
CDecl([("__shared__ cub::BlockReduce<unsigned int, TB_SIZE>::TempStorage", "active_vertices_ts", "")]),
CBlock(["active_vertices.thread_entry()"]),
ForAll("src", G.nodes("__begin", "__end"),
[
CDecl([("bool", "pop", " = src < __end")]),
If("pop", [
]),
UniformConditional(If("!pop", [CBlock("continue")]), uniform_only = False, _only_if_np = True),
ClosureHint(
from gg.ast import *
from gg.lib.graph import Graph
from gg.lib.wl import Worklist
G = Graph("graph")
WL = Worklist()
ast = Module([
CDecl([("extern int", "N_NODES", ""), ("extern int*", "COLORS", ""),
("extern int", "CHROMATIC_NUMBER", "")]),
CBlock([
cgen.Define("NOT_VALID", "0"),
cgen.Define("VALID", "1"),
]),
Kernel(
"check_coloring",
[G.param(), ("int *", "colors"), ("int *", "chromatic_number")], [
ForAll("node", G.nodes(), [
CBlock([("graph.node_data[node] = VALID")]),
ForAll("edge", G.edges("node"), [
CDecl([("int", "dst", "=graph.getAbsDestination(edge)")]),
If("colors[node] == colors[dst]", [
CBlock([("graph.node_data[node] = NOT_VALID")]),
]),
]),
CBlock(["atomicMax(chromatic_number, colors[node])"]),
]),
]),
Kernel("gg_main", [('CSRGraphTy&', 'hg'), ('CSRGraphTy&', 'gg')], [
CDecl([("Shared<int>", "colors", "(hg.nnodes)"),
("Shared<int>", "chromatic_number", "(1)")]),