diff --git a/dace/codegen/compiled_sdfg.py b/dace/codegen/compiled_sdfg.py index 9ee0772eeb..22f95d01d7 100644 --- a/dace/codegen/compiled_sdfg.py +++ b/dace/codegen/compiled_sdfg.py @@ -287,6 +287,7 @@ def get_workspace_sizes(self) -> Dict[dtypes.StorageType, int]: result: Dict[dtypes.StorageType, int] = {} for storage in self.external_memory_types: func = self._lib.get_symbol(f'__dace_get_external_memory_size_{storage.name}') + func.restype = ctypes.c_size_t result[storage] = func(self._libhandle, *self._lastargs[1]) return result diff --git a/dace/codegen/cppunparse.py b/dace/codegen/cppunparse.py index eae0ed229e..77dd34d478 100644 --- a/dace/codegen/cppunparse.py +++ b/dace/codegen/cppunparse.py @@ -78,6 +78,7 @@ import numpy as np import os import tokenize +import warnings import sympy import dace @@ -733,6 +734,21 @@ def _Num(self, t): if isinstance(t.n, complex): dtype = dtypes.DTYPE_TO_TYPECLASS[complex] + # Handle large integer values + if isinstance(t.n, int): + bits = t.n.bit_length() + if bits == 32: # Integer, potentially unsigned + if t.n >= 0: # unsigned + repr_n += 'U' + else: # signed, 64-bit + repr_n += 'LL' + elif 32 < bits <= 63: + repr_n += 'LL' + elif bits == 64 and t.n >= 0: + repr_n += 'ULL' + elif bits >= 64: + warnings.warn(f'Value wider than 64 bits encountered in expression ({t.n}), emitting as-is') + if repr_n.endswith("j"): self.write("%s(0, %s)" % (dtype, repr_n.replace("inf", INFSTR)[:-1])) else: @@ -831,8 +847,23 @@ def _Tuple( self.write(")") unop = {"Invert": "~", "Not": "!", "UAdd": "+", "USub": "-"} + unop_lambda = {'Invert': (lambda x: ~x), 'Not': (lambda x: not x), 'UAdd': (lambda x: +x), 'USub': (lambda x: -x)} def _UnaryOp(self, t): + # Dispatch constants after applying the operation + if sys.version_info[:2] < (3, 8): + if isinstance(t.operand, ast.Num): + newval = self.unop_lambda[t.op.__class__.__name__](t.operand.n) + newnode = ast.Num(n=newval) + self.dispatch(newnode) + return + else: + if isinstance(t.operand, ast.Constant): + newval = self.unop_lambda[t.op.__class__.__name__](t.operand.value) + newnode = ast.Constant(value=newval) + self.dispatch(newnode) + return + self.write("(") self.write(self.unop[t.op.__class__.__name__]) self.write(" ") diff --git a/dace/codegen/targets/cuda.py b/dace/codegen/targets/cuda.py index ee49f04d03..a465d2bbc0 100644 --- a/dace/codegen/targets/cuda.py +++ b/dace/codegen/targets/cuda.py @@ -1939,6 +1939,13 @@ def generate_kernel_scope(self, sdfg: SDFG, dfg_scope: ScopeSubgraphView, state_ kernel_params: list, function_stream: CodeIOStream, kernel_stream: CodeIOStream): node = dfg_scope.source_nodes()[0] + # Get the thread/block index type + ttype = Config.get('compiler', 'cuda', 'thread_id_type') + tidtype = getattr(dtypes, ttype, False) + if not isinstance(tidtype, dtypes.typeclass): + raise ValueError(f'Configured type "{ttype}" for ``thread_id_type`` does not match any DaCe data type. ' + 'See ``dace.dtypes`` for available types (for example ``int32``).') + # allocating shared memory for dynamic threadblock maps if has_dtbmap: kernel_stream.write( @@ -1990,8 +1997,8 @@ def generate_kernel_scope(self, sdfg: SDFG, dfg_scope: ScopeSubgraphView, state_ expr = _topy(bidx[i]).replace('__DAPB%d' % i, block_expr) - kernel_stream.write('int %s = %s;' % (varname, expr), sdfg, state_id, node) - self._dispatcher.defined_vars.add(varname, DefinedType.Scalar, 'int') + kernel_stream.write(f'{tidtype.ctype} {varname} = {expr};', sdfg, state_id, node) + self._dispatcher.defined_vars.add(varname, DefinedType.Scalar, tidtype.ctype) # Delinearize beyond the third dimension if len(krange) > 3: @@ -2010,8 +2017,8 @@ def generate_kernel_scope(self, sdfg: SDFG, dfg_scope: ScopeSubgraphView, state_ ) expr = _topy(bidx[i]).replace('__DAPB%d' % i, block_expr) - kernel_stream.write('int %s = %s;' % (varname, expr), sdfg, state_id, node) - self._dispatcher.defined_vars.add(varname, DefinedType.Scalar, 'int') + kernel_stream.write(f'{tidtype.ctype} {varname} = {expr};', sdfg, state_id, node) + self._dispatcher.defined_vars.add(varname, DefinedType.Scalar, tidtype.ctype) # Dispatch internal code assert CUDACodeGen._in_device_code is False diff --git a/dace/config_schema.yml b/dace/config_schema.yml index e378b6c1f2..08a427aa52 100644 --- a/dace/config_schema.yml +++ b/dace/config_schema.yml @@ -413,6 +413,17 @@ required: a specified larger block size in the third dimension. Default value is derived from hardware limits on common GPUs. + thread_id_type: + type: str + title: Thread/block index data type + default: int32 + description: > + Defines the data type for a thread and block index in the generated code. + The type is based on the type-classes in ``dace.dtypes``. For example, + ``uint64`` is equivalent to ``dace.uint64``. Change this setting when large + index types are needed to address memory offsets that are beyond the 32-bit + range, or to reduce memory usage. + ############################################# # General FPGA flags diff --git a/dace/frontend/fortran/ast_components.py b/dace/frontend/fortran/ast_components.py index a66ee5c0d6..1e5bfb4528 100644 --- a/dace/frontend/fortran/ast_components.py +++ b/dace/frontend/fortran/ast_components.py @@ -1,5 +1,6 @@ # Copyright 2019-2023 ETH Zurich and the DaCe authors. All rights reserved. from fparser.two.Fortran2008 import Fortran2008 as f08 +from fparser.two import Fortran2008 from fparser.two import Fortran2003 as f03 from fparser.two import symbol_table @@ -523,6 +524,31 @@ def declaration_type_spec(self, node: FASTNode): def assumed_shape_spec_list(self, node: FASTNode): return node + def parse_shape_specification(self, dim: f03.Explicit_Shape_Spec, size: List[FASTNode], offset: List[int]): + + dim_expr = [i for i in dim.children if i is not None] + + # handle size definition + if len(dim_expr) == 1: + dim_expr = dim_expr[0] + #now to add the dimension to the size list after processing it if necessary + size.append(self.create_ast(dim_expr)) + offset.append(1) + # Here we support arrays that have size declaration - with initial offset. + elif len(dim_expr) == 2: + # extract offets + for expr in dim_expr: + if not isinstance(expr, f03.Int_Literal_Constant): + raise TypeError("Array offsets must be constant expressions!") + offset.append(int(dim_expr[0].tostr())) + + fortran_size = int(dim_expr[1].tostr()) - int(dim_expr[0].tostr()) + 1 + fortran_ast_size = f03.Int_Literal_Constant(str(fortran_size)) + + size.append(self.create_ast(fortran_ast_size)) + else: + raise TypeError("Array dimension must be at most two expressions") + def type_declaration_stmt(self, node: FASTNode): #decide if its a intrinsic variable type or a derived type @@ -574,33 +600,44 @@ def type_declaration_stmt(self, node: FASTNode): alloc = False symbol = False + attr_size = None + attr_offset = None for i in attributes: if i.string.lower() == "allocatable": alloc = True if i.string.lower() == "parameter": symbol = True + if isinstance(i, Fortran2008.Attr_Spec_List): + + dimension_spec = get_children(i, "Dimension_Attr_Spec") + if len(dimension_spec) == 0: + continue + + attr_size = [] + attr_offset = [] + sizes = get_child(dimension_spec[0], ["Explicit_Shape_Spec_List"]) + + for shape_spec in get_children(sizes, [f03.Explicit_Shape_Spec]): + self.parse_shape_specification(shape_spec, attr_size, attr_offset) + vardecls = [] for var in names: #first handle dimensions size = None + offset = None var_components = self.create_children(var) array_sizes = get_children(var, "Explicit_Shape_Spec_List") actual_name = get_child(var_components, ast_internal_classes.Name_Node) if len(array_sizes) == 1: array_sizes = array_sizes[0] size = [] + offset = [] for dim in array_sizes.children: #sanity check if isinstance(dim, f03.Explicit_Shape_Spec): - dim_expr = [i for i in dim.children if i is not None] - if len(dim_expr) == 1: - dim_expr = dim_expr[0] - #now to add the dimension to the size list after processing it if necessary - size.append(self.create_ast(dim_expr)) - else: - raise TypeError("Array dimension must be a single expression") + self.parse_shape_specification(dim, size, offset) #handle initializiation init = None @@ -615,15 +652,26 @@ def type_declaration_stmt(self, node: FASTNode): if symbol == False: - vardecls.append( - ast_internal_classes.Var_Decl_Node(name=actual_name.name, - type=testtype, - alloc=alloc, - sizes=size, - kind=kind, - line_number=node.item.span)) + if attr_size is None: + vardecls.append( + ast_internal_classes.Var_Decl_Node(name=actual_name.name, + type=testtype, + alloc=alloc, + sizes=size, + offsets=offset, + kind=kind, + line_number=node.item.span)) + else: + vardecls.append( + ast_internal_classes.Var_Decl_Node(name=actual_name.name, + type=testtype, + alloc=alloc, + sizes=attr_size, + offsets=attr_offset, + kind=kind, + line_number=node.item.span)) else: - if size is None: + if size is None and attr_size is None: self.symbols[actual_name.name] = init vardecls.append( ast_internal_classes.Symbol_Decl_Node(name=actual_name.name, @@ -631,16 +679,26 @@ def type_declaration_stmt(self, node: FASTNode): alloc=alloc, init=init, line_number=node.item.span)) + elif attr_size is not None: + vardecls.append( + ast_internal_classes.Symbol_Array_Decl_Node(name=actual_name.name, + type=testtype, + alloc=alloc, + sizes=attr_size, + offsets=attr_offset, + kind=kind, + init=init, + line_number=node.item.span)) else: vardecls.append( ast_internal_classes.Symbol_Array_Decl_Node(name=actual_name.name, type=testtype, alloc=alloc, sizes=size, + offsets=offset, kind=kind, init=init, line_number=node.item.span)) - return ast_internal_classes.Decl_Stmt_Node(vardecl=vardecls, line_number=node.item.span) def entity_decl(self, node: FASTNode): diff --git a/dace/frontend/fortran/ast_internal_classes.py b/dace/frontend/fortran/ast_internal_classes.py index 6bdfb61faf..70a43e21b8 100644 --- a/dace/frontend/fortran/ast_internal_classes.py +++ b/dace/frontend/fortran/ast_internal_classes.py @@ -1,5 +1,5 @@ # Copyright 2019-2023 ETH Zurich and the DaCe authors. All rights reserved. -from typing import Any, List, Tuple, Type, TypeVar, Union, overload +from typing import Any, List, Optional, Tuple, Type, TypeVar, Union, overload # The node class is the base class for all nodes in the AST. It provides attributes including the line number and fields. # Attributes are not used when walking the tree, but are useful for debugging and for code generation. @@ -11,6 +11,14 @@ def __init__(self, *args, **kwargs): # real signature unknown self.integrity_exceptions = [] self.read_vars = [] self.written_vars = [] + self.parent: Optional[ + Union[ + Subroutine_Subprogram_Node, + Function_Subprogram_Node, + Main_Program_Node, + Module_Node + ] + ] = None for k, v in kwargs.items(): setattr(self, k, v) @@ -199,6 +207,7 @@ class Symbol_Array_Decl_Node(Statement_Node): ) _fields = ( 'sizes', + 'offsets' 'typeref', 'init', ) @@ -213,6 +222,7 @@ class Var_Decl_Node(Statement_Node): ) _fields = ( 'sizes', + 'offsets', 'typeref', 'init', ) diff --git a/dace/frontend/fortran/ast_transforms.py b/dace/frontend/fortran/ast_transforms.py index 7e5cd3bf00..e2a7246aed 100644 --- a/dace/frontend/fortran/ast_transforms.py +++ b/dace/frontend/fortran/ast_transforms.py @@ -1,7 +1,7 @@ # Copyright 2023 ETH Zurich and the DaCe authors. All rights reserved. from dace.frontend.fortran import ast_components, ast_internal_classes -from typing import List, Tuple, Set +from typing import Dict, List, Optional, Tuple, Set import copy @@ -310,6 +310,65 @@ def visit_Execution_Part_Node(self, node: ast_internal_classes.Execution_Part_No return ast_internal_classes.Execution_Part_Node(execution=newbody) +class ParentScopeAssigner(NodeVisitor): + """ + For each node, it assigns its parent scope - program, subroutine, function. + + If the parent node is one of the "parent" types, we assign it as the parent. + Otherwise, we look for the parent of my parent to cover nested AST nodes within + a single scope. + """ + def __init__(self): + pass + + def visit(self, node: ast_internal_classes.FNode, parent_node: Optional[ast_internal_classes.FNode] = None): + + parent_node_types = [ + ast_internal_classes.Subroutine_Subprogram_Node, + ast_internal_classes.Function_Subprogram_Node, + ast_internal_classes.Main_Program_Node, + ast_internal_classes.Module_Node + ] + + if parent_node is not None and type(parent_node) in parent_node_types: + node.parent = parent_node + elif parent_node is not None: + node.parent = parent_node.parent + + # Copied from `generic_visit` to recursively parse all leafs + for field, value in iter_fields(node): + if isinstance(value, list): + for item in value: + if isinstance(item, ast_internal_classes.FNode): + self.visit(item, node) + elif isinstance(value, ast_internal_classes.FNode): + self.visit(value, node) + +class ScopeVarsDeclarations(NodeVisitor): + """ + Creates a mapping (scope name, variable name) -> variable declaration. + + The visitor is used to access information on variable dimension, sizes, and offsets. + """ + + def __init__(self): + + self.scope_vars: Dict[Tuple[str, str], ast_internal_classes.FNode] = {} + + def get_var(self, scope: ast_internal_classes.FNode, variable_name: str) -> ast_internal_classes.FNode: + return self.scope_vars[(self._scope_name(scope), variable_name)] + + def visit_Var_Decl_Node(self, node: ast_internal_classes.Var_Decl_Node): + + parent_name = self._scope_name(node.parent) + var_name = node.name + self.scope_vars[(parent_name, var_name)] = node + + def _scope_name(self, scope: ast_internal_classes.FNode) -> str: + if isinstance(scope, ast_internal_classes.Main_Program_Node): + return scope.name.name.name + else: + return scope.name.name class IndexExtractorNodeLister(NodeVisitor): """ @@ -336,9 +395,20 @@ class IndexExtractor(NodeTransformer): Uses the IndexExtractorNodeLister to find all array subscript expressions in the AST node and its children that have to be extracted into independent expressions It then creates a new temporary variable for each of them and replaces the index expression with the variable. + + Before parsing the AST, the transformation first runs: + - ParentScopeAssigner to ensure that each node knows its scope assigner. + - ScopeVarsDeclarations to aggregate all variable declarations for each function. """ - def __init__(self, count=0): + def __init__(self, ast: ast_internal_classes.FNode, normalize_offsets: bool = False, count=0): + self.count = count + self.normalize_offsets = normalize_offsets + + if normalize_offsets: + ParentScopeAssigner().visit(ast) + self.scope_vars = ScopeVarsDeclarations() + self.scope_vars.visit(ast) def visit_Call_Expr_Node(self, node: ast_internal_classes.Call_Expr_Node): if node.name.name in ["sqrt", "exp", "pow", "max", "min", "abs", "tanh"]: @@ -367,9 +437,11 @@ def visit_Execution_Part_Node(self, node: ast_internal_classes.Execution_Part_No lister.visit(child) res = lister.nodes temp = self.count + + if res is not None: for j in res: - for i in j.indices: + for idx, i in enumerate(j.indices): if isinstance(i, ast_internal_classes.ParDecl_Node): continue else: @@ -383,16 +455,34 @@ def visit_Execution_Part_Node(self, node: ast_internal_classes.Execution_Part_No line_number=child.line_number) ], line_number=child.line_number)) - newbody.append( - ast_internal_classes.BinOp_Node( - op="=", - lval=ast_internal_classes.Name_Node(name=tmp_name), - rval=ast_internal_classes.BinOp_Node( - op="-", - lval=i, - rval=ast_internal_classes.Int_Literal_Node(value="1"), - line_number=child.line_number), - line_number=child.line_number)) + if self.normalize_offsets: + + # Find the offset of a variable to which we are assigning + var_name = child.lval.name.name + variable = self.scope_vars.get_var(child.parent, var_name) + offset = variable.offsets[idx] + + newbody.append( + ast_internal_classes.BinOp_Node( + op="=", + lval=ast_internal_classes.Name_Node(name=tmp_name), + rval=ast_internal_classes.BinOp_Node( + op="-", + lval=i, + rval=ast_internal_classes.Int_Literal_Node(value=str(offset)), + line_number=child.line_number), + line_number=child.line_number)) + else: + newbody.append( + ast_internal_classes.BinOp_Node( + op="=", + lval=ast_internal_classes.Name_Node(name=tmp_name), + rval=ast_internal_classes.BinOp_Node( + op="-", + lval=i, + rval=ast_internal_classes.Int_Literal_Node(value="1"), + line_number=child.line_number), + line_number=child.line_number)) newbody.append(self.visit(child)) return ast_internal_classes.Execution_Part_Node(execution=newbody) @@ -646,6 +736,7 @@ def par_Decl_Range_Finder(node: ast_internal_classes.Array_Subscript_Node, rangepos: list, count: int, newbody: list, + scope_vars: ScopeVarsDeclarations, declaration=True, is_sum_to_loop=False): """ @@ -662,16 +753,40 @@ def par_Decl_Range_Finder(node: ast_internal_classes.Array_Subscript_Node, currentindex = 0 indices = [] - for i in node.indices: + offsets = scope_vars.get_var(node.parent, node.name.name).offsets + + for idx, i in enumerate(node.indices): if isinstance(i, ast_internal_classes.ParDecl_Node): + if i.type == "ALL": - ranges.append([ - ast_internal_classes.Int_Literal_Node(value="1"), - ast_internal_classes.Name_Range_Node(name="f2dace_MAX", - type="INTEGER", - arrname=node.name, - pos=currentindex) - ]) + + lower_boundary = None + if offsets[idx] != 1: + lower_boundary = ast_internal_classes.Int_Literal_Node(value=str(offsets[idx])) + else: + lower_boundary = ast_internal_classes.Int_Literal_Node(value="1") + + upper_boundary = ast_internal_classes.Name_Range_Node(name="f2dace_MAX", + type="INTEGER", + arrname=node.name, + pos=currentindex) + """ + When there's an offset, we add MAX_RANGE + offset. + But since the generated loop has `<=` condition, we need to subtract 1. + """ + if offsets[idx] != 1: + upper_boundary = ast_internal_classes.BinOp_Node( + lval=upper_boundary, + op="+", + rval=ast_internal_classes.Int_Literal_Node(value=str(offsets[idx])) + ) + upper_boundary = ast_internal_classes.BinOp_Node( + lval=upper_boundary, + op="-", + rval=ast_internal_classes.Int_Literal_Node(value="1") + ) + ranges.append([lower_boundary, upper_boundary]) + else: ranges.append([i.range[0], i.range[1]]) rangepos.append(currentindex) @@ -693,9 +808,13 @@ class ArrayToLoop(NodeTransformer): """ Transforms the AST by removing array expressions and replacing them with loops """ - def __init__(self): + def __init__(self, ast): self.count = 0 + ParentScopeAssigner().visit(ast) + self.scope_vars = ScopeVarsDeclarations() + self.scope_vars.visit(ast) + def visit_Execution_Part_Node(self, node: ast_internal_classes.Execution_Part_Node): newbody = [] for child in node.execution: @@ -709,7 +828,7 @@ def visit_Execution_Part_Node(self, node: ast_internal_classes.Execution_Part_No val = child.rval ranges = [] rangepos = [] - par_Decl_Range_Finder(current, ranges, rangepos, self.count, newbody, True) + par_Decl_Range_Finder(current, ranges, rangepos, self.count, newbody, self.scope_vars, True) if res_range is not None and len(res_range) > 0: rvals = [i for i in mywalk(val) if isinstance(i, ast_internal_classes.Array_Subscript_Node)] @@ -717,7 +836,7 @@ def visit_Execution_Part_Node(self, node: ast_internal_classes.Execution_Part_No rangeposrval = [] rangesrval = [] - par_Decl_Range_Finder(i, rangesrval, rangeposrval, self.count, newbody, False) + par_Decl_Range_Finder(i, rangesrval, rangeposrval, self.count, newbody, self.scope_vars, False) for i, j in zip(ranges, rangesrval): if i != j: @@ -791,8 +910,11 @@ class SumToLoop(NodeTransformer): """ Transforms the AST by removing array sums and replacing them with loops """ - def __init__(self): + def __init__(self, ast): self.count = 0 + ParentScopeAssigner().visit(ast) + self.scope_vars = ScopeVarsDeclarations() + self.scope_vars.visit(ast) def visit_Execution_Part_Node(self, node: ast_internal_classes.Execution_Part_Node): newbody = [] @@ -811,7 +933,7 @@ def visit_Execution_Part_Node(self, node: ast_internal_classes.Execution_Part_No rangeposrval = [] rangesrval = [] - par_Decl_Range_Finder(val, rangesrval, rangeposrval, self.count, newbody, False, True) + par_Decl_Range_Finder(val, rangesrval, rangeposrval, self.count, newbody, self.scope_vars, False, True) range_index = 0 body = ast_internal_classes.BinOp_Node(lval=current, diff --git a/dace/frontend/fortran/fortran_parser.py b/dace/frontend/fortran/fortran_parser.py index d7112892fe..b15435f4ff 100644 --- a/dace/frontend/fortran/fortran_parser.py +++ b/dace/frontend/fortran/fortran_parser.py @@ -133,7 +133,7 @@ def translate(self, node: ast_internal_classes.FNode, sdfg: SDFG): for i in node: self.translate(i, sdfg) else: - warnings.warn("WARNING:", node.__class__.__name__) + warnings.warn(f"WARNING: {node.__class__.__name__}") def ast2sdfg(self, node: ast_internal_classes.Program_Node, sdfg: SDFG): """ @@ -1015,10 +1015,46 @@ def vardecl2sdfg(self, node: ast_internal_classes.Var_Decl_Node, sdfg: SDFG): if node.name not in self.contexts[sdfg.name].containers: self.contexts[sdfg.name].containers.append(node.name) +def create_ast_from_string( + source_string: str, + sdfg_name: str, + transform: bool = False, + normalize_offsets: bool = False +): + """ + Creates an AST from a Fortran file in a string + :param source_string: The fortran file as a string + :param sdfg_name: The name to be given to the resulting SDFG + :return: The resulting AST + + """ + parser = pf().create(std="f2008") + reader = fsr(source_string) + ast = parser(reader) + tables = SymbolTable + own_ast = ast_components.InternalFortranAst(ast, tables) + program = own_ast.create_ast(ast) + + functions_and_subroutines_builder = ast_transforms.FindFunctionAndSubroutines() + functions_and_subroutines_builder.visit(program) + functions_and_subroutines = functions_and_subroutines_builder.nodes + + if transform: + program = ast_transforms.functionStatementEliminator(program) + program = ast_transforms.CallToArray(functions_and_subroutines_builder.nodes).visit(program) + program = ast_transforms.CallExtractor().visit(program) + program = ast_transforms.SignToIf().visit(program) + program = ast_transforms.ArrayToLoop(program).visit(program) + program = ast_transforms.SumToLoop(program).visit(program) + program = ast_transforms.ForDeclarer().visit(program) + program = ast_transforms.IndexExtractor(program, normalize_offsets).visit(program) + + return (program, own_ast) def create_sdfg_from_string( source_string: str, sdfg_name: str, + normalize_offsets: bool = False ): """ Creates an SDFG from a fortran file in a string @@ -1040,10 +1076,10 @@ def create_sdfg_from_string( program = ast_transforms.CallToArray(functions_and_subroutines_builder.nodes).visit(program) program = ast_transforms.CallExtractor().visit(program) program = ast_transforms.SignToIf().visit(program) - program = ast_transforms.ArrayToLoop().visit(program) - program = ast_transforms.SumToLoop().visit(program) + program = ast_transforms.ArrayToLoop(program).visit(program) + program = ast_transforms.SumToLoop(program).visit(program) program = ast_transforms.ForDeclarer().visit(program) - program = ast_transforms.IndexExtractor().visit(program) + program = ast_transforms.IndexExtractor(program, normalize_offsets).visit(program) ast2sdfg = AST_translator(own_ast, __file__) sdfg = SDFG(sdfg_name) ast2sdfg.top_level = program @@ -1082,10 +1118,10 @@ def create_sdfg_from_fortran_file(source_string: str): program = ast_transforms.CallToArray(functions_and_subroutines_builder.nodes).visit(program) program = ast_transforms.CallExtractor().visit(program) program = ast_transforms.SignToIf().visit(program) - program = ast_transforms.ArrayToLoop().visit(program) - program = ast_transforms.SumToLoop().visit(program) + program = ast_transforms.ArrayToLoop(program).visit(program) + program = ast_transforms.SumToLoop(program).visit(program) program = ast_transforms.ForDeclarer().visit(program) - program = ast_transforms.IndexExtractor().visit(program) + program = ast_transforms.IndexExtractor(program).visit(program) ast2sdfg = AST_translator(own_ast, __file__) sdfg = SDFG(source_string) ast2sdfg.top_level = program diff --git a/dace/libraries/blas/environments/cublas.py b/dace/libraries/blas/environments/cublas.py index d4ab879e61..ef73b511c0 100644 --- a/dace/libraries/blas/environments/cublas.py +++ b/dace/libraries/blas/environments/cublas.py @@ -25,7 +25,7 @@ class cuBLAS: def handle_setup_code(node): location = node.location if not location or "gpu" not in node.location: - location = 0 + location = -1 # -1 means current device else: try: location = int(location["gpu"]) diff --git a/dace/libraries/blas/environments/rocblas.py b/dace/libraries/blas/environments/rocblas.py index 5d752ed690..47e16531ff 100644 --- a/dace/libraries/blas/environments/rocblas.py +++ b/dace/libraries/blas/environments/rocblas.py @@ -25,7 +25,7 @@ class rocBLAS: def handle_setup_code(node): location = node.location if not location or "gpu" not in node.location: - location = 0 + location = -1 # -1 means current device else: try: location = int(location["gpu"]) diff --git a/dace/libraries/blas/include/dace_cublas.h b/dace/libraries/blas/include/dace_cublas.h index 8ec03c2b37..3547a009d2 100644 --- a/dace/libraries/blas/include/dace_cublas.h +++ b/dace/libraries/blas/include/dace_cublas.h @@ -21,8 +21,10 @@ static void CheckCublasError(cublasStatus_t const& status) { } static cublasHandle_t CreateCublasHandle(int device) { - if (cudaSetDevice(device) != cudaSuccess) { - throw std::runtime_error("Failed to set CUDA device."); + if (device >= 0) { + if (cudaSetDevice(device) != cudaSuccess) { + throw std::runtime_error("Failed to set CUDA device."); + } } cublasHandle_t handle; CheckCublasError(cublasCreate(&handle)); @@ -65,8 +67,10 @@ class _CublasConstants { } _CublasConstants(int device) { - if (cudaSetDevice(device) != cudaSuccess) { - throw std::runtime_error("Failed to set CUDA device."); + if (device >= 0) { + if (cudaSetDevice(device) != cudaSuccess) { + throw std::runtime_error("Failed to set CUDA device."); + } } // Allocate constant zero with the largest used size cudaMalloc(&zero_, sizeof(cuDoubleComplex) * 1); diff --git a/dace/libraries/blas/include/dace_rocblas.h b/dace/libraries/blas/include/dace_rocblas.h index 7a7e4a75ee..00469136a3 100644 --- a/dace/libraries/blas/include/dace_rocblas.h +++ b/dace/libraries/blas/include/dace_rocblas.h @@ -24,8 +24,10 @@ static void CheckRocblasError(rocblas_status const& status) { } static rocblas_handle CreateRocblasHandle(int device) { - if (hipSetDevice(device) != hipSuccess) { - throw std::runtime_error("Failed to set HIP device."); + if (device >= 0) { + if (hipSetDevice(device) != hipSuccess) { + throw std::runtime_error("Failed to set HIP device."); + } } rocblas_handle handle; CheckRocblasError(rocblas_create_handle(&handle)); @@ -68,53 +70,55 @@ class _RocblasConstants { } _RocblasConstants(int device) { - if (hipSetDevice(device) != hipSuccess) { - throw std::runtime_error("Failed to set HIP device."); + if (device >= 0) { + if (hipSetDevice(device) != hipSuccess) { + throw std::runtime_error("Failed to set HIP device."); + } } // Allocate constant zero with the largest used size - hipMalloc(&zero_, sizeof(hipDoubleComplex) * 1); - hipMemset(zero_, 0, sizeof(hipDoubleComplex) * 1); + (void)hipMalloc(&zero_, sizeof(hipDoubleComplex) * 1); + (void)hipMemset(zero_, 0, sizeof(hipDoubleComplex) * 1); // Allocate constant one - hipMalloc(&half_pone_, sizeof(__half) * 1); + (void)hipMalloc(&half_pone_, sizeof(__half) * 1); __half half_pone = __float2half(1.0f); - hipMemcpy(half_pone_, &half_pone, sizeof(__half) * 1, + (void)hipMemcpy(half_pone_, &half_pone, sizeof(__half) * 1, hipMemcpyHostToDevice); - hipMalloc(&float_pone_, sizeof(float) * 1); + (void)hipMalloc(&float_pone_, sizeof(float) * 1); float float_pone = 1.0f; - hipMemcpy(float_pone_, &float_pone, sizeof(float) * 1, + (void)hipMemcpy(float_pone_, &float_pone, sizeof(float) * 1, hipMemcpyHostToDevice); - hipMalloc(&double_pone_, sizeof(double) * 1); + (void)hipMalloc(&double_pone_, sizeof(double) * 1); double double_pone = 1.0; - hipMemcpy(double_pone_, &double_pone, sizeof(double) * 1, + (void)hipMemcpy(double_pone_, &double_pone, sizeof(double) * 1, hipMemcpyHostToDevice); - hipMalloc(&complex64_pone_, sizeof(hipComplex) * 1); + (void)hipMalloc(&complex64_pone_, sizeof(hipComplex) * 1); hipComplex complex64_pone = make_hipFloatComplex(1.0f, 0.0f); - hipMemcpy(complex64_pone_, &complex64_pone, sizeof(hipComplex) * 1, + (void)hipMemcpy(complex64_pone_, &complex64_pone, sizeof(hipComplex) * 1, hipMemcpyHostToDevice); - hipMalloc(&complex128_pone_, sizeof(hipDoubleComplex) * 1); + (void)hipMalloc(&complex128_pone_, sizeof(hipDoubleComplex) * 1); hipDoubleComplex complex128_pone = make_hipDoubleComplex(1.0, 0.0); - hipMemcpy(complex128_pone_, &complex128_pone, sizeof(hipDoubleComplex) * 1, + (void)hipMemcpy(complex128_pone_, &complex128_pone, sizeof(hipDoubleComplex) * 1, hipMemcpyHostToDevice); // Allocate custom factors and default to zero - hipMalloc(&custom_alpha_, sizeof(hipDoubleComplex) * 1); - hipMemset(custom_alpha_, 0, sizeof(hipDoubleComplex) * 1); - hipMalloc(&custom_beta_, sizeof(hipDoubleComplex) * 1); - hipMemset(custom_beta_, 0, sizeof(hipDoubleComplex) * 1); + (void)hipMalloc(&custom_alpha_, sizeof(hipDoubleComplex) * 1); + (void)hipMemset(custom_alpha_, 0, sizeof(hipDoubleComplex) * 1); + (void)hipMalloc(&custom_beta_, sizeof(hipDoubleComplex) * 1); + (void)hipMemset(custom_beta_, 0, sizeof(hipDoubleComplex) * 1); } _RocblasConstants(_RocblasConstants const&) = delete; ~_RocblasConstants() { - hipFree(zero_); - hipFree(half_pone_); - hipFree(float_pone_); - hipFree(double_pone_); - hipFree(complex64_pone_); - hipFree(complex128_pone_); - hipFree(custom_alpha_); - hipFree(custom_beta_); + (void)hipFree(zero_); + (void)hipFree(half_pone_); + (void)hipFree(float_pone_); + (void)hipFree(double_pone_); + (void)hipFree(complex64_pone_); + (void)hipFree(complex128_pone_); + (void)hipFree(custom_alpha_); + (void)hipFree(custom_beta_); } _RocblasConstants& operator=(_RocblasConstants const&) = delete; diff --git a/dace/libraries/blas/nodes/gemm.py b/dace/libraries/blas/nodes/gemm.py index 2db2055ae5..83be99d78b 100644 --- a/dace/libraries/blas/nodes/gemm.py +++ b/dace/libraries/blas/nodes/gemm.py @@ -184,11 +184,11 @@ def expansion(node, state, sdfg): code = '' if dtype in (dace.complex64, dace.complex128): code = f''' - {dtype.ctype} alpha = {alpha}; - {dtype.ctype} beta = {beta}; + {dtype.ctype} __alpha = {alpha}; + {dtype.ctype} __beta = {beta}; ''' - opt['alpha'] = '&alpha' - opt['beta'] = '&beta' + opt['alpha'] = '&__alpha' + opt['beta'] = '&__beta' code += ("cblas_{func}(CblasColMajor, {ta}, {tb}, " "{M}, {N}, {K}, {alpha}, {x}, {lda}, {y}, {ldb}, {beta}, " @@ -287,12 +287,12 @@ def expansion(cls, node, state, sdfg): # Set pointer mode to host call_prefix += f'''{cls.set_pointer_mode}(__dace_{cls.backend}blas_handle, {cls.pointer_host}); - {dtype.ctype} alpha = {alpha}; - {dtype.ctype} beta = {beta}; + {dtype.ctype} __alpha = {alpha}; + {dtype.ctype} __beta = {beta}; ''' call_suffix += f'''{cls.set_pointer_mode}(__dace_{cls.backend}blas_handle, {cls.pointer_device});''' - alpha = f'({cdtype} *)&alpha' - beta = f'({cdtype} *)&beta' + alpha = f'({cdtype} *)&__alpha' + beta = f'({cdtype} *)&__beta' else: alpha = constants[node.alpha] beta = constants[node.beta] diff --git a/dace/libraries/lapack/environments/cusolverdn.py b/dace/libraries/lapack/environments/cusolverdn.py index c92c8bf3e7..4daad8062e 100644 --- a/dace/libraries/lapack/environments/cusolverdn.py +++ b/dace/libraries/lapack/environments/cusolverdn.py @@ -24,7 +24,7 @@ class cuSolverDn: def handle_setup_code(node): location = node.location if not location or "gpu" not in node.location: - location = 0 + location = -1 # -1 means current device else: try: location = int(location["gpu"]) diff --git a/dace/libraries/lapack/include/dace_cusolverdn.h b/dace/libraries/lapack/include/dace_cusolverdn.h index 2da65ffa2f..f262541f0b 100644 --- a/dace/libraries/lapack/include/dace_cusolverdn.h +++ b/dace/libraries/lapack/include/dace_cusolverdn.h @@ -21,8 +21,10 @@ static void CheckCusolverDnError(cusolverStatus_t const& status) { } static cusolverDnHandle_t CreateCusolverDnHandle(int device) { - if (cudaSetDevice(device) != cudaSuccess) { - throw std::runtime_error("Failed to set CUDA device."); + if (device >= 0) { + if (cudaSetDevice(device) != cudaSuccess) { + throw std::runtime_error("Failed to set CUDA device."); + } } cusolverDnHandle_t handle; CheckCusolverDnError(cusolverDnCreate(&handle)); diff --git a/dace/libraries/linalg/environments/cutensor.py b/dace/libraries/linalg/environments/cutensor.py index e3572a0673..0022ec1f57 100644 --- a/dace/libraries/linalg/environments/cutensor.py +++ b/dace/libraries/linalg/environments/cutensor.py @@ -24,7 +24,7 @@ class cuTensor: def handle_setup_code(node): location = node.location if not location or "gpu" not in node.location: - location = 0 + location = -1 # -1 means current device else: try: location = int(location["gpu"]) diff --git a/dace/libraries/linalg/include/dace_cutensor.h b/dace/libraries/linalg/include/dace_cutensor.h index 8079892285..ddad2feaa3 100644 --- a/dace/libraries/linalg/include/dace_cutensor.h +++ b/dace/libraries/linalg/include/dace_cutensor.h @@ -20,8 +20,10 @@ static void CheckCuTensorError(cutensorStatus_t const& status) { } static cutensorHandle_t CreateCuTensorHandle(int device) { - if (cudaSetDevice(device) != cudaSuccess) { - throw std::runtime_error("Failed to set CUDA device."); + if (device >= 0) { + if (cudaSetDevice(device) != cudaSuccess) { + throw std::runtime_error("Failed to set CUDA device."); + } } cutensorHandle_t handle; CheckCuTensorError(cutensorInit(&handle)); diff --git a/dace/libraries/sparse/environments/cusparse.py b/dace/libraries/sparse/environments/cusparse.py index 0970557944..a731f75bf7 100644 --- a/dace/libraries/sparse/environments/cusparse.py +++ b/dace/libraries/sparse/environments/cusparse.py @@ -24,7 +24,7 @@ class cuSPARSE: def handle_setup_code(node): location = node.location if not location or "gpu" not in node.location: - location = 0 + location = -1 # -1 means current device else: try: location = int(location["gpu"]) diff --git a/dace/libraries/sparse/include/dace_cusparse.h b/dace/libraries/sparse/include/dace_cusparse.h index 82470089e0..9d28bb4748 100644 --- a/dace/libraries/sparse/include/dace_cusparse.h +++ b/dace/libraries/sparse/include/dace_cusparse.h @@ -20,8 +20,10 @@ static void CheckCusparseError(cusparseStatus_t const& status) { } static cusparseHandle_t CreateCusparseHandle(int device) { - if (cudaSetDevice(device) != cudaSuccess) { - throw std::runtime_error("Failed to set CUDA device."); + if (device >= 0) { + if (cudaSetDevice(device) != cudaSuccess) { + throw std::runtime_error("Failed to set CUDA device."); + } } cusparseHandle_t handle; CheckCusparseError(cusparseCreate(&handle)); diff --git a/tests/fortran/array_attributes_test.py b/tests/fortran/array_attributes_test.py new file mode 100644 index 0000000000..af433905bc --- /dev/null +++ b/tests/fortran/array_attributes_test.py @@ -0,0 +1,117 @@ +# Copyright 2019-2023 ETH Zurich and the DaCe authors. All rights reserved. + +import numpy as np + +from dace.frontend.fortran import fortran_parser + +def test_fortran_frontend_array_attribute_no_offset(): + """ + Tests that the Fortran frontend can parse array accesses and that the accessed indices are correct. + """ + test_string = """ + PROGRAM index_offset_test + implicit none + double precision, dimension(5) :: d + CALL index_test_function(d) + end + + SUBROUTINE index_test_function(d) + double precision, dimension(5) :: d + + do i=1,5 + d(i) = i * 2.0 + end do + + END SUBROUTINE index_test_function + """ + sdfg = fortran_parser.create_sdfg_from_string(test_string, "index_offset_test") + sdfg.simplify(verbose=True) + sdfg.compile() + + assert len(sdfg.data('d').shape) == 1 + assert sdfg.data('d').shape[0] == 5 + assert len(sdfg.data('d').offset) == 1 + assert sdfg.data('d').offset[0] == -1 + + a = np.full([5], 42, order="F", dtype=np.float64) + sdfg(d=a) + for i in range(1,5): + # offset -1 is already added + assert a[i-1] == i * 2 + +def test_fortran_frontend_array_attribute_offset(): + """ + Tests that the Fortran frontend can parse array accesses and that the accessed indices are correct. + """ + test_string = """ + PROGRAM index_offset_test + implicit none + double precision, dimension(50:54) :: d + CALL index_test_function(d) + end + + SUBROUTINE index_test_function(d) + double precision, dimension(50:54) :: d + + do i=50,54 + d(i) = i * 2.0 + end do + + END SUBROUTINE index_test_function + """ + sdfg = fortran_parser.create_sdfg_from_string(test_string, "index_offset_test") + sdfg.simplify(verbose=True) + sdfg.compile() + + assert len(sdfg.data('d').shape) == 1 + assert sdfg.data('d').shape[0] == 5 + assert len(sdfg.data('d').offset) == 1 + assert sdfg.data('d').offset[0] == -1 + + a = np.full([60], 42, order="F", dtype=np.float64) + sdfg(d=a) + for i in range(50,54): + # offset -1 is already added + assert a[i-1] == i * 2 + +def test_fortran_frontend_array_offset(): + """ + Tests that the Fortran frontend can parse array accesses and that the accessed indices are correct. + """ + test_string = """ + PROGRAM index_offset_test + implicit none + double precision d(50:54) + CALL index_test_function(d) + end + + SUBROUTINE index_test_function(d) + double precision d(50:54) + + do i=50,54 + d(i) = i * 2.0 + end do + + END SUBROUTINE index_test_function + """ + sdfg = fortran_parser.create_sdfg_from_string(test_string, "index_offset_test") + sdfg.simplify(verbose=True) + sdfg.compile() + + assert len(sdfg.data('d').shape) == 1 + assert sdfg.data('d').shape[0] == 5 + assert len(sdfg.data('d').offset) == 1 + assert sdfg.data('d').offset[0] == -1 + + a = np.full([60], 42, order="F", dtype=np.float64) + sdfg(d=a) + for i in range(50,54): + # offset -1 is already added + assert a[i-1] == i * 2 + + +if __name__ == "__main__": + + test_fortran_frontend_array_offset() + test_fortran_frontend_array_attribute_no_offset() + test_fortran_frontend_array_attribute_offset() diff --git a/tests/fortran/array_to_loop_offset.py b/tests/fortran/array_to_loop_offset.py new file mode 100644 index 0000000000..43d01d9b6b --- /dev/null +++ b/tests/fortran/array_to_loop_offset.py @@ -0,0 +1,119 @@ +# Copyright 2019-2023 ETH Zurich and the DaCe authors. All rights reserved. + +import numpy as np + +from dace.frontend.fortran import ast_transforms, fortran_parser + +def test_fortran_frontend_arr2loop_without_offset(): + """ + Tests that the generated array map correctly handles offsets. + """ + test_string = """ + PROGRAM index_offset_test + implicit none + double precision, dimension(5,3) :: d + CALL index_test_function(d) + end + + SUBROUTINE index_test_function(d) + double precision, dimension(5,3) :: d + + do i=1,5 + d(i, :) = i * 2.0 + end do + + END SUBROUTINE index_test_function + """ + + # Now test to verify it executes correctly with no offset normalization + + sdfg = fortran_parser.create_sdfg_from_string(test_string, "index_offset_test", False) + sdfg.simplify(verbose=True) + sdfg.compile() + + assert len(sdfg.data('d').shape) == 2 + assert sdfg.data('d').shape[0] == 5 + assert sdfg.data('d').shape[1] == 3 + + a = np.full([5,9], 42, order="F", dtype=np.float64) + sdfg(d=a) + for i in range(1,6): + for j in range(1,4): + assert a[i-1, j-1] == i * 2 + +def test_fortran_frontend_arr2loop_1d_offset(): + """ + Tests that the generated array map correctly handles offsets. + """ + test_string = """ + PROGRAM index_offset_test + implicit none + double precision, dimension(2:6) :: d + CALL index_test_function(d) + end + + SUBROUTINE index_test_function(d) + double precision, dimension(2:6) :: d + + d(:) = 5 + + END SUBROUTINE index_test_function + """ + + # Now test to verify it executes correctly with no offset normalization + + sdfg = fortran_parser.create_sdfg_from_string(test_string, "index_offset_test", False) + sdfg.simplify(verbose=True) + sdfg.compile() + + assert len(sdfg.data('d').shape) == 1 + assert sdfg.data('d').shape[0] == 5 + + a = np.full([6], 42, order="F", dtype=np.float64) + sdfg(d=a) + assert a[0] == 42 + for i in range(2,7): + assert a[i-1] == 5 + +def test_fortran_frontend_arr2loop_2d_offset(): + """ + Tests that the generated array map correctly handles offsets. + """ + test_string = """ + PROGRAM index_offset_test + implicit none + double precision, dimension(5,7:9) :: d + CALL index_test_function(d) + end + + SUBROUTINE index_test_function(d) + double precision, dimension(5,7:9) :: d + + do i=1,5 + d(i, :) = i * 2.0 + end do + + END SUBROUTINE index_test_function + """ + + # Now test to verify it executes correctly with no offset normalization + + sdfg = fortran_parser.create_sdfg_from_string(test_string, "index_offset_test", False) + sdfg.simplify(verbose=True) + sdfg.compile() + + assert len(sdfg.data('d').shape) == 2 + assert sdfg.data('d').shape[0] == 5 + assert sdfg.data('d').shape[1] == 3 + + a = np.full([5,9], 42, order="F", dtype=np.float64) + sdfg(d=a) + for i in range(1,6): + for j in range(7,10): + assert a[i-1, j-1] == i * 2 + +if __name__ == "__main__": + + test_fortran_frontend_arr2loop_1d_offset() + test_fortran_frontend_arr2loop_2d_offset() + test_fortran_frontend_arr2loop_without_offset() diff --git a/tests/fortran/offset_normalizer.py b/tests/fortran/offset_normalizer.py new file mode 100644 index 0000000000..b4138c1cac --- /dev/null +++ b/tests/fortran/offset_normalizer.py @@ -0,0 +1,164 @@ +# Copyright 2019-2023 ETH Zurich and the DaCe authors. All rights reserved. + +import numpy as np + +from dace.frontend.fortran import ast_transforms, fortran_parser + +def test_fortran_frontend_offset_normalizer_1d(): + """ + Tests that the Fortran frontend can parse array accesses and that the accessed indices are correct. + """ + test_string = """ + PROGRAM index_offset_test + implicit none + double precision, dimension(50:54) :: d + CALL index_test_function(d) + end + + SUBROUTINE index_test_function(d) + double precision, dimension(50:54) :: d + + do i=50,54 + d(i) = i * 2.0 + end do + + END SUBROUTINE index_test_function + """ + + # Test to verify that offset is normalized correctly + ast, own_ast = fortran_parser.create_ast_from_string(test_string, "index_offset_test", True, True) + + for subroutine in ast.subroutine_definitions: + + loop = subroutine.execution_part.execution[1] + idx_assignment = loop.body.execution[1] + assert idx_assignment.rval.rval.value == "50" + + # Now test to verify it executes correctly + + sdfg = fortran_parser.create_sdfg_from_string(test_string, "index_offset_test", True) + sdfg.simplify(verbose=True) + sdfg.compile() + + assert len(sdfg.data('d').shape) == 1 + assert sdfg.data('d').shape[0] == 5 + + a = np.full([5], 42, order="F", dtype=np.float64) + sdfg(d=a) + for i in range(0,5): + assert a[i] == (50+i)* 2 + +def test_fortran_frontend_offset_normalizer_2d(): + """ + Tests that the Fortran frontend can parse array accesses and that the accessed indices are correct. + """ + test_string = """ + PROGRAM index_offset_test + implicit none + double precision, dimension(50:54,7:9) :: d + CALL index_test_function(d) + end + + SUBROUTINE index_test_function(d) + double precision, dimension(50:54,7:9) :: d + + do i=50,54 + do j=7,9 + d(i, j) = i * 2.0 + 3 * j + end do + end do + + END SUBROUTINE index_test_function + """ + + # Test to verify that offset is normalized correctly + ast, own_ast = fortran_parser.create_ast_from_string(test_string, "index_offset_test", True, True) + + for subroutine in ast.subroutine_definitions: + + loop = subroutine.execution_part.execution[1] + nested_loop = loop.body.execution[1] + + idx = nested_loop.body.execution[1] + assert idx.lval.name == 'tmp_index_0' + assert idx.rval.rval.value == "50" + + idx2 = nested_loop.body.execution[3] + assert idx2.lval.name == 'tmp_index_1' + assert idx2.rval.rval.value == "7" + + # Now test to verify it executes correctly + + sdfg = fortran_parser.create_sdfg_from_string(test_string, "index_offset_test", True) + sdfg.simplify(verbose=True) + sdfg.compile() + + assert len(sdfg.data('d').shape) == 2 + assert sdfg.data('d').shape[0] == 5 + assert sdfg.data('d').shape[1] == 3 + + a = np.full([5,3], 42, order="F", dtype=np.float64) + sdfg(d=a) + for i in range(0,5): + for j in range(0,3): + assert a[i, j] == (50+i) * 2 + 3 * (7 + j) + +def test_fortran_frontend_offset_normalizer_2d_arr2loop(): + """ + Tests that the Fortran frontend can parse array accesses and that the accessed indices are correct. + """ + test_string = """ + PROGRAM index_offset_test + implicit none + double precision, dimension(50:54,7:9) :: d + CALL index_test_function(d) + end + + SUBROUTINE index_test_function(d) + double precision, dimension(50:54,7:9) :: d + + do i=50,54 + d(i, :) = i * 2.0 + end do + + END SUBROUTINE index_test_function + """ + + # Test to verify that offset is normalized correctly + ast, own_ast = fortran_parser.create_ast_from_string(test_string, "index_offset_test", True, True) + + for subroutine in ast.subroutine_definitions: + + loop = subroutine.execution_part.execution[1] + nested_loop = loop.body.execution[1] + + idx = nested_loop.body.execution[1] + assert idx.lval.name == 'tmp_index_0' + assert idx.rval.rval.value == "50" + + idx2 = nested_loop.body.execution[3] + assert idx2.lval.name == 'tmp_index_1' + assert idx2.rval.rval.value == "7" + + # Now test to verify it executes correctly with no normalization + + sdfg = fortran_parser.create_sdfg_from_string(test_string, "index_offset_test", True) + sdfg.save('test.sdfg') + sdfg.simplify(verbose=True) + sdfg.compile() + + assert len(sdfg.data('d').shape) == 2 + assert sdfg.data('d').shape[0] == 5 + assert sdfg.data('d').shape[1] == 3 + + a = np.full([5,3], 42, order="F", dtype=np.float64) + sdfg(d=a) + for i in range(0,5): + for j in range(0,3): + assert a[i, j] == (50 + i) * 2 + +if __name__ == "__main__": + + test_fortran_frontend_offset_normalizer_1d() + test_fortran_frontend_offset_normalizer_2d() + test_fortran_frontend_offset_normalizer_2d_arr2loop() diff --git a/tests/fortran/parent_test.py b/tests/fortran/parent_test.py new file mode 100644 index 0000000000..b1d08eaf37 --- /dev/null +++ b/tests/fortran/parent_test.py @@ -0,0 +1,91 @@ +# Copyright 2023 ETH Zurich and the DaCe authors. All rights reserved. + +from dace.frontend.fortran import fortran_parser + +import dace.frontend.fortran.ast_transforms as ast_transforms +import dace.frontend.fortran.ast_internal_classes as ast_internal_classes + + +def test_fortran_frontend_parent(): + """ + Tests that the Fortran frontend can parse array accesses and that the accessed indices are correct. + """ + test_string = """ + PROGRAM access_test + implicit none + double precision d(4) + d(1)=0 + CALL array_access_test_function(d) + end + + SUBROUTINE array_access_test_function(d) + double precision d(4) + + d(2)=5.5 + + END SUBROUTINE array_access_test_function + """ + ast, functions = fortran_parser.create_ast_from_string(test_string, "array_access_test") + ast_transforms.ParentScopeAssigner().visit(ast) + + assert ast.parent is None + assert ast.main_program.parent == None + + main_program = ast.main_program + # Both executed lines + for execution in main_program.execution_part.execution: + assert execution.parent == main_program + # call to the function + call_node = main_program.execution_part.execution[1] + assert isinstance(call_node, ast_internal_classes.Call_Expr_Node) + for arg in call_node.args: + assert arg.parent == main_program + + for subroutine in ast.subroutine_definitions: + + assert subroutine.parent == None + assert subroutine.execution_part.parent == subroutine + for execution in subroutine.execution_part.execution: + assert execution.parent == subroutine + +def test_fortran_frontend_module(): + """ + Tests that the Fortran frontend can parse array accesses and that the accessed indices are correct. + """ + test_string = """ + module test_module + implicit none + ! good enough approximation + integer, parameter :: pi = 4 + end module test_module + + PROGRAM access_test + implicit none + double precision d(4) + d(1)=0 + CALL array_access_test_function(d) + end + + SUBROUTINE array_access_test_function(d) + double precision d(4) + + d(2)=5.5 + + END SUBROUTINE array_access_test_function + """ + ast, functions = fortran_parser.create_ast_from_string(test_string, "array_access_test") + ast_transforms.ParentScopeAssigner().visit(ast) + + assert ast.parent is None + assert ast.main_program.parent == None + + module = ast.modules[0] + assert module.parent == None + specification = module.specification_part.specifications[0] + assert specification.parent == module + + +if __name__ == "__main__": + + test_fortran_frontend_parent() + test_fortran_frontend_module() diff --git a/tests/fortran/scope_arrays.py b/tests/fortran/scope_arrays.py new file mode 100644 index 0000000000..0eb0cf44b2 --- /dev/null +++ b/tests/fortran/scope_arrays.py @@ -0,0 +1,47 @@ +# Copyright 2023 ETH Zurich and the DaCe authors. All rights reserved. + +from dace.frontend.fortran import fortran_parser + +import dace.frontend.fortran.ast_transforms as ast_transforms +import dace.frontend.fortran.ast_internal_classes as ast_internal_classes + + +def test_fortran_frontend_parent(): + """ + Tests that the Fortran frontend can parse array accesses and that the accessed indices are correct. + """ + test_string = """ + PROGRAM scope_test + implicit none + double precision d(4) + double precision, dimension(5) :: arr + double precision, dimension(50:54) :: arr3 + CALL scope_test_function(d) + end + + SUBROUTINE scope_test_function(d) + double precision d(4) + double precision, dimension(50:54) :: arr4 + + d(2)=5.5 + + END SUBROUTINE scope_test_function + """ + + ast, functions = fortran_parser.create_ast_from_string(test_string, "array_access_test") + ast_transforms.ParentScopeAssigner().visit(ast) + visitor = ast_transforms.ScopeVarsDeclarations() + visitor.visit(ast) + + for var in ['d', 'arr', 'arr3']: + assert ('scope_test', var) in visitor.scope_vars + assert isinstance(visitor.scope_vars[('scope_test', var)], ast_internal_classes.Var_Decl_Node) + assert visitor.scope_vars[('scope_test', var)].name == var + + for var in ['d', 'arr4']: + assert ('scope_test_function', var) in visitor.scope_vars + assert visitor.scope_vars[('scope_test_function', var)].name == var + +if __name__ == "__main__": + + test_fortran_frontend_parent()