diff --git a/src/psyclone/domain/common/transformations/kernel_module_inline_trans.py b/src/psyclone/domain/common/transformations/kernel_module_inline_trans.py index b5f2b456bc..7a1bc07ed8 100644 --- a/src/psyclone/domain/common/transformations/kernel_module_inline_trans.py +++ b/src/psyclone/domain/common/transformations/kernel_module_inline_trans.py @@ -42,16 +42,17 @@ and move it to psyir/transformations/. ''' +from typing import Any, Optional, Union import warnings from psyclone.psyGen import Transformation, CodedKern from psyclone.psyir.transformations import TransformationError from psyclone.psyir.symbols import ( - ContainerSymbol, ImportInterface, - GenericInterfaceSymbol, RoutineSymbol, Symbol, SymbolError, SymbolTable) + ContainerSymbol, GenericInterfaceSymbol, RoutineSymbol, Symbol, + SymbolError) from psyclone.psyir.nodes import ( - Call, Container, FileContainer, Reference, Routine, ScopingNode, - IntrinsicCall) + Call, Container, FileContainer, IntrinsicCall, Reference, Routine, + ScopingNode) from psyclone.utils import transformation_documentation_wrapper @@ -83,17 +84,17 @@ def __str__(self): "Container of the call site.") # pylint: disable=too-many-branches - def validate(self, node, options=None, **kwargs): + def validate(self, + node: Union[CodedKern, Call], + options: Optional[dict[str, Any]] = None, + **kwargs): ''' Checks that the supplied node is a Kernel or Call and that it is possible to inline its PSyIR into the parent Container. :param node: the kernel or call which is the target of the transformation. - :type node: :py:class:`psyclone.psyGen.CodedKern` | - :py:class:`psyclone.psyir.nodes.Call` :param options: a dictionary with options for transformations. - :type options: Optional[Dict[str, Any]] :raises TransformationError: if the target node is not a sub-class of psyGen.CodedKern or psyir.nodes.Call or is an IntrinsicCall. @@ -104,8 +105,6 @@ def validate(self, node, options=None, **kwargs): and there's no Container at the call site to which to add the interface definition. :raises TransformationError: if the kernel cannot be safely inlined. - :raises TransformationError: if the target of the supplied call is - already module inlined. ''' if not options: @@ -188,9 +187,8 @@ def _validate_schedule(self, node, kname, kern_or_call, kernel_schedule): :raises TransformationError: if the called routine contains accesses to data declared in the same module scope or of unknown origin. - :raises TransformationError: if the called routine has the same - name as an existing Symbol in the calling scope (other than the - one representing the routine itself). + :raises TransformationError: if the called routine contains a local + Symbol that shadows a module name in its outer scope. ''' # We do not support kernels that use symbols representing data @@ -205,51 +203,19 @@ def _validate_schedule(self, node, kname, kern_or_call, kernel_schedule): f"{err.value}") from err # We can't transform subroutines that shadow top-level symbol module - # names, because we won't be able to bring this into the subroutine + # names, because we won't be able to bring them into the subroutine. + # (We could attempt to rename the local symbol.) symtab = kernel_schedule.ancestor(Container).symbol_table + ctr_names = [sym.name.lower() for sym in symtab.containersymbols] for scope in kernel_schedule.walk(ScopingNode): for symbol in scope.symbol_table.symbols: - for mod in symtab.containersymbols: - if (symbol.name.lower() == mod.name.lower() and - not isinstance(symbol, ContainerSymbol)): - raise TransformationError( - f"{kern_or_call} '{kname}' cannot be module-" - f"inlined because the subroutine shadows the " - f"symbol name of the module container " - f"'{symbol.name}'.") - - # If the symbol already exists at the call site it must be referring - # to a Routine - existing_symbol = node.scope.symbol_table.lookup(kernel_schedule.name, - otherwise=None) - if existing_symbol and not isinstance(existing_symbol, RoutineSymbol): - raise TransformationError( - f"Cannot module-inline {kern_or_call} '{kname}' because " - f"symbol '{existing_symbol}' with the same name already " - f"exists and changing the name of module-inlined " - f"subroutines is not supported yet.") - - # Check that the associated Routine isn't already present in the - # Container. Strictly speaking, we should check that the interface of - # any existing Routine matches that required by the Call but for now - # we live with the possibility of a false positive resulting in a - # refusal to module inline. - parent_container = node.ancestor(Container) - for routine in parent_container.walk(Routine, stop_type=Routine): - if routine.name.lower() == kname.lower(): - # Compare the routine to be inlined with the one that - # is already present. - new_rts = self._prepare_code_to_inline([kernel_schedule]) - if len(new_rts) == 1 and routine == new_rts[0]: - # It's the same so we can proceed (although all we need to - # do is update the RoutineSymbol referenced by the Call.) - return - raise TransformationError( - f"{kern_or_call} '{kname}' cannot be module inlined " - f"into Container '{parent_container.name}' because " - f"a *different* routine with that name " - f"already exists and versioning of module-inlined " - f"subroutines is not implemented yet.") + if (symbol.name.lower() in ctr_names and + not isinstance(symbol, ContainerSymbol)): + raise TransformationError( + f"{kern_or_call} '{kname}' cannot be module-" + f"inlined because the subroutine contains a symbol " + f"'{symbol.name}' which shadows the name of a module " + f"in the outer scope.") @staticmethod def _prepare_code_to_inline( @@ -330,117 +296,27 @@ def _prepare_code_to_inline( module_symbol.name) return copied_routines - @staticmethod - def _rm_imported_routine_symbol(symbol: Symbol, - schedule: Routine, - table: SymbolTable): - ''' - If the named symbol is in the supplied table (or an ancestor) *and* is - an import then it is removed. If the Container from which it was being - imported no longer has any imports associated with it then the - ContainerSymbol is also removed. - - TODO #2846 - this method is only required because at present we do not - rename inlined routines. However, not only is removing an imported - routine symbol dangerous, it also will not work if it should happen - that the same Routine is brought into scope through multiple - wildcard imports. - - :param symbol: the symbol to remove. - :param schedule: the Routine that is associated with this symbol. - :type table: the table from which to remove the symbol. - - ''' - if symbol.is_unresolved: - # The symbol is unresolved (at the call site) but we now have the - # source of the target routine and thus can find out which - # Container it is in. - cntr = schedule.ancestor(Container, excluding=FileContainer) - if cntr: - cntr_name = cntr.name - cntr_sym = table.lookup(cntr_name) - # Now we have a ContainerSymbol, we can change the - # interface of called_sym and proceed in exactly the - # same way as if it had been resolved originally. - symbol.interface = ImportInterface(cntr_sym) - - if not symbol.is_import: - return - - # The symbol is in the table (or an outer scope) and is - # imported. We therefore remove it and potentially the ContainerSymbol - # from which it is imported. - csym = symbol.interface.container_symbol - # Find the table containing the symbol we're going to remove. - actual_table = (symbol.find_symbol_table(table.node) if - symbol.name not in table else table) - # Find the table containing the ContainerSymbol from which - # the symbol is imported. - ctable = (csym.find_symbol_table(table.node) if - table.node else table) - remove_csym = (ctable.symbols_imported_from(csym) == [symbol] and - not csym.wildcard_import) - if csym.wildcard_import: - # The Routine is brought into scope via a wildcard import. We have - # to rename it on import to avoid a clash with the newly inlined - # Routine. - # TODO #2846 - if the same Routine is also brought into scope - # through some other wildcard import then this renaming doesn't - # help. The only solution to this is to rename the module-inlined - # Routine (or proceed to fully inline it). - KernelModuleInlineTrans._rename_import(ctable, csym, symbol.name) - # pylint:disable-next=protected-access - actual_table._symbols.pop(symbol.name.lower()) - if remove_csym: - actual_table.remove(csym) - - @staticmethod - def _rename_import(table: SymbolTable, - csym: ContainerSymbol, - name: str) -> None: - ''' - Adds a new RoutineSymbol imported from `csym` with its original name - set to be the supplied name while having an auto-generated actual name. - If there is already such a Symbol then this routine does nothing. - - :param table: the SymbolTable to update. - :param csym: the ContainerSymbol from which the routine symbol is - imported. - :param name: the name of the imported symbol. - - ''' - for isym in table.symbols_imported_from(csym): - if (isym.interface.orig_name and - isym.interface.orig_name.lower() == name.lower()): - # We already have a suitable import so we don't need - # another one. - return - - table.new_symbol( - f"old_{name}", - symbol_type=RoutineSymbol, - interface=ImportInterface( - csym, orig_name=name)) - - def apply(self, node, options=None, **kwargs): + def apply(self, + node: Union[CodedKern, Call], + options: dict[str, Any] = None, + update_all: Optional[bool] = False, + **kwargs): ''' Bring the implementation of this kernel/call into this Container. NOTE: when applying this transformation to a Kernel in a PSyKAl invoke, - *all* Kernels of that name in that invoke are marked as inlined. + by default *only* that Kernel call is updated (and marked as inlined). Similarly, when applied to a Call to a Routine in a particular scope, - all Calls to a routine of that name in that scope are updated. + *only* that Call is updated. This behaviour may be changed using the + `update_all=True` option. :param node: the Kernel or Call to module-inline. - :type node: :py:class:`psyclone.psyGen.CodedKern` | - :py:class:`psyclone.psyir.nodes.Call` :param options: a dictionary with options for transformations. - :type options: Optional[Dict[str, Any]] + :param update_all: by default, only the supplied Kernel or Call is + updated by this transformation. If this argument is set to True then + this behaviour is modified and all calls to this Kernel/Routine in + the associated program unit are updated. ''' - if isinstance(node, CodedKern) and node.module_inline: - # This PSyKal Kernel is already module inlined. - return - if options: # TODO 2668 - options dict is deprecated. warnings.warn(self._deprecation_warning, DeprecationWarning, 2) @@ -492,7 +368,7 @@ def apply(self, node, options=None, **kwargs): # Deal with the RoutineSymbol that is in scope at the call site. sym_in_ctr = None - shadowed_sym = None + #shadowed_sym = None if called_sym and (called_sym.is_import or called_sym.is_unresolved): table = called_sym.find_symbol_table(node) @@ -502,144 +378,88 @@ def apply(self, node, options=None, **kwargs): # update any other Calls to it (at the end of this method). sym_in_ctr = called_sym - self._rm_imported_routine_symbol(called_sym, codes_to_inline[0], - table) + #self._rm_imported_routine_symbol(called_sym, codes_to_inline[0], + # table) # Double check that this import is not shadowing a routine we've # already module-inlined. - if table.node.parent: - # There is a scope outside the one that contained the - # RoutineSymbol. - caller_cntr_table = table.node.parent.scope.symbol_table - # Look to see whether it also contains a symbol matching - # the name of the called routine. - shadowed_sym = caller_cntr_table.lookup(called_sym.name, - otherwise=None) - if shadowed_sym: - caller_cntr_table = shadowed_sym.find_symbol_table( - table.node.parent) - if not isinstance(caller_cntr_table.node, FileContainer): - # It is shadowing an outer symbol that is in a - # Container (not a FileContainer) so we just need to - # update the call to point to the outer symbol. - node.routine.symbol = shadowed_sym - if not (shadowed_sym.is_import or - shadowed_sym.is_unresolved): - # The outer symbol is local to this Container so - # there's nothing else to do. - return +# if table.node.parent: +# # There is a scope outside the one that contained the +# # RoutineSymbol. +# caller_cntr_table = table.node.parent.scope.symbol_table +# # Look to see whether it also contains a symbol matching +# # the name of the called routine. +# shadowed_sym = caller_cntr_table.lookup(called_sym.name, +# otherwise=None) +# if shadowed_sym: +# caller_cntr_table = shadowed_sym.find_symbol_table( +# table.node.parent) +# if not isinstance(caller_cntr_table.node, FileContainer): +# # It is shadowing an outer symbol that is in a +# # Container (not a FileContainer) so we just need to +# # update the call to point to the outer symbol. +# node.routine.symbol = shadowed_sym +# if not (shadowed_sym.is_import or +# shadowed_sym.is_unresolved): +# # The outer symbol is local to this Container so +# # there's nothing else to do. +# return updated_routines = self._prepare_code_to_inline(codes_to_inline) # The Container into which we will inline the Routine(s). container = node.ancestor(Container) + name_map = {} for code_to_inline in updated_routines: - # Does the Container already have this Routine? - sym = container.symbol_table.lookup( - code_to_inline.name, - scope_limit=container, - otherwise=None) - if not sym: - # If it doesn't exist already, module-inline the subroutine by - # inserting the relevant code into the tree. We need to set - # the visibility of the routine's symbol to be private. - sym = code_to_inline.symbol - sym.visibility = Symbol.Visibility.PRIVATE - container.addchild(code_to_inline.detach()) - - elif sym.is_import: - # The RoutineSymbol is imported into the table. We must - # therefore update its interface and potentially remove the - # ContainerSymbol (from which it is imported) altogether. - self._rm_imported_routine_symbol(sym, - code_to_inline, - container.symbol_table) - # Inline the code. This will automatically add the - # associated RoutineSymbol into the Container. - code_to_inline = code_to_inline.detach() - container.addchild(code_to_inline) - newsym = container.symbol_table.lookup(code_to_inline.name) - newsym.visibility = Symbol.Visibility.PRIVATE - - elif sym.is_unresolved: - # The Symbol in the Container scope is unresolved. However, - # we've found the source so we now know where it comes from. - self._rm_imported_routine_symbol(sym, - code_to_inline, - container.symbol_table) - cntr = code_to_inline.ancestor(Container, - excluding=FileContainer) - if cntr: - # Inline the code. This will automatically add the - # associated RoutineSymbol into the Container. - code_to_inline = code_to_inline.detach() - container.addchild(code_to_inline) - newsym = container.symbol_table.lookup(code_to_inline.name) - newsym.visibility = Symbol.Visibility.PRIVATE - - # All Calls to a routine of the same name in the same scope as the - # target node must refer to the same Symbol. - target_name = sym.name.lower() - target_sym = node.scope.symbol_table.lookup(target_name) - for call in node.ancestor(Routine).walk(Call): - name = call.routine.symbol.name.lower() - if name == target_name: - call.routine.symbol = target_sym - # All Calls that referred to this Symbol must also be updated. Take - # care that the name matches as sym_in_ctr might be an interface. - if sym_in_ctr and sym_in_ctr.name == target_sym.name: - for call in container.walk(Call): - if call.routine.symbol is sym_in_ctr: + new_name = code_to_inline.name+"_inlined_" + new_sym = container.symbol_table.new_symbol( + new_name, symbol_type=RoutineSymbol) + new_sym.copy_properties(code_to_inline.symbol, + exclude_interface=True) + new_sym.visibility = Symbol.Visibility.PRIVATE + # Add the routine code into this Container + code_to_inline = code_to_inline.detach() + code_to_inline.symbol = new_sym + container.addchild(code_to_inline) + # Keep a record of the new and original names. + name_map[code_to_inline.name] = new_sym + + if update_all: + # All Calls to a routine of the same name in the same scope as + # the target node must refer to the same Symbol. + #target_name = sym.name.lower() + target_sym = node.scope.symbol_table.lookup(caller_name) + for call in node.ancestor(Routine).walk(Call): + name = call.routine.symbol.name.lower() + if name == caller_name: call.routine.symbol = target_sym + # All Calls that referred to this Symbol must also be updated. + # Take care that the name matches as sym_in_ctr might be an + # interface. + if sym_in_ctr and sym_in_ctr.name == target_sym.name: + for call in container.walk(Call): + if call.routine.symbol is sym_in_ctr: + call.routine.symbol = target_sym if interface_sym: - # Deal with the interface symbol - remove any existing import and - # then make sure the local symbol is private. - self._rm_imported_routine_symbol(interface_sym, - codes_to_inline[0], - callsite_table) - if shadowed_sym: - self._rm_imported_routine_symbol(shadowed_sym, - codes_to_inline[0], - caller_cntr_table) - if caller_name != interface_sym.name: - # If the interface was originally renamed on import, then we - # must create a new symbol with the local name. - new_sym = GenericInterfaceSymbol( - caller_name, routines=[(RoutineSymbol("dummy"), True)]) - new_sym.copy_properties(interface_sym) - else: - # Otherwise we can use the existing symbol. - new_sym = interface_sym - container.symbol_table.add(new_sym) - interface_sym.visibility = Symbol.Visibility.PRIVATE - interface_sym.replace_symbols_using(container.symbol_table) - else: - # No interface but was the original routine symbol renamed - # on import? - if caller_name != external_callee_name: - # It was so we need to rename the inlined routine. - sym = node.scope.symbol_table.lookup(external_callee_name) - table = sym.find_symbol_table(node) - table.rename_symbol(sym, caller_name) - new_sym = sym - else: - new_sym = node.scope.symbol_table.lookup(caller_name) + # Deal with the interface symbol - create a new, local + # private version. + new_sym = container.symbol_table.new_symbol( + interface_sym.name+"_inlined_", + symbol_type=GenericInterfaceSymbol, + routines=[(sym, True) for sym in name_map.values()], + visibility=Symbol.Visibility.PRIVATE) # Update the Kernel to point to the updated PSyIR. if isinstance(node, CodedKern): - cntr = node.ancestor(Container) - # TODO #2846 - since we do not currently rename module-inlined - # routines, inlining just one instance of a Kernel call and - # subsequently transforming it (e.g. by adding ACC ROUTINE) would - # inhibit all further transformations of any other calls to that - # kernel (that relied upon inlining). Therefore, for now, we update - # *all* calls to this particular kernel in the current module to - # point to the module-inlined version. - for kern in cntr.walk(CodedKern, stop_type=CodedKern): - if kern.name == node.name: - # pylint: disable=protected-access - kern._schedules = updated_routines - kern.routine = Reference(new_sym) + node.routine = Reference(new_sym) + # pylint: disable=protected-access + node._schedules = updated_routines + else: + # Update the Call to point to the inlined routine. + node.routine.symbol = new_sym + + # In theory we could remove the import of the original routine but + # that is dangerous and unnecessary so we don't. diff --git a/src/psyclone/domain/lfric/lfric_kern.py b/src/psyclone/domain/lfric/lfric_kern.py index 4db6097a87..e9cf975db3 100644 --- a/src/psyclone/domain/lfric/lfric_kern.py +++ b/src/psyclone/domain/lfric/lfric_kern.py @@ -129,6 +129,9 @@ def __init__(self): api_config = Config.get().api_conf("lfric") self._argument_kinds = {api_config.default_kind["real"], api_config.default_kind["integer"]} + # TODO #2054 - this 'routine' property can be replaced once this + # class sub-classes Call. + self.routine: Optional[Reference] = None def reference_accesses(self) -> VariablesAccessMap: ''' diff --git a/src/psyclone/psyGen.py b/src/psyclone/psyGen.py index 1f8c926592..3b1e19c486 100644 --- a/src/psyclone/psyGen.py +++ b/src/psyclone/psyGen.py @@ -1172,11 +1172,12 @@ def __init__(self, self._opencl_options = {'local_size': 64, 'queue_number': 1} self.arg_descriptors = call.ktype.arg_descriptors - # If we have an ancestor InvokeSchedule then add the necessary - # symbols. # TODO #2054 - this 'routine' property can be replaced once this # class sub-classes Call. self.routine: Optional[Reference] = None + + # If we have an ancestor InvokeSchedule then add the necessary + # symbols. container = self.ancestor(Container) if container: symtab = container.symbol_table @@ -1213,6 +1214,16 @@ def get_callees(self): raise NotImplementedError( f"get_callees() must be overridden in class {self.__class__}") + @property + def name(self) -> str: + ''' + Override Kern.name to use the RoutineSymbol name if one is available. + + ''' + if self.routine: + return self.routine.symbol.name + return super().name + @property def opencl_options(self): ''' @@ -1309,7 +1320,7 @@ def lower_to_language_level(self) -> Node: ''' symtab = self.ancestor(InvokeSchedule).symbol_table - rsymbol = symtab.lookup(self._name) + rsymbol = symtab.lookup(self.name) # Create Call to the rsymbol with the argument expressions as children # of the new node diff --git a/src/psyclone/tests/domain/common/transformations/kernel_module_inline_trans_test.py b/src/psyclone/tests/domain/common/transformations/kernel_module_inline_trans_test.py index 3765cdc59b..27fbcc8936 100644 --- a/src/psyclone/tests/domain/common/transformations/kernel_module_inline_trans_test.py +++ b/src/psyclone/tests/domain/common/transformations/kernel_module_inline_trans_test.py @@ -213,37 +213,36 @@ def test_validate_name_clashes(): ''' Test that if the module-inline transformation finds the kernel name already used in the Container scope, it raises the appropriate error''' # Use LFRic example with a repeated CodedKern - psy, _ = get_invoke("4.6_multikernel_invokes.f90", "lfric", idx=0, - dist_mem=False) - schedule = psy.invokes.invoke_list[0].schedule + psy, invoke = get_invoke("4.6_multikernel_invokes.f90", "lfric", idx=0, + dist_mem=False) + schedule = invoke.schedule coded_kern = schedule.children[0].loop_body[0] inline_trans = KernelModuleInlineTrans() - # Check that name clashes which are not subroutines are detected + # Check that name clashes which are not subroutines are handled schedule.symbol_table.add(DataSymbol("ru_code", ScalarType.real_type())) - with pytest.raises(TransformationError) as err: - inline_trans.apply(coded_kern) - assert ("Cannot module-inline Kernel 'ru_code' because symbol " - "'ru_code: DataSymbol, Automatic>' with " - "the same name already exists and changing the name of " - "module-inlined subroutines is not supported yet." - in str(err.value)) + inline_trans.apply(coded_kern) + assert coded_kern.name == "ru_code_inlined_" - # TODO #898. Manually force removal of previously added symbol - # symbol_table.remove() is not implemented yet. - schedule.symbol_table._symbols.pop("ru_code") - # Also remove the RoutineSymbol representing the Kernel that is - # automatically added to the Container during PSy-layer construction. - schedule.ancestor(Container).symbol_table._symbols.pop("ru_code") - # Check that if a subroutine with the same name already exists and it is - # not identical, it fails. - schedule.parent.addchild(Routine.create("ru_code")) - with pytest.raises(TransformationError) as err: - inline_trans.apply(coded_kern) - assert ("Kernel 'ru_code' cannot be module inlined into Container " - "'multikernel_invokes_7_psy' because a *different* routine with " - "that name already exists and versioning of module-inlined " - "subroutines is not implemented yet.") in str(err.value) + +def test_validate_routine_name_clashes(): + ''' Test that the module-inline transformation copes when the kernel name + is already used by another Routine in the Container scope.''' + inline_trans = KernelModuleInlineTrans() + # Use LFRic example with a repeated CodedKern + psy, invoke = get_invoke("4.6_multikernel_invokes.f90", "lfric", idx=0, + dist_mem=False) + schedule = invoke.schedule + coded_kern = schedule.children[0].loop_body[0] + + # Check that if a subroutine with the same name already exists + # everything still works. + schedule.parent.addchild(Routine.create("ru_code_inlined_")) + inline_trans.apply(coded_kern) + assert coded_kern.name == "ru_code_inlined__1" + + output = str(psy.gen) + assert output.count("end subroutine ru_code_inlined__1") == 1 def test_validate_unsupported_symbol_shadowing(fortran_reader, monkeypatch): @@ -276,8 +275,8 @@ def test_validate_unsupported_symbol_shadowing(fortran_reader, monkeypatch): with pytest.raises(TransformationError) as err: inline_trans.apply(kern_call) assert ("Kernel 'compute_cv_code' cannot be module-inlined because the " - "subroutine shadows the symbol name of the module container " - "'external_mod'." in str(err.value)) + "subroutine contains a symbol 'external_mod' which shadows the " + "name of a module in the outer scope." in str(err.value)) # Repeat the same with a wildcard import psyir = fortran_reader.psyir_from_source(''' @@ -298,8 +297,8 @@ def test_validate_unsupported_symbol_shadowing(fortran_reader, monkeypatch): with pytest.raises(TransformationError) as err: inline_trans.apply(kern_call) assert ("Kernel 'compute_cv_code' cannot be module-inlined because the " - "subroutine shadows the symbol name of the module container " - "'external_mod'." in str(err.value)) + "subroutine contains a symbol 'external_mod' which shadows the " + "name of a module in the outer scope." in str(err.value)) # But it is fine if it shadows itself psyir = fortran_reader.psyir_from_source(''' @@ -322,8 +321,8 @@ def test_validate_unsupported_symbol_shadowing(fortran_reader, monkeypatch): inline_trans.apply(kern_call) - # The RoutineSymbol should no longer be an import. - rsym = container.symbol_table.lookup("compute_cv_code") + # A RoutineSymbol should have been added to the Container symbol table. + rsym = container.symbol_table.lookup("compute_cv_code_inlined_") assert isinstance(rsym, RoutineSymbol) assert not rsym.is_import assert rsym.visibility == Symbol.Visibility.PRIVATE @@ -448,33 +447,34 @@ def test_module_inline_apply_transformation(tmpdir, fortran_writer): # The new inlined routine must now exist and be private. routine_sym = kern_call.ancestor(Container).symbol_table.lookup( - "compute_cv_code") + "compute_cv_code_inlined_") assert routine_sym assert routine_sym.visibility == Symbol.Visibility.PRIVATE - assert kern_call.ancestor(Container).children[1].name == "compute_cv_code" + assert (kern_call.ancestor(Container).children[1].name == + "compute_cv_code_inlined_") assert (kern_call.ancestor(Container).symbol_table. - lookup("compute_cv_code").is_modulevar) + lookup("compute_cv_code_inlined_").is_modulevar) # Generate the code - code = str(psy.gen) - assert 'subroutine compute_cv_code(i, j, cv, p, v)' in code + code = str(psy.gen).lower() + assert 'subroutine compute_cv_code_inlined_(i, j, cv, p, v)' in code - # And the import has been remove from both, so check that the associated - # use no longer exists - assert 'use compute_cv_mod' not in code.lower() + # The import for the original Kern is unchanged + assert 'use compute_cv_mod, only : compute_cv_code\n' in code # Do the check again because repeating the call resets some aspects and we # need to see if the second call still works as expected gen = str(psy.gen) - assert 'subroutine compute_cv_code(i, j, cv, p, v)' in gen - assert 'use compute_cv_mod' not in gen - assert gen.count("subroutine compute_cv_code(") == 1 + assert 'subroutine compute_cv_code_inlined_(i, j, cv, p, v)' in gen + assert 'use compute_cv_mod, only : compute_cv_code\n' in gen + assert gen.count("subroutine compute_cv_code_inlined_(") == 1 # And it is valid code assert GOceanBuild(tmpdir).code_compiles(psy) -def test_module_inline_apply_kernel_in_multiple_invokes(tmpdir): +@pytest.mark.parametrize("do_all", [True, False]) +def test_module_inline_apply_kernel_in_multiple_invokes(tmpdir, do_all): ''' Check that module-inline works as expected when the same kernel is provided in different invokes''' # Use LFRic example with the kernel 'testkern_qr_mod' repeated once in @@ -487,27 +487,38 @@ def test_module_inline_apply_kernel_in_multiple_invokes(tmpdir): assert gen.count("use testkern_qr_mod, only : testkern_qr_code") == 1 assert gen.count("end subroutine testkern_qr_code") == 0 - # Module inline kernel in invoke 1 - inline_trans = KernelModuleInlineTrans() + # Module inline kernel in the first invoke + mod_inline_trans = KernelModuleInlineTrans() artrans = ACCRoutineTrans() schedule1 = psy.invokes.invoke_list[0].schedule for coded_kern in schedule1.walk(CodedKern): if coded_kern.name == "testkern_qr_code": - inline_trans.apply(coded_kern) + mod_inline_trans.apply(coded_kern, update_all=do_all) artrans.apply(coded_kern) + break gen = str(psy.gen) - - # After this, both invokes use the inlined top-level subroutine. - # Module-inlining kernel in invoke 2 should have no effect. - schedule1 = psy.invokes.invoke_list[1].schedule - for coded_kern in schedule1.walk(CodedKern): - if coded_kern.name == "testkern_qr_code": - inline_trans.apply(coded_kern) - gen = str(psy.gen) - # After this, no imports are remaining and both use the same - # top-level implementation - assert gen.count("use testkern_qr_mod, only : testkern_qr_code") == 0 - assert gen.count("end subroutine testkern_qr_code") == 1 + assert gen.count("end subroutine testkern_qr_code_inlined_") == 1 + + # After this, the other calls should still be to the original kernel. + # Transform these next (and check that there are three of them). + if not do_all: + schedule1 = psy.invokes.invoke_list[1].schedule + count = 0 + for coded_kern in schedule1.walk(CodedKern): + if coded_kern.name == "testkern_qr_code": + count += 1 + mod_inline_trans.apply(coded_kern) + assert count == 3 + gen = str(psy.gen) + # After this, the imports remain unchanged and all four calls are now + # to either separate, inlined versions of the routine or just a single + # inlined version (if do_all is True). + assert gen.count("use testkern_qr_mod, only : testkern_qr_code\n") == 1 + assert gen.count("call testkern_qr_code_inlined_") == 4 + if do_all: + assert gen.count("end subroutine testkern_qr_code_inlined") == 1 + else: + assert gen.count("end subroutine testkern_qr_code_inlined") == 4 # And it is valid code assert LFRicBuild(tmpdir).code_compiles(psy) @@ -531,19 +542,14 @@ def test_module_inline_apply_polymorphic_kernel_in_multiple_invokes(tmpdir): # a CodeBlock. artrans.apply(coded_kern, options={"force": True}) output = str(psy.gen).lower() - assert "subroutine mixed_code_32" in output + assert "subroutine mixed_code_32_inlined_" in output assert output.count("!$acc routine seq") == 2 - assert "subroutine mixed_code_64" in output - # Since we don't currently rename module-inlined kernels (TODO #2846), - # module-inlining just one instance means that calls to that same Kernel - # throughout the whole module use the newly-inlined version. - assert """subroutine invoke_1(scalar_r_bl, field_r_bl, \ -operator_r_def, f1, f2, m1, a, m2, istp, qr) - use quadrature_xyoz_mod, only : quadrature_xyoz_proxy_type, \ -quadrature_xyoz_type - use function_space_mod, only : basis, diff_basis - real""" in output - assert "mixed_kernel_mod" not in output + assert "subroutine mixed_code_64_inlined_" in output + assert "call mixed_code_inlined_" in output + # We've only module-inlined one instance so the other calls to that same + # Kernel should use the original version. + assert "mixed_kernel_mod" in output + assert output.count("call mixed_code(") == 1 assert LFRicBuild(tmpdir).code_compiles(psy) @@ -558,7 +564,7 @@ def test_module_inline_apply_with_sub_use(tmpdir): inline_trans.apply(kern_call) gen = str(psy.gen) # check that the subroutine has been inlined - assert 'subroutine bc_ssh_code(ji, jj, istep, ssha, tmask)' in gen + assert 'subroutine bc_ssh_code_inlined_(ji, jj, istep, ssha, tmask)' in gen # check that the use within the subroutine exists assert 'use grid_mod' in gen # check that the associated psy use does not exist @@ -578,18 +584,19 @@ def test_module_inline_apply_same_kernel(tmpdir): inline_trans.apply(kern_calls[0]) gen = str(psy.gen) # check that the subroutine has been inlined - assert 'subroutine compute_cu_code(' in gen - # check that the associated psy "use" does not exist - assert 'use compute_cu_mod' not in gen + assert 'subroutine compute_cu_code_inlined_(' in gen + # check that the associated "use" is unchanged + assert 'use compute_cu_mod' in gen # check that the subroutine has only been inlined once - count = count_lines(gen, "subroutine compute_cu_code(") + count = count_lines(gen, "subroutine compute_cu_code_inlined_(") assert count == 1, "Expecting subroutine to be inlined once" assert GOceanBuild(tmpdir).code_compiles(psy) # Calling the transformation on a second call to the same kernel - # should have no effect. + # should create a new, private copy of the routine. inline_trans.apply(kern_calls[1]) gen2 = str(psy.gen) - assert gen2 == gen + assert 'subroutine compute_cu_code_inlined__1(' in gen2 + assert GOceanBuild(tmpdir).code_compiles(psy) def test_module_inline_apply_bring_in_non_local_symbols( @@ -823,9 +830,9 @@ def test_module_inline_lfric(tmpdir, monkeypatch, annexed, dist_mem): inline_trans.apply(kern_call) gen = str(psy.gen) # check that the subroutine has been inlined - assert 'subroutine ru_code(' in gen - # check that the associated psy "use" does not exist - assert 'use ru_kernel_mod' not in gen + assert 'subroutine ru_code_inlined_(' in gen + # check that the associated psy "use" is still present + assert 'use ru_kernel_mod' in gen # Check that we can subsequently transform the inlined kernel. omptrans = OMPDeclareTargetTrans() omptrans.apply(kern_call) @@ -845,7 +852,7 @@ def test_module_inline_with_interfaces(tmpdir): kern_calls = invoke.schedule.walk(CodedKern) inline_trans = KernelModuleInlineTrans() inline_trans.apply(kern_calls[0]) - sym = kern_calls[0].scope.symbol_table.lookup("mixed_code") + sym = kern_calls[0].scope.symbol_table.lookup("mixed_code_inlined_") # Check that the interface symbol is declared and is private. assert isinstance(sym, GenericInterfaceSymbol) assert sym.visibility == Symbol.Visibility.PRIVATE @@ -855,10 +862,10 @@ def test_module_inline_with_interfaces(tmpdir): gen = str(psy.gen).lower() # Both the caller and the callee are in the file and use the interface # name. - assert "call mixed_code(" in gen - assert "interface mixed_code" in gen - assert "subroutine mixed_code_64(" in gen - assert "subroutine mixed_code_32(" in gen + assert "call mixed_code_inlined_(" in gen + assert "interface mixed_code_inlined_" in gen + assert "subroutine mixed_code_64_inlined_(" in gen + assert "subroutine mixed_code_32_inlined_(" in gen # And it is valid code assert LFRicBuild(tmpdir).code_compiles(psy) @@ -894,8 +901,8 @@ def test_module_inline_with_renamed_import(monkeypatch, psyir = fortran_reader.psyir_from_source(code) intrans.apply(psyir.walk(Call)[0]) result = fortran_writer(psyir) - assert "call local_name(var)" in result - assert "subroutine local_name(arg)" in result + assert "call my_sub_inlined_(var)" in result + assert "subroutine my_sub_inlined_(arg)" in result def test_module_inline_interface_with_renamed_import(monkeypatch, @@ -925,6 +932,7 @@ def test_module_inline_interface_with_renamed_import(monkeypatch, end module my_mod ''') intrans = KernelModuleInlineTrans() + # No parent module in which to insert interface definition. code = '''\ program my_prog implicit none @@ -950,51 +958,9 @@ def test_module_inline_interface_with_renamed_import(monkeypatch, psyir = fortran_reader.psyir_from_source(code) intrans.apply(psyir.walk(Call)[0]) result = fortran_writer(psyir) - assert "interface local_name" in result - assert "call local_name(var)" in result - assert "subroutine my_sub(arg)" in result - - -def test_rm_imported_routine_symbol(fortran_reader): - ''' - Tests for the _rm_imported_routine_symbol() utility method. - ''' - intrans = KernelModuleInlineTrans() - table = SymbolTable() - # Does nothing if the symbol is not an import. - here = table.new_symbol("here") - intrans._rm_imported_routine_symbol(here, None, table) - assert "here" in table - csym = table.new_symbol("from_here", symbol_type=ContainerSymbol) - # Update the symbol so that it is imported from a Container. - here.interface = ImportInterface(csym) - intrans._rm_imported_routine_symbol(here, None, table) - # Both it and the Container should have been removed. - assert "here" not in table - assert "from_here" not in table - # Repeat for the case where the Container has a wildcard import. - csym = table.new_symbol("from_here", symbol_type=ContainerSymbol) - psyir = fortran_reader.psyir_from_source(''' - module from_here - contains - subroutine here - end subroutine here - end module from_here - ''') - sched = psyir.walk(Routine)[0] - csym.wildcard_import = True - here = table.new_symbol("here", interface=ImportInterface(csym)) - intrans._rm_imported_routine_symbol(here, sched, table) - # Only the Symbol should have been removed (not the ContainerSymbol). - assert "here" not in table - assert "from_here" in table - # There should be an import of a new symbol that prevents a clash between - # an inlined routine and the imported symbol of the same name. - assert table.lookup("old_here").interface.orig_name == "here" - # Repeat - we should not get another imported symbol. - here = table.new_symbol("here", interface=ImportInterface(csym)) - intrans._rm_imported_routine_symbol(here, sched, table) - assert len(table.symbols_imported_from(csym)) == 1 + assert "interface my_interface_inlined_" in result + assert "call my_interface_inlined_(var)" in result + assert "subroutine my_sub_inlined_(arg)" in result @pytest.mark.parametrize( @@ -1051,37 +1017,42 @@ def test_psyir_mod_inline(fortran_reader, fortran_writer, tmpdir, intrans.apply(calls[0]) routines = container.walk(Routine) assert len(routines) == 2 - assert routines[0].name in ["a_sub", "my_sub"] - assert routines[1].name in ["a_sub", "my_sub"] + assert routines[0].name in ["a_sub", "my_sub_inlined_"] + assert routines[1].name in ["a_sub", "my_sub_inlined_"] # Local copy of routine must be private and in Container symbol table. - rsym = container.symbol_table.lookup("my_sub") + rsym = container.symbol_table.lookup("my_sub_inlined_") assert rsym.visibility == Symbol.Visibility.PRIVATE output = fortran_writer(psyir) assert "subroutine a_sub" in output - assert "subroutine my_sub" in output - assert "use my_mod, only : my_interface, my_other_sub\n" in output + assert "subroutine my_sub_inlined_" in output + # USE statement left unchanged (for safety). + assert "use my_mod, only : my_interface, my_other_sub, my_sub\n" in output # Module inline the target of the second call. intrans.apply(calls[1]) # Local copy of routine must be private and in Container symbol table. - rsym = container.symbol_table.lookup("my_other_sub") + rsym = container.symbol_table.lookup("my_other_sub_inlined_") assert rsym.visibility == Symbol.Visibility.PRIVATE output = fortran_writer(psyir) - assert "use my_mod, only : my_interface\n" in output + # USE statement should still be unmodified. + assert "use my_mod, only : my_interface, my_other_sub, my_sub\n" in output - # Finally, inline the call to the interface. This should then remove all - # imports from 'my_mod'. + # Finally, inline the call to the interface. intrans.apply(calls[2]) routines = container.walk(Routine) output = fortran_writer(psyir) - assert "use my_mod" not in output - assert "subroutine my_other_sub" in output - assert "interface my_interface" in output + assert "use my_mod, only : my_interface, my_other_sub, my_sub\n" in output + assert "subroutine my_other_sub_inlined_" in output + assert "interface my_interface_inlined_" in output # Check that the calls themselves are unaffected. - assert "call my_sub(a)" in output - assert "call my_other_sub(b)" in output - assert "call my_interface(b)" in output - assert Compile(tmpdir).string_compiles(output) + assert "call my_sub_inlined_(a)" in output + assert "call my_other_sub_inlined_(b)" in output + assert "call my_interface_inlined_(b)" in output + # The USE statement is left for safety but is actually not required so + # remove it to permit compilation check. + to_build = output.replace( + "use my_mod, only : my_interface, my_other_sub, my_sub\n", "") + assert Compile(tmpdir).string_compiles(to_build) @pytest.mark.usefixtures("clear_module_manager_instance") @@ -1120,14 +1091,12 @@ def test_mod_inline_no_container(fortran_reader, fortran_writer, tmpdir, intrans.apply(call) assert len(prog_psyir.children) == 2 - assert set(child.name for child in prog_psyir.children) == {"my_sub", - "my_prog"} + assert set(child.name for child + in prog_psyir.children) == {"my_sub_inlined_", "my_prog"} output = fortran_writer(prog_psyir) assert "use my_mod" not in output - # Now that we've 'privatised' the target of the call, the code can be - # compiled standalone. assert Compile(tmpdir).string_compiles(output) @@ -1173,23 +1142,19 @@ def test_mod_inline_from_wildcard_import(fortran_reader, fortran_writer, calls = prog_psyir.walk(Call) intrans.apply(calls[0]) output = fortran_writer(prog_psyir) - # Wildcard import is preserved but the original routine is renamed - # to avoid a clash. This renaming avoids clashing with any pre-existing - # symbols. - assert "use my_mod, old_my_sub_1=>my_sub" in output - assert "call my_sub" in output + # Wildcard import is preserved + assert "use my_mod\n" in output + assert "call my_sub_inlined_" in output assert ('''end program my_prog -subroutine my_sub(arg)''' in output) - # Two calls in the same scope to a routine of the same name must be to - # the same RoutineSymbol +subroutine my_sub_inlined_(arg)''' in output) new_calls = prog_psyir.walk(Call) - assert new_calls[0].routine.symbol is new_calls[1].routine.symbol + assert new_calls[0].routine.symbol is not new_calls[1].routine.symbol # Apply the transformation to the second call. This should fail to validate # as there's nothing to do. - with pytest.raises(TransformationError) as err: - intrans.validate(calls[1]) - assert ("The target of 'call my_sub(b)' is already module inlined." - in str(err.value)) + intrans.apply(calls[1]) + output = fortran_writer(prog_psyir) + assert "subroutine my_sub_inlined__1" in output + assert "call my_sub_inlined__1(" in output # We can't compile this because of the use statement. @@ -1241,20 +1206,22 @@ def test_inline_of_shadowed_import(tmpdir, monkeypatch, fortran_reader, calls = prog_psyir.walk(Call) intrans.apply(calls[0]) assert (do_it.walk(Call)[0].routine.symbol is - container.symbol_table.lookup("my_sub")) + container.symbol_table.lookup("my_sub_inlined_")) # Call in second subroutine still refers to imported Symbol in local table. again = container.find_routine_psyir("and_again") assert (again.walk(Call)[0].routine.symbol is again.symbol_table.lookup("my_sub", scope_limit=again)) # Apply the transformation to the call in the second routine. intrans.apply(calls[1]) - # Now it should refer to the top-level, inlined RoutineSymbol. - assert calls[1].routine.symbol is container.symbol_table.lookup("my_sub") + # Now it should refer to another top-level, inlined RoutineSymbol. + assert calls[1].routine.symbol is container.symbol_table.lookup( + "my_sub_inlined__1") assert "my_mod" not in again.symbol_table - assert len(container.walk(Routine)) == 3 - # Cannot compile this because we still have a wildcard import from my_mod. + assert len(container.walk(Routine)) == 4 + # Cannot compile this because we still have a wildcard import from my_mod + # in 'do_it'. output = fortran_writer(prog_psyir) - assert "use my_mod, old_my_sub_1=>my_sub" in output + assert "use my_mod\n" in output def test_mod_inline_all_calls_updated(monkeypatch, fortran_reader): @@ -1299,12 +1266,13 @@ def test_mod_inline_all_calls_updated(monkeypatch, fortran_reader): for call in calls: assert call.routine.symbol is rt_sym0 intrans.apply(calls[0]) - # Since all calls previously referenced the same symbol, they should now - # all reference the new one. - rt_sym1 = container.symbol_table.lookup("my_sub") + # All calls previously referenced the same symbol but now the first one + # reference the new one. + rt_sym1 = container.symbol_table.lookup("my_sub_inlined_") assert rt_sym0 is not rt_sym1 - for call in calls: - assert call.routine.symbol is rt_sym1 + assert calls[0].routine.symbol is rt_sym1 + for call in calls[1:]: + assert call.routine.symbol is rt_sym0 def test_mod_inline_unresolved_sym_in_container(monkeypatch, fortran_reader): @@ -1355,11 +1323,8 @@ def test_mod_inline_unresolved_sym_in_container(monkeypatch, fortran_reader): csym = calls[0].scope.symbol_table.lookup("my_mod") rsym.interface = ImportInterface(csym) intrans.apply(calls[0]) - new_rt = container.symbol_table.lookup("my_sub") + new_rt = container.symbol_table.lookup("my_sub_inlined_") assert not (new_rt.is_import or new_rt.is_unresolved) - ctr_sym = container.symbol_table.lookup("my_mod") - (isym,) = container.symbol_table.symbols_imported_from(ctr_sym) - assert isym.interface.orig_name == "my_sub" def test_mod_inline_shared_wildcard_import(monkeypatch, tmp_path, @@ -1414,7 +1379,8 @@ def test_mod_inline_shared_wildcard_import(monkeypatch, tmp_path, # Check that eps20 has the correct interface in the inlined Routine. # The imports in the inlined routine should have been followed and # hence 'eps20' resolved. - inlined = container.find_routine_psyir("my_sub", allow_private=True) + inlined = container.find_routine_psyir("my_sub_inlined_", + allow_private=True) eps_sym = inlined.symbol_table.lookup("eps20") assert not eps_sym.is_unresolved assert eps_sym.is_import diff --git a/src/psyclone/tests/domain/gocean/goloop_test.py b/src/psyclone/tests/domain/gocean/goloop_test.py index af1072414e..a501e77727 100644 --- a/src/psyclone/tests/domain/gocean/goloop_test.py +++ b/src/psyclone/tests/domain/gocean/goloop_test.py @@ -267,6 +267,7 @@ def __init__(self): PSyIR node''' # pylint: disable=super-init-not-called, non-parent-init-called Node.__init__(self) # Ignore hierarchy constructors + self.routine = None # Test Loop containing kernels with different offsets gokern1 = GOKernMock() diff --git a/src/psyclone/tests/domain/gocean/transformations/globalstoargs_test.py b/src/psyclone/tests/domain/gocean/transformations/globalstoargs_test.py index 6d12aa1b7e..e13133ef92 100644 --- a/src/psyclone/tests/domain/gocean/transformations/globalstoargs_test.py +++ b/src/psyclone/tests/domain/gocean/transformations/globalstoargs_test.py @@ -104,8 +104,8 @@ def test_kernelimportstoargumentstrans_no_wildcard_import(): trans = KernelImportsToArguments() with pytest.raises(TransformationError) as err: trans.apply(kernel) - assert ("'kernel_with_use_code' contains accesses to 'rdt' which is " - "unresolved" in str(err.value)) + assert ("'kernel_with_use_code_inlined_' contains accesses to 'rdt' which " + "is unresolved" in str(err.value)) def test_kernelimportstoargumentstrans(fortran_writer): @@ -165,8 +165,8 @@ def test_kernelimportstoargumentstrans(fortran_writer): # Check the kernel code is generated as expected kernel_code = fortran_writer(kernel.get_callees()[0]) - assert ("subroutine kernel_with_use_code(ji, jj, istep, ssha, tmask, " - "rdt, magic)" in kernel_code) + assert ("subroutine kernel_with_use_code_inlined_(ji, jj, istep, ssha, " + "tmask, rdt, magic)" in kernel_code) assert "real(kind=go_wp), intent(in) :: rdt" in kernel_code assert "real(kind=go_wp), intent(inout) :: magic" in kernel_code @@ -174,7 +174,7 @@ def test_kernelimportstoargumentstrans(fortran_writer): # and argument call generated_code = str(psy.gen) assert "use model_mod, only : magic, rdt" in generated_code - assert ("call kernel_with_use_code(i, j, oldu_fld, cu_fld%data, " + assert ("call kernel_with_use_code_inlined_(i, j, oldu_fld, cu_fld%data, " "cu_fld%grid%tmask, rdt, magic)" in generated_code) assert invoke.schedule.symbol_table.lookup("model_mod") assert invoke.schedule.symbol_table.lookup("rdt") @@ -210,8 +210,8 @@ def create_data_symbol(arg): kernels = kernel.get_callees() kernel_code = fortran_writer(kernels[0]) - assert ("subroutine kernel_with_use_code(ji, jj, istep, ssha, tmask, rdt, " - "magic)" in kernel_code) + assert ("subroutine kernel_with_use_code_inlined_(ji, jj, istep, ssha, " + "tmask, rdt, magic)" in kernel_code) assert "integer, intent(in) :: rdt" in kernel_code @@ -261,15 +261,15 @@ def test_kernelimportstoarguments_multiple_kernels(fortran_writer): # The kernels are checked before the psy.gen, so they don't include the # modified suffix. expected = [ - ["subroutine kernel_with_use_code(ji, jj, istep, ssha, tmask, rdt, " - "magic)", + ["subroutine kernel_with_use_code_inlined_(ji, jj, istep, ssha, " + "tmask, rdt, magic)", "real(kind=go_wp), intent(in) :: rdt"], - ["subroutine kernel_with_use2_code(ji, jj, istep, ssha, tmask, cbfr," - " rdt)", + ["subroutine kernel_with_use2_code_inlined_(ji, jj, istep, ssha, " + "tmask, cbfr, rdt)", "real(kind=go_wp), intent(inout) :: cbfr\n real(kind=go_wp), " "intent(in) :: rdt"], - ["subroutine kernel_with_use_code(ji, jj, istep, ssha, tmask, rdt, " - "magic)", + ["subroutine kernel_with_use_code_inlined_(ji, jj, istep, ssha, " + "tmask, rdt, magic)", "real(kind=go_wp), intent(in) :: rdt\n real(kind=go_wp), " "intent(inout) :: magic"]] diff --git a/src/psyclone/tests/domain/gocean/transformations/gocean1p0_transformations_test.py b/src/psyclone/tests/domain/gocean/transformations/gocean1p0_transformations_test.py index fda730c02c..eaf0298341 100644 --- a/src/psyclone/tests/domain/gocean/transformations/gocean1p0_transformations_test.py +++ b/src/psyclone/tests/domain/gocean/transformations/gocean1p0_transformations_test.py @@ -1530,8 +1530,8 @@ def raise_gen_error(): monkeypatch.setattr(kern, "get_callees", raise_gen_error) with pytest.raises(TransformationError) as err: rtrans.apply(kern) - assert ("Failed to create PSyIR for kernel 'continuity_code'. Cannot " - "transform such a kernel." in str(err.value)) + assert ("Failed to create PSyIR for kernel 'continuity_code_inlined_'. " + "Cannot transform such a kernel." in str(err.value)) def test_accroutinetrans_with_routine(fortran_writer): diff --git a/src/psyclone/tests/domain/gocean/transformations/gocean_move_iteration_boundaries_inside_kernel_trans_test.py b/src/psyclone/tests/domain/gocean/transformations/gocean_move_iteration_boundaries_inside_kernel_trans_test.py index c7bdc20fd3..65755239fc 100644 --- a/src/psyclone/tests/domain/gocean/transformations/gocean_move_iteration_boundaries_inside_kernel_trans_test.py +++ b/src/psyclone/tests/domain/gocean/transformations/gocean_move_iteration_boundaries_inside_kernel_trans_test.py @@ -209,8 +209,8 @@ def test_go_move_iteration_boundaries_inside_kernel_two_kernels_apply_twice( ystop = cu_fld%internal%ystop do j = 1, SIZE(cu_fld%data, dim=2), 1 do i = 1, SIZE(cu_fld%data, dim=1), 1 - call compute_cu_code(i, j, cu_fld%data, p_fld%data, u_fld%data, xstart, \ -xstop, ystart, ystop) + call compute_cu_code_inlined_(i, j, cu_fld%data, p_fld%data, u_fld%data,\ + xstart, xstop, ystart, ystop) enddo enddo xstart_1 = 1 @@ -219,8 +219,8 @@ def test_go_move_iteration_boundaries_inside_kernel_two_kernels_apply_twice( ystop_1 = SIZE(uold_fld%data, dim=2) do j = 1, SIZE(uold_fld%data, dim=2), 1 do i = 1, SIZE(uold_fld%data, dim=1), 1 - call time_smooth_code(i, j, cu_fld%data, unew_fld%data, uold_fld%data, \ -xstart_1, xstop_1, ystart_1, ystop_1) + call time_smooth_code_inlined_(i, j, cu_fld%data, unew_fld%data, \ +uold_fld%data, xstart_1, xstop_1, ystart_1, ystop_1) enddo enddo diff --git a/src/psyclone/tests/domain/gocean/transformations/gocean_opencl_trans_test.py b/src/psyclone/tests/domain/gocean/transformations/gocean_opencl_trans_test.py index 72923b557b..08f24ef94e 100644 --- a/src/psyclone/tests/domain/gocean/transformations/gocean_opencl_trans_test.py +++ b/src/psyclone/tests/domain/gocean/transformations/gocean_opencl_trans_test.py @@ -39,6 +39,7 @@ ''' Module containing tests for the PSyclone GOcean OpenCL transformation.''' import os +import re import pytest from psyclone.configuration import Config @@ -263,7 +264,8 @@ def test_invoke_opencl_initialisation(kernel_outputdir, fortran_writer): if (first_time) then call psy_init() cmd_queues => get_cmd_queues() - kernel_compute_cu_code = get_kernel_by_name('compute_cu_code') + kernel_compute_cu_code_inlined_ = get_kernel_by_name(\ +'compute_cu_code_inlined_') call initialise_device_buffer(cu_fld) call initialise_device_buffer(p_fld) call initialise_device_buffer(u_fld) @@ -272,8 +274,9 @@ def test_invoke_opencl_initialisation(kernel_outputdir, fortran_writer): cu_fld_cl_mem = transfer(cu_fld%device_ptr, cu_fld_cl_mem) p_fld_cl_mem = transfer(p_fld%device_ptr, p_fld_cl_mem) u_fld_cl_mem = transfer(u_fld%device_ptr, u_fld_cl_mem) - call compute_cu_code_set_args(kernel_compute_cu_code, cu_fld_cl_mem, \ -p_fld_cl_mem, u_fld_cl_mem, xstart - 1, xstop - 1, ystart - 1, ystop - 1) + call compute_cu_code_inlined__set_args(kernel_compute_cu_code_inlined_, \ +cu_fld_cl_mem, p_fld_cl_mem, u_fld_cl_mem, xstart - 1, xstop - 1, ystart - 1, \ +ystop - 1) ! write data to the device''' assert expected in generated_code @@ -388,7 +391,8 @@ def test_invoke_opencl_initialisation_grid(kernel_outputdir): if (first_time) then call psy_init() cmd_queues => get_cmd_queues() - kernel_compute_kernel_code = get_kernel_by_name('compute_kernel_code') + kernel_compute_kernel_code_inlined_ = get_kernel_by_name(\ +'compute_kernel_code_inlined_') call initialise_device_buffer(out_fld) call initialise_device_buffer(in_out_fld) call initialise_device_buffer(in_fld) @@ -403,7 +407,8 @@ def test_invoke_opencl_initialisation_grid(kernel_outputdir): dx_cl_mem = transfer(dx%device_ptr, dx_cl_mem) gphiu_cl_mem = transfer(in_fld%grid%gphiu_device, \ gphiu_cl_mem) - call compute_kernel_code_set_args(kernel_compute_kernel_code, \ + call compute_kernel_code_inlined__set_args(\ +kernel_compute_kernel_code_inlined_, \ out_fld_cl_mem, in_out_fld_cl_mem, in_fld_cl_mem, dx_cl_mem, \ in_fld%grid%dx, gphiu_cl_mem, xstart - 1, xstop - 1, ystart - 1, \ ystop - 1) @@ -594,7 +599,7 @@ def test_psy_init_defaults(kernel_outputdir): if (.not.initialised) then initialised = .true. call ocl_env_init(1, ocl_device_num, .false., .false.) - kernel_names(1) = 'compute_cu_code' + kernel_names(1) = 'compute_cu_code_inlined_' call add_kernels(1, kernel_names) end if @@ -626,18 +631,17 @@ def test_psy_init_multiple_kernels(kernel_outputdir): generated_code = str(psy.gen) # Check that the kernel_names has enough space for all kernels - assert "character(len=30), dimension(2) :: kernel_names" in generated_code - - # The order doesn't matter as far as the two kernels are loaded - assert ("kernel_names(1) = 'kernel_with_use_code'" in generated_code or - "kernel_names(2) = 'kernel_with_use_code'" in generated_code) + assert "character(len=30), dimension(3) :: kernel_names" in generated_code - assert ("kernel_names(1) = 'kernel_with_use2_code'" in generated_code or - "kernel_names(2) = 'kernel_with_use2_code'" in generated_code) - assert "kernel_names(3)" not in generated_code + # The order doesn't matter as long as the three kernels are loaded + for name in ['kernel_with_use_code_inlined_', + 'kernel_with_use_code_inlined__1', + 'kernel_with_use2_code_inlined_']: + assert re.search( + rf"kernel_names\([1-3]\) = '{name}'", generated_code) # Check that add_kernels is provided with the total number of kernels - assert "call add_kernels(2, kernel_names)" in generated_code + assert "call add_kernels(3, kernel_names)" in generated_code assert GOceanOpenCLBuild(kernel_outputdir).code_compiles( psy, dependencies=["model_mod.f90"]) @@ -653,7 +657,7 @@ def test_psy_init_multiple_devices_per_node(kernel_outputdir, monkeypatch): trans = GOMoveIterationBoundariesInsideKernelTrans() mod_inline_trans = KernelModuleInlineTrans() for kernel in sched.coded_kernels(): - mod_inline_trans.apply(kernel) + mod_inline_trans.apply(kernel, update_all=True) trans.apply(kernel) # Test with a different configuration value for OCL_DEVICES_PER_NODE @@ -677,7 +681,7 @@ def test_psy_init_multiple_devices_per_node(kernel_outputdir, monkeypatch): initialised = .true. ocl_device_num = mod(get_rank() - 1, 2) + 1 call ocl_env_init(1, ocl_device_num, .false., .false.) - kernel_names(1) = 'compute_cu_code' + kernel_names(1) = 'compute_cu_code_inlined_' call add_kernels(1, kernel_names) end if @@ -749,7 +753,7 @@ def test_invoke_opencl_kernel_call(kernel_outputdir, monkeypatch, debug_mode): # the kernel expected += ''' ierr = clFinish(cmd_queues(1)) - call check_status('Errors before compute_cu_code launch', ierr)''' + call check_status('Errors before compute_cu_code_inlined_ launch', ierr)''' # Cast dl_esm_inf pointers to cl_mem handlers expected += ''' @@ -760,7 +764,7 @@ def test_invoke_opencl_kernel_call(kernel_outputdir, monkeypatch, debug_mode): # Call the set_args subroutine with the boundaries corrected for the # OpenCL 0-indexing expected += ''' - call compute_cu_code_set_args(kernel_compute_cu_code, \ + call compute_cu_code_inlined__set_args(kernel_compute_cu_code_inlined_, \ cu_fld_cl_mem, p_fld_cl_mem, u_fld_cl_mem, \ xstart - 1, xstop - 1, \ ystart - 1, ystop - 1) @@ -768,7 +772,7 @@ def test_invoke_opencl_kernel_call(kernel_outputdir, monkeypatch, debug_mode): expected += ''' ! Launch the kernel - ierr = clEnqueueNDRangeKernel(cmd_queues(1), kernel_compute_cu_code, \ + ierr = clEnqueueNDRangeKernel(cmd_queues(1), kernel_compute_cu_code_inlined_, \ 2, C_NULL_PTR, C_LOC(globalsize), C_LOC(localsize), 0, C_NULL_PTR, \ C_NULL_PTR)''' @@ -776,9 +780,9 @@ def test_invoke_opencl_kernel_call(kernel_outputdir, monkeypatch, debug_mode): # Check that there are no errors during the kernel launch or during # the execution of the kernel. expected += ''' - call check_status('compute_cu_code clEnqueueNDRangeKernel', ierr) + call check_status('compute_cu_code_inlined_ clEnqueueNDRangeKernel', ierr) ierr = clFinish(cmd_queues(1)) - call check_status('Errors during compute_cu_code', ierr)''' + call check_status('Errors during compute_cu_code_inlined_', ierr)''' assert expected in generated_code assert GOceanOpenCLBuild(kernel_outputdir).code_compiles(psy) @@ -924,7 +928,7 @@ def test_opencl_options_effects(): generated_code = str(psy.gen) assert "localsize = (/64, 1/)" in generated_code assert "ierr = clEnqueueNDRangeKernel(cmd_queues(1), " \ - "kernel_compute_cu_code, 2, C_NULL_PTR, C_LOC(globalsize), " \ + "kernel_compute_cu_code_inlined_, 2, C_NULL_PTR, C_LOC(globalsize), " \ "C_LOC(localsize), 0, C_NULL_PTR, C_NULL_PTR)" in generated_code assert "ierr = clFinish(cmd_queues(1))" in generated_code assert "ierr = clFinish(cmd_queues(2))" not in generated_code @@ -958,7 +962,7 @@ def test_opencl_options_effects(): otrans.apply(sched) generated_code = str(psy.gen) assert "ierr = clEnqueueNDRangeKernel(cmd_queues(2), " \ - "kernel_compute_cu_code, 2, C_NULL_PTR, C_LOC(globalsize), " \ + "kernel_compute_cu_code_inlined_, 2, C_NULL_PTR, C_LOC(globalsize), " \ "C_LOC(localsize), 0, C_NULL_PTR, C_NULL_PTR)" in generated_code assert " ierr = clFinish(cmd_queues(1))\n" \ " ierr = clFinish(cmd_queues(2))\n" in generated_code @@ -1046,7 +1050,7 @@ def test_set_kern_args(kernel_outputdir): generated_code = str(psy.gen) # Check we've only generated one set-args routine with arguments: # kernel object + kernel arguments + boundary values - assert generated_code.count("subroutine compute_cu_code_set_args(" + assert generated_code.count("subroutine compute_cu_code_inlined__set_args(" "kernel_obj, cu_fld, p_fld, u_fld, xstart, " "xstop, ystart, ystop)") == 1 # Declarations @@ -1066,32 +1070,32 @@ def test_set_kern_args(kernel_outputdir): assert expected in generated_code expected = '''\ ierr = clSetKernelArg(kernel_obj, 0, C_SIZEOF(cu_fld), C_LOC(cu_fld)) - call check_status('clSetKernelArg: arg 0 of compute_cu_code', ierr) + call check_status('clSetKernelArg: arg 0 of compute_cu_code_inlined_', ierr) ierr = clSetKernelArg(kernel_obj, 1, C_SIZEOF(p_fld), C_LOC(p_fld)) - call check_status('clSetKernelArg: arg 1 of compute_cu_code', ierr) + call check_status('clSetKernelArg: arg 1 of compute_cu_code_inlined_', ierr) ierr = clSetKernelArg(kernel_obj, 2, C_SIZEOF(u_fld), C_LOC(u_fld)) - call check_status('clSetKernelArg: arg 2 of compute_cu_code', ierr) + call check_status('clSetKernelArg: arg 2 of compute_cu_code_inlined_', ierr) ierr = clSetKernelArg(kernel_obj, 3, C_SIZEOF(xstart), C_LOC(xstart)) - call check_status('clSetKernelArg: arg 3 of compute_cu_code', ierr) + call check_status('clSetKernelArg: arg 3 of compute_cu_code_inlined_', ierr) ierr = clSetKernelArg(kernel_obj, 4, C_SIZEOF(xstop), C_LOC(xstop)) - call check_status('clSetKernelArg: arg 4 of compute_cu_code', ierr) + call check_status('clSetKernelArg: arg 4 of compute_cu_code_inlined_', ierr) ierr = clSetKernelArg(kernel_obj, 5, C_SIZEOF(ystart), C_LOC(ystart)) - call check_status('clSetKernelArg: arg 5 of compute_cu_code', ierr) + call check_status('clSetKernelArg: arg 5 of compute_cu_code_inlined_', ierr) ierr = clSetKernelArg(kernel_obj, 6, C_SIZEOF(ystop), C_LOC(ystop)) - call check_status('clSetKernelArg: arg 6 of compute_cu_code', ierr) + call check_status('clSetKernelArg: arg 6 of compute_cu_code_inlined_', ierr) - end subroutine compute_cu_code_set_args''' + end subroutine compute_cu_code_inlined__set_args''' assert expected in generated_code # The call to the set_args matches the expected kernel signature with # the boundary values converted to 0-indexing - assert ("call compute_cu_code_set_args(kernel_compute_cu_code, " + assert ("call compute_cu_code_inlined__set_args(kernel_compute_cu_code_inlined_, " "cu_fld_cl_mem, p_fld_cl_mem, u_fld_cl_mem, " "xstart - 1, xstop - 1, " "ystart - 1, ystop - 1)" in generated_code) # There is also only one version of the set_args for the second kernel - assert generated_code.count("subroutine time_smooth_code_set_args(" + assert generated_code.count("subroutine time_smooth_code_inlined__set_args(" "kernel_obj, cu_fld, unew_fld, uold_fld, " "xstart_1, xstop_1, ystart_1, ystop_1)") == 1 assert GOceanOpenCLBuild(kernel_outputdir).code_compiles(psy) @@ -1116,7 +1120,7 @@ def test_set_kern_args_real_grid_property(kernel_outputdir): generated_code = str(psy.gen) expected = '''\ - subroutine compute_kernel_code_set_args(kernel_obj, out_fld, in_out_fld, \ + subroutine compute_kernel_code_inlined__set_args(kernel_obj, out_fld, in_out_fld, \ in_fld, dx, dx_1, gphiu, xstart, xstop, ystart, ystop) use clfortran, only : clSetKernelArg use iso_c_binding, only : C_LOC, C_SIZEOF, c_intptr_t @@ -1155,7 +1159,7 @@ def test_set_kern_float_arg(kernel_outputdir): # This set_args has a name clash on xstop (one is a grid property and the # other a loop boundary). One of they should appear as 'xstop_1' expected = '''\ - subroutine bc_ssh_code_set_args(kernel_obj, a_scalar, ssh_fld, xstop, \ + subroutine bc_ssh_code_inlined__set_args(kernel_obj, a_scalar, ssh_fld, xstop, \ tmask, xstart, xstop_1, ystart, ystop) use clfortran, only : clSetKernelArg use iso_c_binding, only : C_LOC, C_SIZEOF, c_intptr_t @@ -1174,23 +1178,23 @@ def test_set_kern_float_arg(kernel_outputdir): assert expected in generated_code expected = '''\ ierr = clSetKernelArg(kernel_obj, 0, C_SIZEOF(a_scalar), C_LOC(a_scalar)) - call check_status('clSetKernelArg: arg 0 of bc_ssh_code', ierr) + call check_status('clSetKernelArg: arg 0 of bc_ssh_code_inlined_', ierr) ierr = clSetKernelArg(kernel_obj, 1, C_SIZEOF(ssh_fld), C_LOC(ssh_fld)) - call check_status('clSetKernelArg: arg 1 of bc_ssh_code', ierr) + call check_status('clSetKernelArg: arg 1 of bc_ssh_code_inlined_', ierr) ierr = clSetKernelArg(kernel_obj, 2, C_SIZEOF(xstop), C_LOC(xstop)) - call check_status('clSetKernelArg: arg 2 of bc_ssh_code', ierr) + call check_status('clSetKernelArg: arg 2 of bc_ssh_code_inlined_', ierr) ierr = clSetKernelArg(kernel_obj, 3, C_SIZEOF(tmask), C_LOC(tmask)) - call check_status('clSetKernelArg: arg 3 of bc_ssh_code', ierr) + call check_status('clSetKernelArg: arg 3 of bc_ssh_code_inlined_', ierr) ierr = clSetKernelArg(kernel_obj, 4, C_SIZEOF(xstart), C_LOC(xstart)) - call check_status('clSetKernelArg: arg 4 of bc_ssh_code', ierr) + call check_status('clSetKernelArg: arg 4 of bc_ssh_code_inlined_', ierr) ierr = clSetKernelArg(kernel_obj, 5, C_SIZEOF(xstop_1), C_LOC(xstop_1)) - call check_status('clSetKernelArg: arg 5 of bc_ssh_code', ierr) + call check_status('clSetKernelArg: arg 5 of bc_ssh_code_inlined_', ierr) ierr = clSetKernelArg(kernel_obj, 6, C_SIZEOF(ystart), C_LOC(ystart)) - call check_status('clSetKernelArg: arg 6 of bc_ssh_code', ierr) + call check_status('clSetKernelArg: arg 6 of bc_ssh_code_inlined_', ierr) ierr = clSetKernelArg(kernel_obj, 7, C_SIZEOF(ystop), C_LOC(ystop)) - call check_status('clSetKernelArg: arg 7 of bc_ssh_code', ierr) + call check_status('clSetKernelArg: arg 7 of bc_ssh_code_inlined_', ierr) - end subroutine bc_ssh_code_set_args''' + end subroutine bc_ssh_code_inlined__set_args''' assert expected in generated_code assert GOceanOpenCLBuild(kernel_outputdir).code_compiles(psy) diff --git a/src/psyclone/tests/domain/lfric/lfric_kern_test.py b/src/psyclone/tests/domain/lfric/lfric_kern_test.py index d827e86bb0..7cfeb284dc 100644 --- a/src/psyclone/tests/domain/lfric/lfric_kern_test.py +++ b/src/psyclone/tests/domain/lfric/lfric_kern_test.py @@ -308,6 +308,7 @@ def test_validate_kernel_code_arg(monkeypatch): kernel = LFRicKern() # Kernel name needs to be set when testing exceptions. kernel._name = "dummy" + read_access = ArgumentInterface(ArgumentInterface.Access.READ) real_scalar_symbol = DataSymbol( diff --git a/src/psyclone/tests/psyGen_test.py b/src/psyclone/tests/psyGen_test.py index e852b77d18..390a5d231d 100644 --- a/src/psyclone/tests/psyGen_test.py +++ b/src/psyclone/tests/psyGen_test.py @@ -772,7 +772,8 @@ def test_codedkern_node_str(): assert expected_output in out -def test_codedkern_module_inline_getter(): +@pytest.mark.parametrize("do_all", [True, False]) +def test_codedkern_module_inline_getter(do_all): ''' Check that the module_inline property of CodedKern. ''' # Use LFRic example with a repeated CodedKern _, invoke = get_invoke("4.6_multikernel_invokes.f90", api="lfric", idx=0) @@ -782,10 +783,10 @@ def test_codedkern_module_inline_getter(): assert ckerns[0].name == ckerns[1].name assert ckerns[0].module_inline is False mod_inline_trans = KernelModuleInlineTrans() - mod_inline_trans.apply(ckerns[0]) - # Module inlining one should have updated both. + mod_inline_trans.apply(ckerns[0], update_all=do_all) + # Module inlining one will only have updated both if `update_all` is True assert ckerns[0].module_inline is True - assert ckerns[1].module_inline is True + assert ckerns[1].module_inline is do_all def test_codedkern_lower_to_language_level(monkeypatch): diff --git a/src/psyclone/tests/psyir/transformations/inline_trans_test.py b/src/psyclone/tests/psyir/transformations/inline_trans_test.py index 63e03d88f6..d0724ed29f 100644 --- a/src/psyclone/tests/psyir/transformations/inline_trans_test.py +++ b/src/psyclone/tests/psyir/transformations/inline_trans_test.py @@ -1557,7 +1557,7 @@ def test_apply_raw_subroutine( if start: modinline_trans = KernelModuleInlineTrans() modinline_trans.apply(call) - assert "sub" in psyir.children[0].symbol_table + assert "sub_inlined_" in psyir.children[0].symbol_table inline_trans = InlineTrans() inline_trans.apply(call) output = fortran_writer(psyir) diff --git a/src/psyclone/tests/psyir/transformations/kernel_transformation_test.py b/src/psyclone/tests/psyir/transformations/kernel_transformation_test.py index 1f2c647fd6..650a16e357 100644 --- a/src/psyclone/tests/psyir/transformations/kernel_transformation_test.py +++ b/src/psyclone/tests/psyir/transformations/kernel_transformation_test.py @@ -95,12 +95,12 @@ def test_transform_kern_with_interface(tmp_path, fortran_writer): invoke.setup_psy_layer_symbols() contents = fortran_writer(sched.ancestor(Container)) # Check that the interface name has been updated. - assert "interface mixed_code" in contents - assert ("module procedure :: mixed_code_32, mixed_code_64" - in contents) - # Check that the subroutines themselves haven't been renamed. - assert "subroutine mixed_code_32" in contents - assert "subroutine mixed_code_64" in contents + assert "interface mixed_code_inlined_" in contents + assert ("module procedure :: mixed_code_32_inlined_, " + "mixed_code_64_inlined_" in contents) + # Check that the subroutines themselves havet been renamed. + assert "subroutine mixed_code_32_inlined_" in contents + assert "subroutine mixed_code_64_inlined_" in contents # But they have been transformed. assert ('''real*4, dimension(op_ncell_3d,ndf_w0,ndf_w0), intent(in) :: op @@ -337,16 +337,17 @@ def test_1kern_trans(tmp_path): sched = invoke.schedule kernels = sched.coded_kernels() kern = kernels[1] - # We have to module-inline the kernel before we can transform it and that - # will affect all calls to that kernel in the invoke. + # We have to module-inline the kernel before we can transform it KernelModuleInlineTrans().apply(kern) rtrans = ACCRoutineTrans() rtrans.apply(kern) # Generate the code code = str(psy.gen).lower() - assert 'use testkern_mod' not in code - assert code.count("call testkern_code(") == 2 - assert "private :: testkern_code" in code + # The original import is left unchanged + assert 'use testkern_mod' in code + # The call is now to the inlined version + assert code.count("call testkern_code_inlined_(") == 1 + assert "private :: testkern_code_inlined_" in code assert LFRicBuild(tmp_path).code_compiles(psy) @@ -368,8 +369,9 @@ def test_2kern_trans(tmp_path): code = str(psy.gen).lower() # Check that the old module re-naming no longer happens. assert not re.match('use testkern_any_space_2(.+?)_mod', code) - assert "use testkern_any_space_2_mod, only" not in code - assert "call testkern_any_space_2_code(" in code + # use statements are unchanged + assert "use testkern_any_space_2_mod, only" in code + assert "call testkern_any_space_2_code_inlined_(" in code assert LFRicBuild(tmp_path).code_compiles(psy)