Loop doacross

dace/codegen/codegen.py

@@ -1,7 +1,7 @@
 # Copyright 2019-2021 ETH Zurich and the DaCe authors. All rights reserved.
 import functools
 import os
-from typing import List, Set
+from typing import Dict, List, Set
 
 import dace
 from dace import dtypes
@@ -95,7 +95,7 @@ def _get_codegen_targets(sdfg: SDFG, frame: framecode.DaCeCodeGenerator):
     for node, parent in sdfg.all_nodes_recursive():
         # Query nodes and scopes
         if isinstance(node, SDFGState):
-            frame.targets.add(disp.get_state_dispatcher(parent, node))
+            frame.targets.add(disp.get_state_dispatcher(node.sdfg, node))
         elif isinstance(node, dace.nodes.EntryNode):
             frame.targets.add(disp.get_scope_dispatcher(node.schedule))
         elif isinstance(node, dace.nodes.Node):
@@ -149,7 +149,7 @@ def _get_codegen_targets(sdfg: SDFG, frame: framecode.DaCeCodeGenerator):
         disp.instrumentation[sdfg.instrument] = provider_mapping[sdfg.instrument]
 
 
-def generate_code(sdfg, validate=True) -> List[CodeObject]:
+def generate_code(sdfg: SDFG, validate=True) -> List[CodeObject]:
     """
     Generates code as a list of code objects for a given SDFG.
 
@@ -189,6 +189,7 @@ def generate_code(sdfg, validate=True) -> List[CodeObject]:
     # Convert any loop constructs with hierarchical loop regions into simple 1-level state machine loops.
     # TODO (later): Adapt codegen to deal with hierarchical CFGs instead.
     sdutils.inline_loop_blocks(sdfg)
+    sdutils.inline_control_flow_regions(sdfg)
 
     # Before generating the code, run type inference on the SDFG connectors
     infer_types.infer_connector_types(sdfg)
@@ -207,12 +208,12 @@ def generate_code(sdfg, validate=True) -> List[CodeObject]:
 
     # Instantiate CPU first (as it is used by the other code generators)
     # TODO: Refactor the parts used by other code generators out of CPU
-    default_target = cpu.CPUCodeGen
+    default_target: TargetCodeGenerator = cpu.CPUCodeGen
     for k, v in TargetCodeGenerator.extensions().items():
         # If another target has already been registered as CPU, use it instead
         if v['name'] == 'cpu':
             default_target = k
-    targets = {'cpu': default_target(frame, sdfg)}
+    targets: Dict[str, TargetCodeGenerator] = {'cpu': default_target(frame, sdfg)}
 
     # Instantiate the rest of the targets
     targets.update(

dace/codegen/control_flow.py

@@ -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)
@@ -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,
@@ -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'
@@ -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
 
@@ -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
@@ -390,7 +390,8 @@ def as_cpp(self, codegen, symbols) -> str:
 
         update = ''
         if self.update is not None:
-            update = f'{self.itervar} = {self.update}'
+            cppupdate = unparse_interstate_edge(self.update, sdfg, codegen=codegen)
+            update = f'{self.itervar} = {cppupdate}'
 
         expr = f'{preinit}\nfor ({init}; {cond}; {update}) {{\n'
         expr += _clean_loop_body(self.body.as_cpp(codegen, symbols))
@@ -475,7 +476,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
 
@@ -841,7 +842,7 @@ def structured_control_flow_tree(sdfg: SDFG, dispatch_state: Callable[[SDFGState
     from dace.sdfg.analysis import cfg
 
     # Get parent states and back-edges
-    ptree = cfg.state_parent_tree(sdfg)
+    ptree = cfg.control_flow_block_parent_tree(sdfg)
     back_edges = cfg.back_edges(sdfg)
 
     # Annotate branches

dace/codegen/dispatcher.py

@@ -1,4 +1,4 @@
-# Copyright 2019-2021 ETH Zurich and the DaCe authors. All rights reserved.
+# Copyright 2019-2024 ETH Zurich and the DaCe authors. All rights reserved.
 """
 Contains the DaCe code generator target dispatcher, which is responsible for
 flexible code generation with multiple backends by dispatching certain

dace/codegen/instrumentation/data/data_dump.py

@@ -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 = '', ''
@@ -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 = '', ''

dace/codegen/instrumentation/gpu_events.py

@@ -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)
 

dace/codegen/instrumentation/likwid.py

@@ -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
 {{
@@ -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++)
@@ -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});
             }}
         }}
     }}
@@ -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
@@ -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
@@ -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
 {{
@@ -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
@@ -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));
@@ -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);
@@ -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}");
@@ -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}");
@@ -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}");
 
@@ -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}");

dace/codegen/instrumentation/provider.py

@@ -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:

dace/codegen/instrumentation/report.py

@@ -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:
@@ -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):

dace/codegen/instrumentation/timer.py

@@ -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):

dace/codegen/prettycode.py

@@ -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: