{ "cells": [ { "cell_type": "markdown", "metadata": {}, "source": [ "# Customizing Code Generation\n", "\n", "Writing custom backends in DaCe is a powerful feature that allows users to specify their own backends. This could range from customizing different support libraries, through new target architectures and use of specialized hardware, to outputting a completely different language than C.\n", "\n", "This is made possible due to the modular enumerations and code generation architecture in DaCe. Any enumeration (e.g., for storage types and schedulers) can be extended through the API, in order to enable custom, project-specific behavior, and code generators (similarly to transformations) have a registry that can also be extended at runtime. \n", "\n", "Code generation generally follows a recursive procedure, starting from the top-level SDFG:\n", " * Generate frame-code and entry point for SDFG\n", " * Generate code for array allocation/deallocation based on lifetime (SDFG-wide and persistent lifetime)\n", " * Generate control flow between states\n", " * Generate each state/scope:\n", " * Generate additional code for array allocation/deallocation based on lifetime (state- or scope-wide)\n", " * Generate data movement code for each memlet path\n", " * Generate each node in a topological, hierarchical order (descending into scopes such as Maps)\n", " * If a node is a scope entry, generate code for the scope as a subgraph\n", " * If a node is a nested SDFG, generate code for that SDFG as a separate function and call it\n", " \n", "In this tutorial, we will show how to customize the code generation procedure by creating a new map scheduler and generating custom code for it. First, we will import dace and some classes to make the rest of the code cleaner:" ] }, { "cell_type": "code", "execution_count": 1, "metadata": {}, "outputs": [ { "data": { "text/html": [ "\n" ], "text/plain": [ "" ] }, "metadata": {}, "output_type": "display_data" } ], "source": [ "import dace\n", "from dace import registry\n", "from dace.sdfg.scope import ScopeSubgraphView\n", "from dace.codegen.prettycode import CodeIOStream\n", "from dace.codegen.targets.target import TargetCodeGenerator\n", "from dace.codegen.targets.framecode import DaCeCodeGenerator\n", "from dace.codegen.targets.cpp import sym2cpp" ] }, { "cell_type": "markdown", "metadata": {}, "source": [ "Next, we can define some simple program to work with:" ] }, { "cell_type": "code", "execution_count": 2, "metadata": {}, "outputs": [ { "data": { "text/html": [ "\n", "
\n", "
\n", "
\n", "\n", "" ], "text/plain": [ "SDFG (simple)" ] }, "execution_count": 2, "metadata": {}, "output_type": "execute_result" } ], "source": [ "@dace.program\n", "def simple(A: dace.float64[20, 30]):\n", " for i, j in dace.map[0:20:2, 0:30]:\n", " A[i, j] += A[i, j]\n", "\n", "# Preview SDFG\n", "sdfg = simple.to_sdfg()\n", "sdfg" ] }, { "cell_type": "markdown", "metadata": {}, "source": [ "If we observe the generated code, by default our map would be scheduled to an OpenMP multi-core loop:" ] }, { "cell_type": "code", "execution_count": 3, "metadata": {}, "outputs": [ { "data": { "text/html": [ "
/* DaCe AUTO-GENERATED FILE. DO NOT MODIFY */\n",
       "#include <dace/dace.h>\n",
       "\n",
       "void __program_simple_internal(double * __restrict__ A)\n",
       "{\n",
       "\n",
       "    {\n",
       "        \n",
       "        \n",
       "        {\n",
       "            #pragma omp parallel for\n",
       "            for (auto i = 0; i < 20; i += 2) {\n",
       "                for (auto j = 0; j < 30; j += 1) {\n",
       "                    {\n",
       "                        double __in1 = A[((30 * i) + j)];\n",
       "                        double __in2 = A[((30 * i) + j)];\n",
       "                        double __out;\n",
       "\n",
       "                        ///////////////////\n",
       "                        // Tasklet code (augassign_4_8)\n",
       "                        __out = (__in1 + __in2);\n",
       "                        ///////////////////\n",
       "\n",
       "                        A[((30 * i) + j)] = __out;\n",
       "                    }\n",
       "                }\n",
       "            }\n",
       "        }\n",
       "    }\n",
       "}\n",
       "\n",
       "DACE_EXPORTED void __program_simple(double * __restrict__ A)\n",
       "{\n",
       "    __program_simple_internal(A);\n",
       "}\n",
       "\n",
       "DACE_EXPORTED int __dace_init_simple(double * __restrict__ A)\n",
       "{\n",
       "    int __result = 0;\n",
       "\n",
       "    return __result;\n",
       "}\n",
       "\n",
       "DACE_EXPORTED void __dace_exit_simple(double * __restrict__ A)\n",
       "{\n",
       "}\n",
       "
\n" ], "text/latex": [ "\\begin{Verbatim}[commandchars=\\\\\\{\\}]\n", "\\PY{c+cm}{/* DaCe AUTO\\PYZhy{}GENERATED FILE. DO NOT MODIFY */}\n", "\\PY{c+cp}{\\PYZsh{}}\\PY{c+cp}{include} \\PY{c+cpf}{\\PYZlt{}dace/dace.h\\PYZgt{}}\n", "\n", "\\PY{k+kt}{void} \\PY{n+nf}{\\PYZus{}\\PYZus{}program\\PYZus{}simple\\PYZus{}internal}\\PY{p}{(}\\PY{k+kt}{double} \\PY{o}{*} \\PY{n}{\\PYZus{}\\PYZus{}restrict\\PYZus{}\\PYZus{}} \\PY{n}{A}\\PY{p}{)}\n", "\\PY{p}{\\PYZob{}}\n", "\n", " \\PY{p}{\\PYZob{}}\n", " \n", " \n", " \\PY{p}{\\PYZob{}}\n", " \\PY{c+cp}{\\PYZsh{}}\\PY{c+cp}{pragma omp parallel for}\n", " \\PY{k}{for} \\PY{p}{(}\\PY{k}{auto} \\PY{n}{i} \\PY{o}{=} \\PY{l+m+mi}{0}\\PY{p}{;} \\PY{n}{i} \\PY{o}{\\PYZlt{}} \\PY{l+m+mi}{20}\\PY{p}{;} \\PY{n}{i} \\PY{o}{+}\\PY{o}{=} \\PY{l+m+mi}{2}\\PY{p}{)} \\PY{p}{\\PYZob{}}\n", " \\PY{k}{for} \\PY{p}{(}\\PY{k}{auto} \\PY{n}{j} \\PY{o}{=} \\PY{l+m+mi}{0}\\PY{p}{;} \\PY{n}{j} \\PY{o}{\\PYZlt{}} \\PY{l+m+mi}{30}\\PY{p}{;} \\PY{n}{j} \\PY{o}{+}\\PY{o}{=} \\PY{l+m+mi}{1}\\PY{p}{)} \\PY{p}{\\PYZob{}}\n", " \\PY{p}{\\PYZob{}}\n", " \\PY{k+kt}{double} \\PY{n}{\\PYZus{}\\PYZus{}in1} \\PY{o}{=} \\PY{n}{A}\\PY{p}{[}\\PY{p}{(}\\PY{p}{(}\\PY{l+m+mi}{30} \\PY{o}{*} \\PY{n}{i}\\PY{p}{)} \\PY{o}{+} \\PY{n}{j}\\PY{p}{)}\\PY{p}{]}\\PY{p}{;}\n", " \\PY{k+kt}{double} \\PY{n}{\\PYZus{}\\PYZus{}in2} \\PY{o}{=} \\PY{n}{A}\\PY{p}{[}\\PY{p}{(}\\PY{p}{(}\\PY{l+m+mi}{30} \\PY{o}{*} \\PY{n}{i}\\PY{p}{)} \\PY{o}{+} \\PY{n}{j}\\PY{p}{)}\\PY{p}{]}\\PY{p}{;}\n", " \\PY{k+kt}{double} \\PY{n}{\\PYZus{}\\PYZus{}out}\\PY{p}{;}\n", "\n", " \\PY{c+c1}{///////////////////}\n", " \\PY{c+c1}{// Tasklet code (augassign\\PYZus{}4\\PYZus{}8)}\n", " \\PY{n}{\\PYZus{}\\PYZus{}out} \\PY{o}{=} \\PY{p}{(}\\PY{n}{\\PYZus{}\\PYZus{}in1} \\PY{o}{+} \\PY{n}{\\PYZus{}\\PYZus{}in2}\\PY{p}{)}\\PY{p}{;}\n", " \\PY{c+c1}{///////////////////}\n", "\n", " \\PY{n}{A}\\PY{p}{[}\\PY{p}{(}\\PY{p}{(}\\PY{l+m+mi}{30} \\PY{o}{*} \\PY{n}{i}\\PY{p}{)} \\PY{o}{+} \\PY{n}{j}\\PY{p}{)}\\PY{p}{]} \\PY{o}{=} \\PY{n}{\\PYZus{}\\PYZus{}out}\\PY{p}{;}\n", " \\PY{p}{\\PYZcb{}}\n", " \\PY{p}{\\PYZcb{}}\n", " \\PY{p}{\\PYZcb{}}\n", " \\PY{p}{\\PYZcb{}}\n", " \\PY{p}{\\PYZcb{}}\n", "\\PY{p}{\\PYZcb{}}\n", "\n", "\\PY{n}{DACE\\PYZus{}EXPORTED} \\PY{k+kt}{void} \\PY{n+nf}{\\PYZus{}\\PYZus{}program\\PYZus{}simple}\\PY{p}{(}\\PY{k+kt}{double} \\PY{o}{*} \\PY{n}{\\PYZus{}\\PYZus{}restrict\\PYZus{}\\PYZus{}} \\PY{n}{A}\\PY{p}{)}\n", "\\PY{p}{\\PYZob{}}\n", " \\PY{n}{\\PYZus{}\\PYZus{}program\\PYZus{}simple\\PYZus{}internal}\\PY{p}{(}\\PY{n}{A}\\PY{p}{)}\\PY{p}{;}\n", "\\PY{p}{\\PYZcb{}}\n", "\n", "\\PY{n}{DACE\\PYZus{}EXPORTED} \\PY{k+kt}{int} \\PY{n+nf}{\\PYZus{}\\PYZus{}dace\\PYZus{}init\\PYZus{}simple}\\PY{p}{(}\\PY{k+kt}{double} \\PY{o}{*} \\PY{n}{\\PYZus{}\\PYZus{}restrict\\PYZus{}\\PYZus{}} \\PY{n}{A}\\PY{p}{)}\n", "\\PY{p}{\\PYZob{}}\n", " \\PY{k+kt}{int} \\PY{n}{\\PYZus{}\\PYZus{}result} \\PY{o}{=} \\PY{l+m+mi}{0}\\PY{p}{;}\n", "\n", " \\PY{k}{return} \\PY{n}{\\PYZus{}\\PYZus{}result}\\PY{p}{;}\n", "\\PY{p}{\\PYZcb{}}\n", "\n", "\\PY{n}{DACE\\PYZus{}EXPORTED} \\PY{k+kt}{void} \\PY{n+nf}{\\PYZus{}\\PYZus{}dace\\PYZus{}exit\\PYZus{}simple}\\PY{p}{(}\\PY{k+kt}{double} \\PY{o}{*} \\PY{n}{\\PYZus{}\\PYZus{}restrict\\PYZus{}\\PYZus{}} \\PY{n}{A}\\PY{p}{)}\n", "\\PY{p}{\\PYZob{}}\n", "\\PY{p}{\\PYZcb{}}\n", "\\end{Verbatim}\n" ], "text/plain": [ "/* DaCe AUTO-GENERATED FILE. DO NOT MODIFY */\n", "#include \n", "\n", "void __program_simple_internal(double * __restrict__ A)\n", "{\n", "\n", " {\n", " \n", " \n", " {\n", " #pragma omp parallel for\n", " for (auto i = 0; i < 20; i += 2) {\n", " for (auto j = 0; j < 30; j += 1) {\n", " {\n", " double __in1 = A[((30 * i) + j)];\n", " double __in2 = A[((30 * i) + j)];\n", " double __out;\n", "\n", " ///////////////////\n", " // Tasklet code (augassign_4_8)\n", " __out = (__in1 + __in2);\n", " ///////////////////\n", "\n", " A[((30 * i) + j)] = __out;\n", " }\n", " }\n", " }\n", " }\n", " }\n", "}\n", "\n", "DACE_EXPORTED void __program_simple(double * __restrict__ A)\n", "{\n", " __program_simple_internal(A);\n", "}\n", "\n", "DACE_EXPORTED int __dace_init_simple(double * __restrict__ A)\n", "{\n", " int __result = 0;\n", "\n", " return __result;\n", "}\n", "\n", "DACE_EXPORTED void __dace_exit_simple(double * __restrict__ A)\n", "{\n", "}\n" ] }, "execution_count": 3, "metadata": {}, "output_type": "execute_result" } ], "source": [ "from IPython.display import Code\n", "Code(sdfg.generate_code()[0].clean_code, language='cpp')" ] }, { "cell_type": "markdown", "metadata": {}, "source": [ "## Custom Scope Generator\n", "\n", "To begin, we need to add our own enum entries for our new map schedule. Here are the current schedule types in our extensible enumeration:" ] }, { "cell_type": "code", "execution_count": 4, "metadata": {}, "outputs": [ { "data": { "text/plain": [ "[,\n", " ,\n", " ,\n", " ,\n", " ,\n", " ,\n", " ,\n", " ,\n", " ]" ] }, "execution_count": 4, "metadata": {}, "output_type": "execute_result" } ], "source": [ "list(dace.ScheduleType)" ] }, { "cell_type": "markdown", "metadata": {}, "source": [ "Registering a new value is just a matter of calling `register`:" ] }, { "cell_type": "code", "execution_count": 5, "metadata": {}, "outputs": [ { "data": { "text/plain": [ "[,\n", " ,\n", " ,\n", " ,\n", " ,\n", " ,\n", " ,\n", " ,\n", " ,\n", " ]" ] }, "execution_count": 5, "metadata": {}, "output_type": "execute_result" } ], "source": [ "dace.ScheduleType.register('LoopyLoop')\n", "list(dace.ScheduleType)" ] }, { "cell_type": "markdown", "metadata": {}, "source": [ "It can also be used directly from now on:" ] }, { "cell_type": "code", "execution_count": 6, "metadata": {}, "outputs": [ { "data": { "text/plain": [ "" ] }, "execution_count": 6, "metadata": {}, "output_type": "execute_result" } ], "source": [ "dace.ScheduleType.LoopyLoop" ] }, { "cell_type": "markdown", "metadata": {}, "source": [ "One additional step for code generation is to tell the code generator which arrays and scopes lie inside by default:" ] }, { "cell_type": "code", "execution_count": 7, "metadata": {}, "outputs": [], "source": [ "dace.SCOPEDEFAULT_SCHEDULE[dace.ScheduleType.LoopyLoop] = dace.ScheduleType.Sequential\n", "dace.SCOPEDEFAULT_STORAGE[dace.ScheduleType.LoopyLoop] = dace.StorageType.CPU_Heap" ] }, { "cell_type": "markdown", "metadata": {}, "source": [ "Now we can register and create a matching code generator:" ] }, { "cell_type": "code", "execution_count": 8, "metadata": {}, "outputs": [], "source": [ "@registry.autoregister_params(name='loopy')\n", "class MyCustomLoop(TargetCodeGenerator):\n", " def __init__(self, frame_codegen: DaCeCodeGenerator, sdfg: dace.SDFG):\n", " ################################################################\n", " # Define some locals:\n", " # Can be used to call back to the frame-code generator\n", " self.frame = frame_codegen\n", " # Can be used to dispatch other code generators for allocation/nodes\n", " self.dispatcher = frame_codegen.dispatcher\n", " \n", " ################################################################\n", " # Register handlers/hooks through dispatcher: Can be used for \n", " # nodes, memory copy/allocation, scopes, states, and more.\n", " \n", " # In this case, register scopes\n", " self.dispatcher.register_map_dispatcher(dace.ScheduleType.LoopyLoop, self)\n", " \n", " # You can similarly use register_{array,copy,node,state}_dispatcher\n", " \n", " # A scope dispatcher will trigger a method called generate_scope whenever \n", " # an SDFG has a scope with that schedule\n", " def generate_scope(self, sdfg: dace.SDFG, scope: ScopeSubgraphView,\n", " state_id: int, function_stream: CodeIOStream,\n", " callsite_stream: CodeIOStream):\n", " # The parameters here are:\n", " # sdfg: The SDFG we are currently generating.\n", " # scope: The subgraph of the state containing only the scope (map contents)\n", " # we want to generate the code for.\n", " # state_id: The state in the SDFG the subgraph is taken from (i.e., \n", " # `sdfg.node(state_id)` is the same as `scope.graph`)\n", " # function_stream: A cursor to the global code (which can be used to define\n", " # functions, hence the name).\n", " # callsite_stream: A cursor to the current location in the code, most of\n", " # the code is generated here.\n", " \n", " # We can get the map entry node from the scope graph\n", " entry_node = scope.source_nodes()[0]\n", " \n", " # First, generate an opening brace (for instrumentation and dynamic map ranges)\n", " callsite_stream.write('{', sdfg, state_id, entry_node)\n", " \n", " ################################################################\n", " # Generate specific code: We will generate a reversed loop with a \n", " # comment for each dimension of the map. For the sake of simplicity,\n", " # dynamic map ranges are not supported.\n", " \n", " for param, rng in zip(entry_node.map.params, entry_node.map.range):\n", " # We use the sym2cpp function from the cpp support functions\n", " # to convert symbolic expressions to proper C++\n", " begin, end, stride = (sym2cpp(r) for r in rng)\n", " \n", " # Every write is optionally (but recommended to be) tagged with\n", " # 1-3 extra arguments, serving as line information to match\n", " # SDFG, state, and graph nodes/edges to written code.\n", " callsite_stream.write(f'''// Loopy-loop {param}\n", " for (auto {param} = {end}; {param} >= {begin}; {param} -= {stride}) {{''',\n", " sdfg, state_id, entry_node\n", " )\n", " \n", " # NOTE: CodeIOStream will automatically take care of indentation for us.\n", " \n", " \n", " # Now that the loops have been defined, use the dispatcher to invoke any\n", " # code generator (including this one) that is registered to deal with\n", " # the internal nodes in the subgraph. We skip the MapEntry node.\n", " self.dispatcher.dispatch_subgraph(sdfg, scope, state_id,\n", " function_stream, callsite_stream,\n", " skip_entry_node=True)\n", " \n", " # NOTE: Since skip_exit_node above is set to False, closing braces will\n", " # be automatically generated" ] }, { "cell_type": "markdown", "metadata": {}, "source": [ "After the code generator has been registered, all that's left is to change the map schedule and generate new code:" ] }, { "cell_type": "code", "execution_count": 9, "metadata": {}, "outputs": [ { "data": { "text/html": [ "
/* DaCe AUTO-GENERATED FILE. DO NOT MODIFY */\n",
       "#include <dace/dace.h>\n",
       "\n",
       "void __program_simple_internal(double * __restrict__ A)\n",
       "{\n",
       "\n",
       "    {\n",
       "        \n",
       "        \n",
       "        // Loopy-loop i\n",
       "        for (auto i = 19; i >= 0; i -= 2) {\n",
       "            // Loopy-loop j\n",
       "            for (auto j = 29; j >= 0; j -= 1) {\n",
       "                {\n",
       "                    double __in1 = A[((30 * i) + j)];\n",
       "                    double __in2 = A[((30 * i) + j)];\n",
       "                    double __out;\n",
       "\n",
       "                    ///////////////////\n",
       "                    // Tasklet code (augassign_4_8)\n",
       "                    __out = (__in1 + __in2);\n",
       "                    ///////////////////\n",
       "\n",
       "                    A[((30 * i) + j)] = __out;\n",
       "                }\n",
       "            }\n",
       "        }\n",
       "    }\n",
       "}\n",
       "}\n",
       "\n",
       "DACE_EXPORTED void __program_simple(double * __restrict__ A)\n",
       "{\n",
       "__program_simple_internal(A);\n",
       "}\n",
       "\n",
       "DACE_EXPORTED int __dace_init_simple(double * __restrict__ A)\n",
       "{\n",
       "int __result = 0;\n",
       "\n",
       "return __result;\n",
       "}\n",
       "\n",
       "DACE_EXPORTED void __dace_exit_simple(double * __restrict__ A)\n",
       "{\n",
       "}\n",
       "
\n" ], "text/latex": [ "\\begin{Verbatim}[commandchars=\\\\\\{\\}]\n", "\\PY{c+cm}{/* DaCe AUTO\\PYZhy{}GENERATED FILE. DO NOT MODIFY */}\n", "\\PY{c+cp}{\\PYZsh{}}\\PY{c+cp}{include} \\PY{c+cpf}{\\PYZlt{}dace/dace.h\\PYZgt{}}\n", "\n", "\\PY{k+kt}{void} \\PY{n+nf}{\\PYZus{}\\PYZus{}program\\PYZus{}simple\\PYZus{}internal}\\PY{p}{(}\\PY{k+kt}{double} \\PY{o}{*} \\PY{n}{\\PYZus{}\\PYZus{}restrict\\PYZus{}\\PYZus{}} \\PY{n}{A}\\PY{p}{)}\n", "\\PY{p}{\\PYZob{}}\n", "\n", " \\PY{p}{\\PYZob{}}\n", " \n", " \n", " \\PY{c+c1}{// Loopy\\PYZhy{}loop i}\n", " \\PY{k}{for} \\PY{p}{(}\\PY{k}{auto} \\PY{n}{i} \\PY{o}{=} \\PY{l+m+mi}{19}\\PY{p}{;} \\PY{n}{i} \\PY{o}{\\PYZgt{}}\\PY{o}{=} \\PY{l+m+mi}{0}\\PY{p}{;} \\PY{n}{i} \\PY{o}{\\PYZhy{}}\\PY{o}{=} \\PY{l+m+mi}{2}\\PY{p}{)} \\PY{p}{\\PYZob{}}\n", " \\PY{c+c1}{// Loopy\\PYZhy{}loop j}\n", " \\PY{k}{for} \\PY{p}{(}\\PY{k}{auto} \\PY{n}{j} \\PY{o}{=} \\PY{l+m+mi}{29}\\PY{p}{;} \\PY{n}{j} \\PY{o}{\\PYZgt{}}\\PY{o}{=} \\PY{l+m+mi}{0}\\PY{p}{;} \\PY{n}{j} \\PY{o}{\\PYZhy{}}\\PY{o}{=} \\PY{l+m+mi}{1}\\PY{p}{)} \\PY{p}{\\PYZob{}}\n", " \\PY{p}{\\PYZob{}}\n", " \\PY{k+kt}{double} \\PY{n}{\\PYZus{}\\PYZus{}in1} \\PY{o}{=} \\PY{n}{A}\\PY{p}{[}\\PY{p}{(}\\PY{p}{(}\\PY{l+m+mi}{30} \\PY{o}{*} \\PY{n}{i}\\PY{p}{)} \\PY{o}{+} \\PY{n}{j}\\PY{p}{)}\\PY{p}{]}\\PY{p}{;}\n", " \\PY{k+kt}{double} \\PY{n}{\\PYZus{}\\PYZus{}in2} \\PY{o}{=} \\PY{n}{A}\\PY{p}{[}\\PY{p}{(}\\PY{p}{(}\\PY{l+m+mi}{30} \\PY{o}{*} \\PY{n}{i}\\PY{p}{)} \\PY{o}{+} \\PY{n}{j}\\PY{p}{)}\\PY{p}{]}\\PY{p}{;}\n", " \\PY{k+kt}{double} \\PY{n}{\\PYZus{}\\PYZus{}out}\\PY{p}{;}\n", "\n", " \\PY{c+c1}{///////////////////}\n", " \\PY{c+c1}{// Tasklet code (augassign\\PYZus{}4\\PYZus{}8)}\n", " \\PY{n}{\\PYZus{}\\PYZus{}out} \\PY{o}{=} \\PY{p}{(}\\PY{n}{\\PYZus{}\\PYZus{}in1} \\PY{o}{+} \\PY{n}{\\PYZus{}\\PYZus{}in2}\\PY{p}{)}\\PY{p}{;}\n", " \\PY{c+c1}{///////////////////}\n", "\n", " \\PY{n}{A}\\PY{p}{[}\\PY{p}{(}\\PY{p}{(}\\PY{l+m+mi}{30} \\PY{o}{*} \\PY{n}{i}\\PY{p}{)} \\PY{o}{+} \\PY{n}{j}\\PY{p}{)}\\PY{p}{]} \\PY{o}{=} \\PY{n}{\\PYZus{}\\PYZus{}out}\\PY{p}{;}\n", " \\PY{p}{\\PYZcb{}}\n", " \\PY{p}{\\PYZcb{}}\n", " \\PY{p}{\\PYZcb{}}\n", " \\PY{p}{\\PYZcb{}}\n", "\\PY{p}{\\PYZcb{}}\n", "\\PY{p}{\\PYZcb{}}\n", "\n", "\\PY{n}{DACE\\PYZus{}EXPORTED} \\PY{k+kt}{void} \\PY{n}{\\PYZus{}\\PYZus{}program\\PYZus{}simple}\\PY{p}{(}\\PY{k+kt}{double} \\PY{o}{*} \\PY{n}{\\PYZus{}\\PYZus{}restrict\\PYZus{}\\PYZus{}} \\PY{n}{A}\\PY{p}{)}\n", "\\PY{p}{\\PYZob{}}\n", "\\PY{n}{\\PYZus{}\\PYZus{}program\\PYZus{}simple\\PYZus{}internal}\\PY{p}{(}\\PY{n}{A}\\PY{p}{)}\\PY{p}{;}\n", "\\PY{p}{\\PYZcb{}}\n", "\n", "\\PY{n}{DACE\\PYZus{}EXPORTED} \\PY{k+kt}{int} \\PY{n}{\\PYZus{}\\PYZus{}dace\\PYZus{}init\\PYZus{}simple}\\PY{p}{(}\\PY{k+kt}{double} \\PY{o}{*} \\PY{n}{\\PYZus{}\\PYZus{}restrict\\PYZus{}\\PYZus{}} \\PY{n}{A}\\PY{p}{)}\n", "\\PY{p}{\\PYZob{}}\n", "\\PY{k+kt}{int} \\PY{n}{\\PYZus{}\\PYZus{}result} \\PY{o}{=} \\PY{l+m+mi}{0}\\PY{p}{;}\n", "\n", "\\PY{k}{return} \\PY{n}{\\PYZus{}\\PYZus{}result}\\PY{p}{;}\n", "\\PY{p}{\\PYZcb{}}\n", "\n", "\\PY{n}{DACE\\PYZus{}EXPORTED} \\PY{k+kt}{void} \\PY{n}{\\PYZus{}\\PYZus{}dace\\PYZus{}exit\\PYZus{}simple}\\PY{p}{(}\\PY{k+kt}{double} \\PY{o}{*} \\PY{n}{\\PYZus{}\\PYZus{}restrict\\PYZus{}\\PYZus{}} \\PY{n}{A}\\PY{p}{)}\n", "\\PY{p}{\\PYZob{}}\n", "\\PY{p}{\\PYZcb{}}\n", "\\end{Verbatim}\n" ], "text/plain": [ "/* DaCe AUTO-GENERATED FILE. DO NOT MODIFY */\n", "#include \n", "\n", "void __program_simple_internal(double * __restrict__ A)\n", "{\n", "\n", " {\n", " \n", " \n", " // Loopy-loop i\n", " for (auto i = 19; i >= 0; i -= 2) {\n", " // Loopy-loop j\n", " for (auto j = 29; j >= 0; j -= 1) {\n", " {\n", " double __in1 = A[((30 * i) + j)];\n", " double __in2 = A[((30 * i) + j)];\n", " double __out;\n", "\n", " ///////////////////\n", " // Tasklet code (augassign_4_8)\n", " __out = (__in1 + __in2);\n", " ///////////////////\n", "\n", " A[((30 * i) + j)] = __out;\n", " }\n", " }\n", " }\n", " }\n", "}\n", "}\n", "\n", "DACE_EXPORTED void __program_simple(double * __restrict__ A)\n", "{\n", "__program_simple_internal(A);\n", "}\n", "\n", "DACE_EXPORTED int __dace_init_simple(double * __restrict__ A)\n", "{\n", "int __result = 0;\n", "\n", "return __result;\n", "}\n", "\n", "DACE_EXPORTED void __dace_exit_simple(double * __restrict__ A)\n", "{\n", "}\n" ] }, "execution_count": 9, "metadata": {}, "output_type": "execute_result" } ], "source": [ "# Change schedule\n", "for node, _ in sdfg.all_nodes_recursive():\n", " if isinstance(node, dace.nodes.MapEntry):\n", " node.schedule = dace.ScheduleType.LoopyLoop\n", "\n", "Code(sdfg.generate_code()[0].clean_code, language='cpp')" ] }, { "cell_type": "markdown", "metadata": {}, "source": [ "and the code is generated appropriately." ] } ], "metadata": { "kernelspec": { "display_name": "Python 3", "language": "python", "name": "python3" }, "language_info": { "codemirror_mode": { "name": "ipython", "version": 3 }, "file_extension": ".py", "mimetype": "text/x-python", "name": "python", "nbconvert_exporter": "python", "pygments_lexer": "ipython3", "version": "3.12.1" } }, "nbformat": 4, "nbformat_minor": 4 }