('DynamicBitset&', 'bitset_comp_current'),
('HGAccumulator<unsigned int>', 'DGAccumulator_accum')
], [
CDecl([
("__shared__ cub::BlockReduce<unsigned int, TB_SIZE>::TempStorage",
"DGAccumulator_accum_ts", "")
]),
CBlock(["DGAccumulator_accum.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(
ForAll("jj", G.edges("src"), [
CDecl([("index_type", "dst", "")]),
CBlock(["dst = graph.getAbsDestination(jj)"]),
CDecl([("uint32_t", "new_comp", "")]),
CBlock(["new_comp = p_comp_current[dst]"]),
CDecl([("uint32_t", "old_comp", "")]),
CBlock([
"old_comp = atomicTestMin(&p_comp_current[src], new_comp)"
]),
If("old_comp > new_comp", [
CBlock(["bitset_comp_current.set(src)"]),
CBlock(["DGAccumulator_accum.reduce( 1)"]),
]),
]), ),
]),
CBlock([
ForAll("src", G.nodes("__begin", "__end"), [
CDecl([("bool", "pop", " = src < __end")]),
If("pop", [
If("p_delta[src] > 0", [
CBlock(["_delta = p_delta[src]"]),
CBlock(["p_delta[src] = 0"]),
CBlock(["active_vertices.reduce( 1)"]),
], [
CBlock(["pop = false"]),
]),
]),
UniformConditional(If("!pop", [CBlock("continue")]),
uniform_only=False,
_only_if_np=True),
ClosureHint(
ForAll("nbr", G.edges("src"), [
CDecl([("index_type", "dst", "")]),
CBlock(["dst = graph.getAbsDestination(nbr)"]),
CBlock(["atomicTestAdd(&p_residual[dst], _delta)"]),
CBlock(["bitset_residual.set(dst)"]),
]), ),
]),
CBlock([
"active_vertices.thread_exit<cub::BlockReduce<unsigned int, TB_SIZE> >(active_vertices_ts)"
],
parse=False),
]),
Kernel("PageRankSanity", [
G.param(), ('unsigned int', '__begin'), ('unsigned int', '__end'),
('float', 'local_tolerance'), ('float *', 'p_residual'),
('float *', 'p_value'),
]),
]),
]),
Kernel("FirstIterationSSSP", [
G.param(), ('unsigned int', '__begin'), ('unsigned int', '__end'),
('uint32_t *', 'p_current_length'),
('DynamicBitset&', 'bitset_current_length')
], [
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(
ForAll("current_edge", G.edges("src"), [
CDecl([("index_type", "dst", "")]),
CBlock(["dst = graph.getAbsDestination(current_edge)"]),
CDecl([("int", "edge_weight", "")]),
CBlock(["edge_weight = 1"]),
CBlock(["edge_weight += graph.getAbsWeight(current_edge)"
]),
CDecl([("uint32_t", "new_dist", "")]),
CBlock(["new_dist = edge_weight + p_current_length[src]"]),
CBlock(["atomicMin(&p_current_length[dst], new_dist)"]),
CBlock(["bitset_current_length.set(dst)"]),
]), ),
]),
]),
Kernel("SSSP", [
G.param(), ('unsigned int', '__begin'), ('unsigned int', '__end'),
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)")]),
CFor(CDecl(("int", "i", "= 0")), "i < N_NODES", "i++",
[CBlock([("colors.cpu_wr_ptr()[i] = COLORS[i]")])]),
Invoke("check_coloring",
["gg", "colors.gpu_rd_ptr()", "chromatic_number.zero_gpu()"]),
