Skip to content

Commit

Permalink
Refactor SDFG List to CFG List (#1511)
Browse files Browse the repository at this point in the history
Refactors the SDFG list and SDFG ids to a more general CFG list and CFG
id.
This change provides a further step towards compatibility with the
[hierarchical, multi-level
SDFGs](#1404).

The change is accompanied by a deprecation warning on the old interfaces
(`sdfg_list`, `sdfg_id`, `reset_sdfg_list`, and `update_sdfg_list`).
Until removal, the legacy methods continue to work as expected for
single level SDFGs.

---------

Co-authored-by: alexnick83 <[email protected]>
  • Loading branch information
phschaad and alexnick83 authored Feb 23, 2024
1 parent 9ee470c commit b5ec059
Show file tree
Hide file tree
Showing 69 changed files with 450 additions and 358 deletions.
12 changes: 6 additions & 6 deletions dace/codegen/control_flow.py
Original file line number Diff line number Diff line change
Expand Up @@ -126,7 +126,7 @@ class SingleState(ControlFlow):
def as_cpp(self, codegen, symbols) -> str:
sdfg = self.state.parent

expr = '__state_{}_{}:;\n'.format(sdfg.sdfg_id, self.state.label)
expr = '__state_{}_{}:;\n'.format(sdfg.cfg_id, self.state.label)
if self.state.number_of_nodes() > 0:
expr += '{\n'
expr += self.dispatch_state(self.state)
Expand All @@ -138,7 +138,7 @@ def as_cpp(self, codegen, symbols) -> str:

# If any state has no children, it should jump to the end of the SDFG
if not self.last_state and sdfg.out_degree(self.state) == 0:
expr += 'goto __state_exit_{};\n'.format(sdfg.sdfg_id)
expr += 'goto __state_exit_{};\n'.format(sdfg.cfg_id)
return expr

def generate_transition(self,
Expand Down Expand Up @@ -175,7 +175,7 @@ def generate_transition(self,

if (not edge.data.is_unconditional()
or ((successor is None or edge.dst is not successor) and not assignments_only)):
expr += 'goto __state_{}_{};\n'.format(sdfg.sdfg_id, edge.dst.label)
expr += 'goto __state_{}_{};\n'.format(sdfg.cfg_id, edge.dst.label)

if not edge.data.is_unconditional() and not assignments_only:
expr += '}\n'
Expand Down Expand Up @@ -257,7 +257,7 @@ def as_cpp(self, codegen, symbols) -> str:
# One unconditional edge
if (len(out_edges) == 1 and out_edges[0].data.is_unconditional()):
continue
expr += f'goto __state_exit_{sdfg.sdfg_id};\n'
expr += f'goto __state_exit_{sdfg.cfg_id};\n'

return expr

Expand Down Expand Up @@ -326,7 +326,7 @@ def as_cpp(self, codegen, symbols) -> str:
# execution should end, so we emit an "else goto exit" here.
if len(self.body) > 0:
expr += ' else {\n'
expr += 'goto __state_exit_{};\n'.format(self.sdfg.sdfg_id)
expr += 'goto __state_exit_{};\n'.format(self.sdfg.cfg_id)
if len(self.body) > 0:
expr += '\n}'
return expr
Expand Down Expand Up @@ -475,7 +475,7 @@ def as_cpp(self, codegen, symbols) -> str:
expr += f'case {case}: {{\n'
expr += body.as_cpp(codegen, symbols)
expr += 'break;\n}\n'
expr += f'default: goto __state_exit_{self.sdfg.sdfg_id};'
expr += f'default: goto __state_exit_{self.sdfg.cfg_id};'
expr += '\n}\n'
return expr

Expand Down
4 changes: 2 additions & 2 deletions dace/codegen/instrumentation/data/data_dump.py
Original file line number Diff line number Diff line change
Expand Up @@ -161,7 +161,7 @@ def on_node_end(self, sdfg: SDFG, state: SDFGState, node: nodes.AccessNode, oute
# Create UUID
state_id = sdfg.node_id(state)
node_id = state.node_id(node)
uuid = f'{sdfg.sdfg_id}_{state_id}_{node_id}'
uuid = f'{sdfg.cfg_id}_{state_id}_{node_id}'

# Get optional pre/postamble for instrumenting device data
preamble, postamble = '', ''
Expand Down Expand Up @@ -277,7 +277,7 @@ def on_node_begin(self, sdfg: SDFG, state: SDFGState, node: nodes.AccessNode, ou
# Create UUID
state_id = sdfg.node_id(state)
node_id = state.node_id(node)
uuid = f'{sdfg.sdfg_id}_{state_id}_{node_id}'
uuid = f'{sdfg.cfg_id}_{state_id}_{node_id}'

# Get optional pre/postamble for instrumenting device data
preamble, postamble = '', ''
Expand Down
4 changes: 2 additions & 2 deletions dace/codegen/instrumentation/gpu_events.py
Original file line number Diff line number Diff line change
Expand Up @@ -65,11 +65,11 @@ def _report(self, timer_name: str, sdfg=None, state=None, node=None):
int __dace_micros_{id} = (int) (__dace_ms_{id} * 1000.0);
unsigned long int __dace_ts_end_{id} = std::chrono::duration_cast<std::chrono::microseconds>(std::chrono::high_resolution_clock::now().time_since_epoch()).count();
unsigned long int __dace_ts_start_{id} = __dace_ts_end_{id} - __dace_micros_{id};
__state->report.add_completion("{timer_name}", "GPU", __dace_ts_start_{id}, __dace_ts_end_{id}, {sdfg_id}, {state_id}, {node_id});'''.format(
__state->report.add_completion("{timer_name}", "GPU", __dace_ts_start_{id}, __dace_ts_end_{id}, {cfg_id}, {state_id}, {node_id});'''.format(
id=idstr,
timer_name=timer_name,
backend=self.backend,
sdfg_id=sdfg.sdfg_id,
cfg_id=sdfg.cfg_id,
state_id=state_id,
node_id=node_id)

Expand Down
52 changes: 26 additions & 26 deletions dace/codegen/instrumentation/likwid.py
Original file line number Diff line number Diff line change
Expand Up @@ -169,7 +169,7 @@ def on_sdfg_end(self, sdfg, local_stream, global_stream):
'''
local_stream.write(outer_code, sdfg)

for region, sdfg_id, state_id, node_id in self._regions:
for region, cfg_id, state_id, node_id in self._regions:
report_code = f'''
#pragma omp parallel
{{
Expand All @@ -187,7 +187,7 @@ def on_sdfg_end(self, sdfg, local_stream, global_stream):
for (int t = 0; t < num_threads; t++)
{{
__state->report.add_completion("Timer", "likwid", 0, time[t] * 1000 * 1000, t, {sdfg_id}, {state_id}, {node_id});
__state->report.add_completion("Timer", "likwid", 0, time[t] * 1000 * 1000, t, {cfg_id}, {state_id}, {node_id});
}}
for (int i = 0; i < nevents; i++)
Expand All @@ -196,7 +196,7 @@ def on_sdfg_end(self, sdfg, local_stream, global_stream):
for (int t = 0; t < num_threads; t++)
{{
__state->report.add_counter("{region}", "likwid", event_name, events[t][i], t, {sdfg_id}, {state_id}, {node_id});
__state->report.add_counter("{region}", "likwid", event_name, events[t][i], t, {cfg_id}, {state_id}, {node_id});
}}
}}
}}
Expand All @@ -214,11 +214,11 @@ def on_state_begin(self, sdfg, state, local_stream, global_stream):
return

if state.instrument == dace.InstrumentationType.LIKWID_CPU:
sdfg_id = sdfg.sdfg_id
cfg_id = sdfg.cfg_id
state_id = sdfg.node_id(state)
node_id = -1
region = f"state_{sdfg_id}_{state_id}_{node_id}"
self._regions.append((region, sdfg_id, state_id, node_id))
region = f"state_{cfg_id}_{state_id}_{node_id}"
self._regions.append((region, cfg_id, state_id, node_id))

marker_code = f'''
#pragma omp parallel
Expand Down Expand Up @@ -250,10 +250,10 @@ def on_state_end(self, sdfg, state, local_stream, global_stream):
return

if state.instrument == dace.InstrumentationType.LIKWID_CPU:
sdfg_id = sdfg.sdfg_id
cfg_id = sdfg.cfg_id
state_id = sdfg.node_id(state)
node_id = -1
region = f"state_{sdfg_id}_{state_id}_{node_id}"
region = f"state_{cfg_id}_{state_id}_{node_id}"

marker_code = f'''
#pragma omp parallel
Expand All @@ -272,12 +272,12 @@ def on_scope_entry(self, sdfg, state, node, outer_stream, inner_stream, global_s
elif node.schedule not in LIKWIDInstrumentationCPU.perf_whitelist_schedules:
raise TypeError("Unsupported schedule on scope")

sdfg_id = sdfg.sdfg_id
cfg_id = sdfg.cfg_id
state_id = sdfg.node_id(state)
node_id = state.node_id(node)
region = f"scope_{sdfg_id}_{state_id}_{node_id}"
region = f"scope_{cfg_id}_{state_id}_{node_id}"

self._regions.append((region, sdfg_id, state_id, node_id))
self._regions.append((region, cfg_id, state_id, node_id))
marker_code = f'''
#pragma omp parallel
{{
Expand All @@ -294,10 +294,10 @@ def on_scope_exit(self, sdfg, state, node, outer_stream, inner_stream, global_st
if not self._likwid_used or entry_node.instrument != dace.InstrumentationType.LIKWID_CPU:
return

sdfg_id = sdfg.sdfg_id
cfg_id = sdfg.cfg_id
state_id = sdfg.node_id(state)
node_id = state.node_id(entry_node)
region = f"scope_{sdfg_id}_{state_id}_{node_id}"
region = f"scope_{cfg_id}_{state_id}_{node_id}"

marker_code = f'''
#pragma omp parallel
Expand Down Expand Up @@ -366,7 +366,7 @@ def on_sdfg_end(self, sdfg, local_stream, global_stream):
if not self._likwid_used or sdfg.parent is not None:
return

for region, sdfg_id, state_id, node_id in self._regions:
for region, cfg_id, state_id, node_id in self._regions:
report_code = f'''
{{
double *events = (double*) malloc(MAX_NUM_EVENTS * sizeof(double));
Expand All @@ -377,14 +377,14 @@ def on_sdfg_end(self, sdfg, local_stream, global_stream):
LIKWID_NVMARKER_GET("{region}", &ngpus, &nevents, &events, &time, &count);
__state->report.add_completion("Timer", "likwid_gpu", 0, time * 1000 * 1000, 0, {sdfg_id}, {state_id}, {node_id});
__state->report.add_completion("Timer", "likwid_gpu", 0, time * 1000 * 1000, 0, {cfg_id}, {state_id}, {node_id});
int gid = nvmon_getIdOfActiveGroup();
for (int i = 0; i < nevents; i++)
{{
char* event_name = nvmon_getEventName(gid, i);
__state->report.add_counter("{region}", "likwid_gpu", event_name, events[i], 0, {sdfg_id}, {state_id}, {node_id});
__state->report.add_counter("{region}", "likwid_gpu", event_name, events[i], 0, {cfg_id}, {state_id}, {node_id});
}}
free(events);
Expand All @@ -402,11 +402,11 @@ def on_state_begin(self, sdfg, state, local_stream, global_stream):
return

if state.instrument == dace.InstrumentationType.LIKWID_GPU:
sdfg_id = sdfg.sdfg_id
cfg_id = sdfg.cfg_id
state_id = sdfg.node_id(state)
node_id = -1
region = f"state_{sdfg_id}_{state_id}_{node_id}"
self._regions.append((region, sdfg_id, state_id, node_id))
region = f"state_{cfg_id}_{state_id}_{node_id}"
self._regions.append((region, cfg_id, state_id, node_id))

marker_code = f'''
LIKWID_NVMARKER_REGISTER("{region}");
Expand All @@ -424,10 +424,10 @@ def on_state_end(self, sdfg, state, local_stream, global_stream):
return

if state.instrument == dace.InstrumentationType.LIKWID_GPU:
sdfg_id = sdfg.sdfg_id
cfg_id = sdfg.cfg_id
state_id = sdfg.node_id(state)
node_id = -1
region = f"state_{sdfg_id}_{state_id}_{node_id}"
region = f"state_{cfg_id}_{state_id}_{node_id}"

marker_code = f'''
LIKWID_NVMARKER_STOP("{region}");
Expand All @@ -443,12 +443,12 @@ def on_scope_entry(self, sdfg, state, node, outer_stream, inner_stream, global_s
elif node.schedule not in LIKWIDInstrumentationGPU.perf_whitelist_schedules:
raise TypeError("Unsupported schedule on scope")

sdfg_id = sdfg.sdfg_id
cfg_id = sdfg.cfg_id
state_id = sdfg.node_id(state)
node_id = state.node_id(node)
region = f"scope_{sdfg_id}_{state_id}_{node_id}"
region = f"scope_{cfg_id}_{state_id}_{node_id}"

self._regions.append((region, sdfg_id, state_id, node_id))
self._regions.append((region, cfg_id, state_id, node_id))
marker_code = f'''
LIKWID_NVMARKER_REGISTER("{region}");
Expand All @@ -465,10 +465,10 @@ def on_scope_exit(self, sdfg, state, node, outer_stream, inner_stream, global_st
if not self._likwid_used or entry_node.instrument != dace.InstrumentationType.LIKWID_GPU:
return

sdfg_id = sdfg.sdfg_id
cfg_id = sdfg.cfg_id
state_id = sdfg.node_id(state)
node_id = state.node_id(entry_node)
region = f"scope_{sdfg_id}_{state_id}_{node_id}"
region = f"scope_{cfg_id}_{state_id}_{node_id}"

marker_code = f'''
LIKWID_NVMARKER_STOP("{region}");
Expand Down
2 changes: 1 addition & 1 deletion dace/codegen/instrumentation/provider.py
Original file line number Diff line number Diff line change
Expand Up @@ -27,7 +27,7 @@ class types, given the currently-registered extensions of this class.

def _idstr(self, sdfg, state, node):
""" Returns a unique identifier string from a node or state. """
result = str(sdfg.sdfg_id)
result = str(sdfg.cfg_id)
if state is not None:
result += '_' + str(sdfg.node_id(state))
if node is not None:
Expand Down
8 changes: 4 additions & 4 deletions dace/codegen/instrumentation/report.py
Original file line number Diff line number Diff line change
Expand Up @@ -16,7 +16,7 @@
def _uuid_to_dict(uuid: UUIDType) -> Dict[str, int]:
result = {}
if uuid[0] != -1:
result['sdfg_id'] = uuid[0]
result['cfg_id'] = uuid[0]
if uuid[1] != -1:
result['state_id'] = uuid[1]
if uuid[2] != -1:
Expand Down Expand Up @@ -83,13 +83,13 @@ def get_event_uuid_and_other_info(event) -> Tuple[UUIDType, Dict[str, Any]]:
other_info = {}
if 'args' in event:
args = event['args']
if 'sdfg_id' in args and args['sdfg_id'] is not None:
uuid = (args['sdfg_id'], -1, -1)
if 'cfg_id' in args and args['cfg_id'] is not None:
uuid = (args['cfg_id'], -1, -1)
if 'state_id' in args and args['state_id'] is not None:
uuid = (uuid[0], args['state_id'], -1)
if 'id' in args and args['id'] is not None:
uuid = (uuid[0], uuid[1], args['id'])
other_info = {k: v for k, v in args.items() if k not in ('sdfg_id', 'state_id', 'id')}
other_info = {k: v for k, v in args.items() if k not in ('cfg_id', 'state_id', 'id')}
return uuid, other_info

def __init__(self, filename: str):
Expand Down
4 changes: 2 additions & 2 deletions dace/codegen/instrumentation/timer.py
Original file line number Diff line number Diff line change
Expand Up @@ -40,8 +40,8 @@ def on_tend(self, timer_name: str, stream: CodeIOStream, sdfg=None, state=None,
stream.write('''auto __dace_tend_{id} = std::chrono::high_resolution_clock::now();
unsigned long int __dace_ts_start_{id} = std::chrono::duration_cast<std::chrono::microseconds>(__dace_tbegin_{id}.time_since_epoch()).count();
unsigned long int __dace_ts_end_{id} = std::chrono::duration_cast<std::chrono::microseconds>(__dace_tend_{id}.time_since_epoch()).count();
__state->report.add_completion("{timer_name}", "Timer", __dace_ts_start_{id}, __dace_ts_end_{id}, {sdfg_id}, {state_id}, {node_id});'''
.format(timer_name=timer_name, id=idstr, sdfg_id=sdfg.sdfg_id, state_id=state_id, node_id=node_id))
__state->report.add_completion("{timer_name}", "Timer", __dace_ts_start_{id}, __dace_ts_end_{id}, {cfg_id}, {state_id}, {node_id});'''
.format(timer_name=timer_name, id=idstr, cfg_id=sdfg.cfg_id, state_id=state_id, node_id=node_id))

# Code generation hooks
def on_state_begin(self, sdfg, state, local_stream, global_stream):
Expand Down
2 changes: 1 addition & 1 deletion dace/codegen/prettycode.py
Original file line number Diff line number Diff line change
Expand Up @@ -30,7 +30,7 @@ def write(self, contents, sdfg=None, state_id=None, node_id=None):

# If SDFG/state/node location is given, annotate this line
if sdfg is not None:
location_identifier = ' ////__DACE:%d' % sdfg.sdfg_id
location_identifier = ' ////__DACE:%d' % sdfg.cfg_id
if state_id is not None:
location_identifier += ':' + str(state_id)
if node_id is not None:
Expand Down
10 changes: 5 additions & 5 deletions dace/codegen/targets/cpp.py
Original file line number Diff line number Diff line change
Expand Up @@ -254,15 +254,15 @@ def ptr(name: str, desc: data.Data, sdfg: SDFG = None, framecode=None) -> str:
from dace.codegen.targets.cuda import CUDACodeGen # Avoid import loop

if desc.storage == dtypes.StorageType.CPU_ThreadLocal: # Use unambiguous name for thread-local arrays
return f'__{sdfg.sdfg_id}_{name}'
return f'__{sdfg.cfg_id}_{name}'
elif not CUDACodeGen._in_device_code: # GPU kernels cannot access state
return f'__state->__{sdfg.sdfg_id}_{name}'
return f'__state->__{sdfg.cfg_id}_{name}'
elif (sdfg, name) in framecode.where_allocated and framecode.where_allocated[(sdfg, name)] is not sdfg:
return f'__{sdfg.sdfg_id}_{name}'
return f'__{sdfg.cfg_id}_{name}'
elif (desc.transient and sdfg is not None and framecode is not None and (sdfg, name) in framecode.where_allocated
and framecode.where_allocated[(sdfg, name)] is not sdfg):
# Array allocated for another SDFG, use unambiguous name
return f'__{sdfg.sdfg_id}_{name}'
return f'__{sdfg.cfg_id}_{name}'

return name

Expand Down Expand Up @@ -914,7 +914,7 @@ def unparse_tasklet(sdfg, state_id, dfg, node, function_stream, callsite_stream,
# Doesn't cause crashes due to missing pyMLIR if a MLIR tasklet is not present
from dace.codegen.targets.mlir import utils

mlir_func_uid = "_" + str(sdfg.sdfg_id) + "_" + str(state_id) + "_" + str(dfg.node_id(node))
mlir_func_uid = "_" + str(sdfg.cfg_id) + "_" + str(state_id) + "_" + str(dfg.node_id(node))

mlir_ast = utils.get_ast(node.code.code)
mlir_is_generic = utils.is_generic(mlir_ast)
Expand Down
12 changes: 6 additions & 6 deletions dace/codegen/targets/cpu.py
Original file line number Diff line number Diff line change
Expand Up @@ -964,7 +964,7 @@ def process_out_memlets(self,
shared_data_name = edge.data.data
if not shared_data_name:
# Very unique name. TODO: Make more intuitive
shared_data_name = '__dace_%d_%d_%d_%d_%s' % (sdfg.sdfg_id, state_id, dfg.node_id(node),
shared_data_name = '__dace_%d_%d_%d_%d_%s' % (sdfg.cfg_id, state_id, dfg.node_id(node),
dfg.node_id(dst_node), edge.src_conn)

result.write(
Expand Down Expand Up @@ -1373,7 +1373,7 @@ def _generate_Tasklet(self, sdfg, dfg, state_id, node, function_stream, callsite
shared_data_name = edge.data.data
if not shared_data_name:
# Very unique name. TODO: Make more intuitive
shared_data_name = '__dace_%d_%d_%d_%d_%s' % (sdfg.sdfg_id, state_id, dfg.node_id(src_node),
shared_data_name = '__dace_%d_%d_%d_%d_%s' % (sdfg.cfg_id, state_id, dfg.node_id(src_node),
dfg.node_id(node), edge.src_conn)

# Read variable from shared storage
Expand Down Expand Up @@ -1442,7 +1442,7 @@ def _generate_Tasklet(self, sdfg, dfg, state_id, node, function_stream, callsite
local_name = edge.data.data
if not local_name:
# Very unique name. TODO: Make more intuitive
local_name = '__dace_%d_%d_%d_%d_%s' % (sdfg.sdfg_id, state_id, dfg.node_id(node),
local_name = '__dace_%d_%d_%d_%d_%s' % (sdfg.cfg_id, state_id, dfg.node_id(node),
dfg.node_id(dst_node), edge.src_conn)

# Allocate variable type
Expand Down Expand Up @@ -1545,7 +1545,7 @@ def generate_nsdfg_header(self, sdfg, state, state_id, node, memlet_references,
arguments = []

if state_struct:
toplevel_sdfg: SDFG = sdfg.sdfg_list[0]
toplevel_sdfg: SDFG = sdfg.cfg_list[0]
arguments.append(f'{cpp.mangle_dace_state_struct_name(toplevel_sdfg)} *__state')

# Add "__restrict__" keywords to arguments that do not alias with others in the context of this SDFG
Expand Down Expand Up @@ -1672,7 +1672,7 @@ def _generate_NestedSDFG(
# If the SDFG has a unique name, use it
sdfg_label = node.unique_name
else:
sdfg_label = "%s_%d_%d_%d" % (node.sdfg.name, sdfg.sdfg_id, state_id, dfg.node_id(node))
sdfg_label = "%s_%d_%d_%d" % (node.sdfg.name, sdfg.cfg_id, state_id, dfg.node_id(node))

code_already_generated = False
if unique_functions and not inline:
Expand Down Expand Up @@ -2063,7 +2063,7 @@ def _generate_ConsumeEntry(
ctype = node.out_connectors[edge.src_conn].ctype
if not local_name:
# Very unique name. TODO: Make more intuitive
local_name = '__dace_%d_%d_%d_%d_%s' % (sdfg.sdfg_id, state_id, dfg.node_id(
local_name = '__dace_%d_%d_%d_%d_%s' % (sdfg.cfg_id, state_id, dfg.node_id(
edge.src), dfg.node_id(edge.dst), edge.src_conn)

# Allocate variable type
Expand Down
Loading

0 comments on commit b5ec059

Please sign in to comment.