Beispiel #1
0
def example_psyir_nary():
    '''Utility function that creates a PSyIR tree containing a nary MIN
    intrinsic operator and returns the operator.

    :returns: PSyIR MIN operator instance.
    :rtype: :py:class:`psyclone.psyGen.NaryOperation`

    '''
    symbol_table = SymbolTable()
    arg1 = symbol_table.new_symbol("arg",
                                   symbol_type=DataSymbol,
                                   datatype=REAL_TYPE,
                                   interface=ArgumentInterface(
                                       ArgumentInterface.Access.READWRITE))
    arg2 = symbol_table.new_symbol("arg",
                                   symbol_type=DataSymbol,
                                   datatype=REAL_TYPE,
                                   interface=ArgumentInterface(
                                       ArgumentInterface.Access.READWRITE))
    arg3 = symbol_table.new_symbol("arg",
                                   symbol_type=DataSymbol,
                                   datatype=REAL_TYPE,
                                   interface=ArgumentInterface(
                                       ArgumentInterface.Access.READWRITE))
    arg4 = symbol_table.new_symbol(symbol_type=DataSymbol, datatype=REAL_TYPE)
    symbol_table.specify_argument_list([arg1, arg2, arg3])
    var1 = Reference(arg1)
    var2 = Reference(arg2)
    var3 = Reference(arg3)
    var4 = Reference(arg4)
    oper = NaryOperation.Operator.MIN
    operation = NaryOperation.create(oper, [var1, var2, var3])
    assign = Assignment.create(var4, operation)
    _ = KernelSchedule.create("min_example", symbol_table, [assign])
    return operation
Beispiel #2
0
def test_argumentinterface_access_values(access):
    '''Check that all the ArgumentInterface access values are supported.

    '''
    argument_interface = ArgumentInterface()
    argument_interface.access = access
    assert argument_interface.access == access
Beispiel #3
0
def example_psyir(create_expression):
    '''Utility function that creates a PSyIR tree containing a SIGN
    intrinsic operator and returns the operator.

    :param function create_expresssion: function used to create the \
        content of the first argument of the SIGN operator.

    :returns: PSyIR SIGN operator instance.
    :rtype: :py:class:`psyclone.psyGen.BinaryOperation`

    '''
    symbol_table = SymbolTable()
    arg1 = symbol_table.new_symbol("arg",
                                   symbol_type=DataSymbol,
                                   datatype=REAL_TYPE,
                                   interface=ArgumentInterface(
                                       ArgumentInterface.Access.READWRITE))
    arg2 = symbol_table.new_symbol("arg",
                                   symbol_type=DataSymbol,
                                   datatype=REAL_TYPE,
                                   interface=ArgumentInterface(
                                       ArgumentInterface.Access.READWRITE))
    arg3 = symbol_table.new_symbol(symbol_type=DataSymbol, datatype=REAL_TYPE)
    symbol_table.specify_argument_list([arg1, arg2])
    var1 = Reference(arg1)
    var2 = Reference(arg2)
    var3 = Reference(arg3)
    oper = BinaryOperation.Operator.SIGN
    operation = BinaryOperation.create(oper, create_expression(var1), var2)
    assign = Assignment.create(var3, operation)
    _ = KernelSchedule.create("sign_example", symbol_table, [assign])
    return operation
Beispiel #4
0
def test_symbol_interface_setter():
    '''Test that the Symbol interface setter behaves as expected,
    including raising an exception if the input is of the wrong
    type. Also use this to test the is_local, is_global and
    is_argument and is_unresolved properties.

    '''
    symbol = Symbol('sym1')
    assert symbol.is_local
    assert not symbol.is_global
    assert not symbol.is_argument
    assert not symbol.is_unresolved

    symbol.interface = GlobalInterface(ContainerSymbol("my_mod"))
    assert not symbol.is_local
    assert symbol.is_global
    assert not symbol.is_argument
    assert not symbol.is_unresolved

    symbol.interface = ArgumentInterface()
    assert not symbol.is_local
    assert not symbol.is_global
    assert symbol.is_argument
    assert not symbol.is_unresolved

    symbol.interface = UnresolvedInterface()
    assert not symbol.is_local
    assert not symbol.is_global
    assert not symbol.is_argument
    assert symbol.is_unresolved

    with pytest.raises(TypeError) as info:
        symbol.interface = "hello"
    assert ("The interface to a Symbol must be a SymbolInterface but got "
            "'str'" in str(info.value))
Beispiel #5
0
def test_oclw_gen_declaration():
    '''Check the OpenCLWriter class gen_declaration method produces
    the expected declarations.

    '''
    oclwriter = OpenCLWriter()

    # Basic entry - Scalar are passed by value and don't have additional
    # qualifiers.
    symbol = DataSymbol("dummy1", INTEGER_TYPE)
    result = oclwriter.gen_declaration(symbol)
    assert result == "int dummy1"

    # Array argument has a memory qualifier (only __global for now)
    array_type = ArrayType(INTEGER_TYPE, [2, ArrayType.Extent.ATTRIBUTE, 2])
    symbol = DataSymbol("dummy2", array_type)
    result = oclwriter.gen_declaration(symbol)
    assert result == "__global int * restrict dummy2"

    # Array with unknown intent
    array_type = ArrayType(INTEGER_TYPE, [2, ArrayType.Extent.ATTRIBUTE, 2])
    symbol = DataSymbol("dummy2",
                        array_type,
                        interface=ArgumentInterface(
                            ArgumentInterface.Access.UNKNOWN))
    result = oclwriter.gen_declaration(symbol)
    assert result == "__global int * restrict dummy2"
Beispiel #6
0
def test_datasymbol_copy_properties():
    '''Test that the DataSymbol copy_properties method works as expected.'''
    array_type = ArrayType(REAL_SINGLE_TYPE, [1, 2])
    symbol = DataSymbol("myname",
                        array_type,
                        constant_value=None,
                        interface=ArgumentInterface(
                            ArgumentInterface.Access.READWRITE))

    # Check an exception is raised if an incorrect argument is passed in
    with pytest.raises(TypeError) as excinfo:
        symbol.copy_properties(None)
    assert ("Argument should be of type 'DataSymbol' but found 'NoneType'."
            "") in str(excinfo.value)

    new_symbol = DataSymbol("other_name",
                            INTEGER_SINGLE_TYPE,
                            constant_value=7)

    symbol.copy_properties(new_symbol)

    assert symbol.name == "myname"
    assert symbol.datatype.intrinsic == ScalarType.Intrinsic.INTEGER
    assert symbol.datatype.precision == ScalarType.Precision.SINGLE
    assert symbol.is_local
    assert isinstance(symbol.constant_value, Literal)
    assert symbol.constant_value.value == "7"
    assert (
        symbol.constant_value.datatype.intrinsic == symbol.datatype.intrinsic)
    assert (
        symbol.constant_value.datatype.precision == symbol.datatype.precision)
Beispiel #7
0
def example_psyir(create_expression):
    '''Utility function that creates a PSyIR tree containing an ABS
    intrinsic operator and returns the operator.

    :param function create_expresssion: function used to create the \
        content of the ABS operator.

    :returns: PSyIR ABS operator instance.
    :rtype: :py:class:`psyclone.psyGen.UnaryOperation`

    '''
    symbol_table = SymbolTable()
    name1 = symbol_table.new_symbol_name("arg")
    arg1 = DataSymbol(name1,
                      REAL_TYPE,
                      interface=ArgumentInterface(
                          ArgumentInterface.Access.READWRITE))
    symbol_table.add(arg1)
    name2 = symbol_table.new_symbol_name()
    local = DataSymbol(name2, REAL_TYPE)
    symbol_table.add(local)
    symbol_table.specify_argument_list([arg1])
    var1 = Reference(arg1)
    var2 = Reference(local)
    oper = UnaryOperation.Operator.ABS
    operation = UnaryOperation.create(oper, create_expression(var1))
    assign = Assignment.create(var2, operation)
    _ = KernelSchedule.create("abs_example", symbol_table, [assign])
    return operation
Beispiel #8
0
def test_generic_scalars(data_type, symbol, intrinsic, precision):
    '''Test the generated generic scalar datatypes and symbols are created
    correctly.

    '''
    # datatype
    lfric_datatype = data_type()
    assert lfric_datatype.intrinsic == intrinsic
    assert lfric_datatype.precision is precision
    # precision can be set explicitly
    lfric_datatype = data_type(precision=4)
    assert lfric_datatype.precision == 4
    # symbol
    lfric_symbol = symbol("symbol")
    assert lfric_symbol.name == "symbol"
    assert isinstance(lfric_symbol.interface, LocalInterface)
    assert isinstance(lfric_symbol.datatype, data_type)
    lfric_symbol = symbol("symbol",
                          interface=ArgumentInterface(
                              ArgumentInterface.Access.READ))
    assert isinstance(lfric_symbol.interface, ArgumentInterface)
    assert lfric_symbol.interface.access == ArgumentInterface.Access.READ
    # precision can be set explicitly
    lfric_symbol = symbol("symbol", precision=4)
    assert lfric_symbol.datatype.precision == 4
Beispiel #9
0
    def scalar(self, scalar_arg, var_accesses=None):
        '''Create an LFRic scalar argument and add it to the symbol table and
        argument list.

        :param scalar_arg: the scalar to add.
        :type scalar_arg: :py:class:`psyclone.dynamo0p3.DynKernelArgument`
        :param var_accesses: an unused optional argument that stores \
            information about variable accesses.
        :type var_accesses: :\
            py:class:`psyclone.core.access_info.VariablesAccessInfo`

        :raises NotImplementedError: if the datatype of the scalar is \
            not supported.

        '''
        mapping = {
            "integer": lfric_psyir.LfricIntegerScalarDataSymbol,
            "real": lfric_psyir.LfricRealScalarDataSymbol,
            "logical": lfric_psyir.LfricLogicalScalarDataSymbol
        }
        try:
            symbol = self._symbol_table.symbol_from_tag(
                scalar_arg.name,
                symbol_type=mapping[scalar_arg.intrinsic_type],
                interface=ArgumentInterface(INTENT_MAPPING[scalar_arg.intent]))
        except KeyError as info:
            message = (
                "scalar of type '{0}' not implemented in KernelInterface "
                "class.".format(scalar_arg.intrinsic_type))
            six.raise_from(NotImplementedError(message), info)
        self._arglist.append(symbol)
Beispiel #10
0
def test_cw_gen_declaration():
    '''Check the CWriter class gen_declaration method produces
    the expected declarations.

    '''
    cwriter = CWriter()

    # Basic entries
    symbol = DataSymbol("dummy1", INTEGER_TYPE)
    result = cwriter.gen_declaration(symbol)
    assert result == "int dummy1"

    symbol = DataSymbol("dummy1", CHARACTER_TYPE)
    result = cwriter.gen_declaration(symbol)
    assert result == "char dummy1"

    symbol = DataSymbol("dummy1", BOOLEAN_TYPE)
    result = cwriter.gen_declaration(symbol)
    assert result == "bool dummy1"

    # Array argument
    array_type = ArrayType(REAL_TYPE, [2, ArrayType.Extent.ATTRIBUTE, 2])
    symbol = DataSymbol("dummy2",
                        array_type,
                        interface=ArgumentInterface(
                            ArgumentInterface.Access.READ))
    result = cwriter.gen_declaration(symbol)
    assert result == "double * restrict dummy2"

    # Array with unknown access
    array_type = ArrayType(INTEGER_TYPE, [2, ArrayType.Extent.ATTRIBUTE, 2])
    symbol = DataSymbol("dummy2",
                        array_type,
                        interface=ArgumentInterface(
                            ArgumentInterface.Access.UNKNOWN))
    result = cwriter.gen_declaration(symbol)
    assert result == "int * restrict dummy2"

    # Check invalid datatype produces and error
    symbol._datatype = "invalid"
    with pytest.raises(NotImplementedError) as error:
        _ = cwriter.gen_declaration(symbol)
    assert "Could not generate the C definition for the variable 'dummy2', " \
        "type 'invalid' is currently not supported." in str(error.value)
Beispiel #11
0
def test_datasymbol_copy():
    '''Test that the DataSymbol copy method produces a faithful separate copy
    of the original symbol.

    '''
    array_type = ArrayType(REAL_SINGLE_TYPE, [1, 2])
    symbol = DataSymbol("myname",
                        array_type,
                        constant_value=None,
                        interface=ArgumentInterface(
                            ArgumentInterface.Access.READWRITE))
    new_symbol = symbol.copy()

    # Check the new symbol has the same properties as the original
    assert symbol.name == new_symbol.name
    assert symbol.datatype == new_symbol.datatype
    assert symbol.shape == new_symbol.shape
    assert symbol.constant_value == new_symbol.constant_value
    assert symbol.interface == new_symbol.interface

    # Change the properties of the new symbol and check the original
    # is not affected. Can't check constant_value yet as we have a
    # shape value
    new_symbol._name = "new"
    new_symbol.datatype = ArrayType(
        ScalarType(ScalarType.Intrinsic.INTEGER, ScalarType.Precision.DOUBLE),
        [3, 4])
    new_symbol._interface = LocalInterface()

    assert symbol.name == "myname"
    assert symbol.datatype.intrinsic == ScalarType.Intrinsic.REAL
    assert symbol.datatype.precision == ScalarType.Precision.SINGLE
    assert len(symbol.shape) == 2
    assert isinstance(symbol.shape[0], Literal)
    assert symbol.shape[0].value == "1"
    assert symbol.shape[0].datatype.intrinsic == ScalarType.Intrinsic.INTEGER
    assert symbol.shape[0].datatype.precision == ScalarType.Precision.UNDEFINED
    assert isinstance(symbol.shape[1], Literal)
    assert symbol.shape[1].value == "2"
    assert symbol.shape[1].datatype.intrinsic == ScalarType.Intrinsic.INTEGER
    assert symbol.shape[1].datatype.precision == ScalarType.Precision.UNDEFINED
    assert not symbol.constant_value

    # Now check constant_value
    new_symbol.constant_value = 3

    assert isinstance(symbol.shape[0], Literal)
    assert symbol.shape[0].value == "1"
    assert symbol.shape[0].datatype.intrinsic == ScalarType.Intrinsic.INTEGER
    assert symbol.shape[0].datatype.precision == ScalarType.Precision.UNDEFINED
    assert isinstance(symbol.shape[1], Literal)
    assert symbol.shape[1].value == "2"
    assert symbol.shape[1].datatype.intrinsic == ScalarType.Intrinsic.INTEGER
    assert symbol.shape[1].datatype.precision == ScalarType.Precision.UNDEFINED
    assert not symbol.constant_value
Beispiel #12
0
def test_datasymbol_initialisation():
    '''Test that a DataSymbol instance can be created when valid arguments are
    given, otherwise raise relevant exceptions.'''

    # Test with valid arguments
    assert isinstance(DataSymbol('a', REAL_SINGLE_TYPE), DataSymbol)
    assert isinstance(DataSymbol('a', REAL_DOUBLE_TYPE), DataSymbol)
    assert isinstance(DataSymbol('a', REAL4_TYPE), DataSymbol)
    kind = DataSymbol('r_def', INTEGER_SINGLE_TYPE)
    real_kind_type = ScalarType(ScalarType.Intrinsic.REAL, kind)
    assert isinstance(DataSymbol('a', real_kind_type),
                      DataSymbol)
    # real constants are not currently supported
    assert isinstance(DataSymbol('a', INTEGER_SINGLE_TYPE), DataSymbol)
    assert isinstance(DataSymbol('a', INTEGER_DOUBLE_TYPE, constant_value=0),
                      DataSymbol)
    assert isinstance(DataSymbol('a', INTEGER4_TYPE),
                      DataSymbol)

    assert isinstance(DataSymbol('a', CHARACTER_TYPE), DataSymbol)
    assert isinstance(DataSymbol('a', CHARACTER_TYPE,
                                 constant_value="hello"), DataSymbol)
    assert isinstance(DataSymbol('a', BOOLEAN_TYPE), DataSymbol)
    assert isinstance(DataSymbol('a', BOOLEAN_TYPE,
                                 constant_value=False),
                      DataSymbol)
    array_type = ArrayType(REAL_SINGLE_TYPE, [ArrayType.Extent.ATTRIBUTE])
    assert isinstance(DataSymbol('a', array_type), DataSymbol)

    array_type = ArrayType(REAL_SINGLE_TYPE, [3])
    assert isinstance(DataSymbol('a', array_type), DataSymbol)
    array_type = ArrayType(REAL_SINGLE_TYPE, [3, ArrayType.Extent.ATTRIBUTE])
    assert isinstance(DataSymbol('a', array_type), DataSymbol)
    assert isinstance(DataSymbol('a', REAL_SINGLE_TYPE), DataSymbol)
    assert isinstance(DataSymbol('a', REAL8_TYPE), DataSymbol)
    dim = DataSymbol('dim', INTEGER_SINGLE_TYPE,
                     interface=UnresolvedInterface())
    array_type = ArrayType(REAL_SINGLE_TYPE, [Reference(dim)])
    assert isinstance(DataSymbol('a', array_type), DataSymbol)
    array_type = ArrayType(REAL_SINGLE_TYPE,
                           [3, Reference(dim), ArrayType.Extent.ATTRIBUTE])
    assert isinstance(DataSymbol('a', array_type), DataSymbol)
    assert isinstance(
        DataSymbol('a', REAL_SINGLE_TYPE,
                   interface=ArgumentInterface(
                       ArgumentInterface.Access.READWRITE)), DataSymbol)
    assert isinstance(
        DataSymbol('a', REAL_SINGLE_TYPE,
                   visibility=Symbol.Visibility.PRIVATE), DataSymbol)
    assert isinstance(DataSymbol('field', DataTypeSymbol("field_type",
                                                         DeferredType())),
                      DataSymbol)
Beispiel #13
0
def test_datasymbol_constant_value_setter_invalid():
    '''Test that a DataSymbol constant value setter raises the appropriate
    error if an invalid value and/or datatype are given.'''

    # Test with invalid constant values
    sym = DataSymbol('a', DeferredType())
    with pytest.raises(ValueError) as error:
        sym.constant_value = 1.0
    assert ("Error setting constant value for symbol 'a'. A DataSymbol with "
            "a constant value must be a scalar or an array but found "
            "'DeferredType'." in str(error.value))

    # Test with invalid constant expressions
    ct_expr = Return()
    with pytest.raises(ValueError) as error:
        _ = DataSymbol('a', INTEGER_SINGLE_TYPE, constant_value=ct_expr)
    assert "Error setting constant value for symbol 'a'. PSyIR static " \
        "expressions can only contain PSyIR literal, operation or reference" \
        " nodes but found:" in str(error.value)

    with pytest.raises(ValueError) as error:
        DataSymbol('a',
                   INTEGER_SINGLE_TYPE,
                   interface=ArgumentInterface(),
                   constant_value=9)
    assert ("Error setting constant value for symbol 'a'. A DataSymbol with "
            "an ArgumentInterface can not have a constant value."
            in str(error.value))

    with pytest.raises(ValueError) as error:
        DataSymbol('a', INTEGER_SINGLE_TYPE, constant_value=9.81)
    assert ("Error setting constant value for symbol 'a'. This DataSymbol "
            "instance datatype is 'Scalar<INTEGER, SINGLE>' which "
            "means the constant value is expected to be") in str(error.value)
    assert "'int'>' but found " in str(error.value)
    assert "'float'>'." in str(error.value)

    with pytest.raises(ValueError) as error:
        DataSymbol('a', CHARACTER_TYPE, constant_value=42)
    assert ("Error setting constant value for symbol 'a'. This DataSymbol "
            "instance datatype is 'Scalar<CHARACTER, UNDEFINED>' which "
            "means the constant value is expected to be") in str(error.value)
    assert "'str'>' but found " in str(error.value)
    assert "'int'>'." in str(error.value)

    with pytest.raises(ValueError) as error:
        DataSymbol('a', BOOLEAN_TYPE, constant_value="hello")
    assert ("Error setting constant value for symbol 'a'. This DataSymbol "
            "instance datatype is 'Scalar<BOOLEAN, UNDEFINED>' which "
            "means the constant value is expected to be") in str(error.value)
    assert "'bool'>' but found " in str(error.value)
    assert "'str'>'." in str(error.value)
Beispiel #14
0
    def operator(self, arg, var_accesses=None):
        '''Create an LFRic operator argument and an ncells argument and add
        them to the symbol table and argument list. Also declare the
        associated 'fs_from', 'fs_to' symbols if they have not already
        been declared so that they can be used to dimension the
        operator symbol (as well as ncells).

        :param arg: the operator to add.
        :type arg: :py:class:`psyclone.dynamo0p3.DynKernelArgument`
        :param var_accesses: an unused optional argument that stores \
            information about variable accesses.
        :type var_accesses: :\
            py:class:`psyclone.core.access_info.VariablesAccessInfo`

        :raises NotImplementedError: if the datatype of the field is \
            not supported.

        '''
        fs_from_name = arg.function_space_from.orig_name
        ndf_symbol_from = self._symbol_table.symbol_from_tag(
            "ndf_{0}".format(fs_from_name),
            fs=fs_from_name,
            symbol_type=lfric_psyir.NumberOfDofsDataSymbol,
            interface=self._read_access)
        fs_to_name = arg.function_space_to.orig_name
        ndf_symbol_to = self._symbol_table.symbol_from_tag(
            "ndf_{0}".format(fs_to_name),
            fs=fs_to_name,
            symbol_type=lfric_psyir.NumberOfDofsDataSymbol,
            interface=self._read_access)

        ncells = lfric_psyir.NumberOfCellsDataSymbol(
            "ncell_3d", interface=self._read_access)
        self._symbol_table.add(ncells)
        self._arglist.append(ncells)

        op_arg_symbol = self._symbol_table.symbol_from_tag(
            arg.name,
            symbol_type=lfric_psyir.OperatorDataSymbol,
            dims=[
                Reference(ndf_symbol_from),
                Reference(ndf_symbol_to),
                Reference(ncells)
            ],
            fs_from=fs_from_name,
            fs_to=fs_to_name,
            interface=ArgumentInterface(INTENT_MAPPING[arg.intent]))
        self._arglist.append(op_arg_symbol)
Beispiel #15
0
def test_specific_scalar_symbols(symbol, generic_symbol, attribute_map):
    '''Test the generated specific scalar symbols are
    created correctly.

    '''
    args = ["symbol"] + list(attribute_map.values())
    lfric_symbol = symbol(*args)
    assert isinstance(lfric_symbol, generic_symbol)
    assert lfric_symbol.name == "symbol"
    assert isinstance(lfric_symbol.interface, LocalInterface)
    for attribute in attribute_map:
        assert getattr(lfric_symbol, attribute) == attribute_map[attribute]
    lfric_symbol = symbol(*args,
                          interface=ArgumentInterface(
                              ArgumentInterface.Access.READ))
    assert isinstance(lfric_symbol.interface, ArgumentInterface)
    assert lfric_symbol.interface.access == ArgumentInterface.Access.READ
Beispiel #16
0
def test_field_vector(monkeypatch):
    '''Test the KernelInterface class field_vector method. We want to
    check that the correct symbol is referenced for the dimension of
    the vector field symbols so the simplest solution is to use one of
    the Fortran test examples. Also check that the expected exception
    is raised if the type of the vector field is not supported.

    '''
    kernel_interface = KernelInterface(None)
    _, invoke_info = parse(os.path.join(BASE_PATH, "8_vector_field.f90"),
                           api="dynamo0.3")
    psy = PSyFactory("dynamo0.3", distributed_memory=False).create(invoke_info)
    schedule = psy.invokes.invoke_list[0].schedule
    kernel = schedule[0].loop_body[0]
    vector_arg = kernel.args[1]
    kernel_interface.field_vector(vector_arg)

    # undf symbol declared
    undf_tag = "undf_{0}".format(vector_arg.function_space.orig_name)
    undf_symbol = kernel_interface._symbol_table.lookup(undf_tag)
    assert isinstance(undf_symbol, lfric_psyir.NumberOfUniqueDofsDataSymbol)
    assert isinstance(undf_symbol.interface, ArgumentInterface)
    assert (
        undf_symbol.interface.access == kernel_interface._read_access.access)

    # vector fields declared, added to argument list, correct function
    # space specified and dimensioned correctly
    for idx in range(vector_arg.vector_size):
        tag = "{0}_v{1}".format(vector_arg.name, idx)
        symbol = kernel_interface._symbol_table.lookup(tag)
        assert isinstance(symbol, lfric_psyir.RealVectorFieldDataDataSymbol)
        assert isinstance(symbol.interface, ArgumentInterface)
        assert (symbol.interface.access == ArgumentInterface(
            INTENT_MAPPING[vector_arg.intent]).access)
        assert kernel_interface._arglist[idx - 3] is symbol
        assert symbol.fs == vector_arg.function_space.orig_name
        assert len(symbol.shape) == 1
        assert isinstance(symbol.shape[0], Reference)
        assert symbol.shape[0].symbol is undf_symbol

    # Force an exception by setting the intrinsic type to an unsupported value
    monkeypatch.setattr(vector_arg, "_intrinsic_type", "unsupported")
    with pytest.raises(NotImplementedError) as info:
        kernel_interface.field_vector(vector_arg)
    assert ("kernel interface does not support a vector field of type "
            "'unsupported'." in str(info.value))
Beispiel #17
0
    def field_vector(self, argvect, var_accesses=None):
        '''Create LFRic field vector arguments and add them to the symbol
        table and argument list. Also declare the associated "undf"
        symbol if it has not already been declared so that it can be
        used to dimension the field vector arguments.

        :param argvect: the field vector to add.
        :type argvect: :py:class:`psyclone.dynamo0p3.DynKernelArgument`
        :param var_accesses: an unused optional argument that stores \
            information about variable accesses.
        :type var_accesses: :\
            py:class:`psyclone.core.access_info.VariablesAccessInfo`

        :raises NotImplementedError: if the datatype of the vector \
            field is not supported.

        '''
        fs_name = argvect.function_space.orig_name
        undf_symbol = self._symbol_table.symbol_from_tag(
            "undf_{0}".format(fs_name),
            fs=fs_name,
            symbol_type=lfric_psyir.NumberOfUniqueDofsDataSymbol,
            interface=self._read_access)

        interface = ArgumentInterface(INTENT_MAPPING[argvect.intent])
        try:
            field_class = self.vector_field_mapping[argvect.intrinsic_type]
        except KeyError as info:
            message = ("kernel interface does not support a vector field of "
                       "type '{0}'.".format(argvect.intrinsic_type))
            six.raise_from(NotImplementedError(message), info)
        for idx in range(argvect.vector_size):
            tag = "{0}_v{1}".format(argvect.name, idx)
            field_data_symbol = self._symbol_table.symbol_from_tag(
                tag,
                symbol_type=field_class,
                dims=[Reference(undf_symbol)],
                fs=fs_name,
                interface=interface)
            self._arglist.append(field_data_symbol)
Beispiel #18
0
def test_arrays(data_type, symbol, scalar_type, dims, attribute_map):
    '''Test the generated array datatypes and datasymbols are created
    correctly. This includes field datatypes and symbols which are
    kept as a separate list in psyir.py

    '''
    # Datatype creation
    lfric_datatype = data_type(dims)
    assert isinstance(lfric_datatype, ArrayType)
    assert isinstance(lfric_datatype._datatype, scalar_type)
    for idx, dim in enumerate(lfric_datatype.shape):
        if isinstance(dim, Literal):
            assert dim.value == str(dims[idx])
        elif isinstance(dim, Reference):
            assert dim is dims[idx]
            assert dim.symbol is dims[idx].symbol
        else:
            assert False, "unexpected type of dimension found"
    # Wrong number of dims
    with pytest.raises(TypeError) as info:
        _ = data_type([])
    assert ("{0} expected the number of supplied dimensions to be {1} but "
            "found 0.".format(type(lfric_datatype).__name__, len(dims))
            in str(info.value))
    # Datasymbol creation
    args = list(attribute_map.values())
    lfric_symbol = symbol("symbol", dims, *args)
    assert isinstance(lfric_symbol, DataSymbol)
    assert lfric_symbol.name == "symbol"
    assert isinstance(lfric_symbol.interface, LocalInterface)
    assert isinstance(lfric_symbol.datatype, data_type)
    lfric_symbol = symbol("symbol",
                          dims,
                          *args,
                          interface=ArgumentInterface(
                              ArgumentInterface.Access.READ))
    assert isinstance(lfric_symbol.interface, ArgumentInterface)
    assert lfric_symbol.interface.access == ArgumentInterface.Access.READ
Beispiel #19
0
def test_argumentinterface_init():
    '''Check that the ArgumentInterface can be created successfully and
    has the expected values. Also checks the access property and that
    an exception is raised if the supplied access value is the wrong
    type.

    '''
    argument_interface = ArgumentInterface()
    assert argument_interface._access == ArgumentInterface.Access.UNKNOWN
    assert argument_interface.access == argument_interface._access
    assert argument_interface._pass_by_value is False

    argument_interface = ArgumentInterface(ArgumentInterface.Access.READ)
    assert argument_interface._access == ArgumentInterface.Access.READ
    assert argument_interface.access == argument_interface._access

    with pytest.raises(TypeError) as info:
        _ = ArgumentInterface("hello")
    assert ("SymbolInterface.access must be an 'ArgumentInterface.Access' but "
            "got 'str'." in str(info.value))
    with pytest.raises(TypeError) as info:
        argument_interface.access = "hello"
    assert ("SymbolInterface.access must be an 'ArgumentInterface.Access' but "
            "got 'str'." in str(info.value))
Beispiel #20
0
Once you have psyclone installed, this script may be run by doing:

>>> python create.py

This should output a Fortran representation of the LFRic-PSyIR.

'''
# pylint: disable=no-member
from __future__ import print_function
from psyclone.psyir.nodes import Call, Reference, Container, KernelSchedule
from psyclone.psyir.symbols import RoutineSymbol, SymbolTable, \
    ArgumentInterface
from psyclone.domain.lfric import psyir as lfric_psyir
from psyclone.psyir.backend.fortran import FortranWriter

READ_ARG = ArgumentInterface(ArgumentInterface.Access.READ)

# Add LFRic precision symbols and the module in which they are
# contained to the symbol table
SYMBOL_TABLE = SymbolTable()
for symbol in [
        lfric_psyir.I_DEF, lfric_psyir.R_DEF, lfric_psyir.CONSTANTS_MOD
]:
    SYMBOL_TABLE.add(symbol)

# Create LFRic ndf and undf symbols and add them to the symbol table
NDF_W3 = lfric_psyir.NumberOfDofsDataSymbol("ndf_w3", "w3", interface=READ_ARG)
UNDF_W3 = lfric_psyir.NumberOfUniqueDofsDataSymbol("undf_w3",
                                                   "w3",
                                                   interface=READ_ARG)
for symbol in [NDF_W3, UNDF_W3]:
Beispiel #21
0
def test_symtab_implementation_for_opencl():
    ''' Tests that the GOcean specialised Symbol Table implements the
    abstract properties needed to generate OpenCL.
    '''
    kschedule = GOKernelSchedule('test')

    # Test symbol table without any kernel argument
    with pytest.raises(GenerationError) as err:
        _ = kschedule.symbol_table.iteration_indices
    assert ("GOcean 1.0 API kernels should always have at least two "
            "arguments representing the iteration indices but the Symbol "
            "Table for kernel 'test' has only 0 argument(s)."
            in str(err.value))

    # Test symbol table with 1 kernel argument
    arg1 = DataSymbol("arg1",
                      INTEGER_TYPE,
                      interface=ArgumentInterface(
                          ArgumentInterface.Access.READ))
    kschedule.symbol_table.add(arg1)
    kschedule.symbol_table.specify_argument_list([arg1])
    with pytest.raises(GenerationError) as err:
        _ = kschedule.symbol_table.iteration_indices
    assert ("GOcean 1.0 API kernels should always have at least two "
            "arguments representing the iteration indices but the Symbol "
            "Table for kernel 'test' has only 1 argument(s)."
            in str(err.value))

    # Test symbol table with 2 kernel argument
    arg2 = DataSymbol("arg2",
                      INTEGER_TYPE,
                      interface=ArgumentInterface(
                          ArgumentInterface.Access.READ))
    kschedule.symbol_table.add(arg2)
    kschedule.symbol_table.specify_argument_list([arg1, arg2])
    iteration_indices = kschedule.symbol_table.iteration_indices
    assert iteration_indices[0] is arg1
    assert iteration_indices[1] is arg2

    # Test symbol table with 3 kernel argument
    array_type = ArrayType(REAL_TYPE, [10, 10])
    arg3 = DataSymbol("buffer1",
                      array_type,
                      interface=ArgumentInterface(
                          ArgumentInterface.Access.READ))
    kschedule.symbol_table.add(arg3)
    kschedule.symbol_table.specify_argument_list([arg1, arg2, arg3])
    iteration_indices = kschedule.symbol_table.iteration_indices
    data_args = kschedule.symbol_table.data_arguments
    assert iteration_indices[0] is arg1
    assert iteration_indices[1] is arg2
    assert data_args[0] is arg3

    # Test gen_ocl with wrong iteration indices types and shapes.
    arg1._datatype._intrinsic = ScalarType.Intrinsic.REAL
    with pytest.raises(GenerationError) as err:
        _ = kschedule.symbol_table.iteration_indices
    assert ("GOcean 1.0 API kernels first argument should be a scalar "
            "integer but got 'Scalar<REAL, UNDEFINED>' for kernel 'test'."
            in str(err.value))

    arg1._datatype._intrinsic = ScalarType.Intrinsic.INTEGER  # restore
    arg2._datatype = ArrayType(INTEGER_TYPE, [10])
    with pytest.raises(GenerationError) as err:
        _ = kschedule.symbol_table.iteration_indices
    assert ("GOcean 1.0 API kernels second argument should be a scalar integer"
            " but got 'Array<Scalar<INTEGER, UNDEFINED>, shape=[10]>' for "
            "kernel 'test'." in str(err.value))
Beispiel #22
0
def test_scoping_node_copy_hierarchy():
    ''' Test that the ScopingNode copy() method creates a new symbol table
    with copied symbols and updates the children references.

    This test has 2 ScopingNodes, and the copy will only be applied to the
    inner one. This means that the References to the symbols on the outer
    scope should not be duplicated. Also it contains argument symbols and
    a reference inside another reference to make sure all these are copied
    appropriately.
    '''
    parent_node = Container("module")
    symbol_b = parent_node.symbol_table.new_symbol("b",
                                                   symbol_type=DataSymbol,
                                                   datatype=ArrayType(
                                                       INTEGER_TYPE, [5]))
    schedule = Routine("routine")
    parent_node.addchild(schedule)
    symbol_a = schedule.symbol_table.new_symbol(
        "a",
        symbol_type=DataSymbol,
        datatype=INTEGER_TYPE,
        interface=ArgumentInterface(ArgumentInterface.Access.READWRITE))
    schedule.symbol_table.specify_argument_list([symbol_a])
    symbol_i = schedule.symbol_table.new_symbol("i",
                                                symbol_type=DataSymbol,
                                                datatype=INTEGER_TYPE)

    schedule.addchild(
        Assignment.create(
            Reference(symbol_a),
            ArrayReference.create(symbol_b, [Reference(symbol_i)])))

    new_schedule = schedule.copy()

    # Check that the symbol_table has been deep copied
    assert new_schedule.symbol_table is not schedule.symbol_table
    assert new_schedule.symbol_table.lookup("i") is not \
        schedule.symbol_table.lookup("i")
    assert new_schedule.symbol_table.lookup("a") is not \
        schedule.symbol_table.lookup("a")

    # Check that 'a' and 'i' have been copied to the new symbol table.
    assert new_schedule[0].lhs.symbol not in schedule.symbol_table.symbols
    assert new_schedule[0].lhs.symbol in new_schedule.symbol_table.symbols
    assert new_schedule[0].rhs.children[0].symbol not in \
        schedule.symbol_table.symbols
    assert new_schedule[0].rhs.children[0].symbol in \
        new_schedule.symbol_table.symbols

    # Add the "_new" suffix to all symbol in the copied schedule
    for symbol in new_schedule.symbol_table.symbols:
        new_schedule.symbol_table.rename_symbol(symbol, symbol.name + "_new")

    # An update to a symbol in the outer scope must affect both copies of the
    # inner schedule.
    parent_node.symbol_table.rename_symbol(symbol_b, symbol_b.name + "_global")

    # Insert the schedule back to the original container
    parent_node.addchild(new_schedule)

    # Check that the expected code is generated
    # TODO #1200: the new 'routine' RoutineSymbol also needs to change.
    expected = '''\
module module
  implicit none
  integer, dimension(5) :: b_global

  contains
  subroutine routine(a)
    integer, intent(inout) :: a
    integer :: i

    a = b_global(i)

  end subroutine routine
  subroutine routine(a_new)
    integer, intent(inout) :: a_new
    integer :: i_new

    a_new = b_global(i_new)

  end subroutine routine

end module module
'''
    writer = FortranWriter()
    output = writer(parent_node)
    assert expected == output
Beispiel #23
0
def create_psyir_tree():
    ''' Create an example PSyIR Tree.

    :returns: an example PSyIR tree.
    :rtype: :py:class:`psyclone.psyir.nodes.Container`

    '''
    # Symbol table, symbols and scalar datatypes
    symbol_table = SymbolTable()
    arg1 = symbol_table.new_symbol(symbol_type=DataSymbol,
                                   datatype=REAL_TYPE,
                                   interface=ArgumentInterface(
                                       ArgumentInterface.Access.READWRITE))
    symbol_table.specify_argument_list([arg1])
    tmp_symbol = symbol_table.new_symbol(symbol_type=DataSymbol,
                                         datatype=REAL_DOUBLE_TYPE)
    index_symbol = symbol_table.new_symbol(root_name="i",
                                           symbol_type=DataSymbol,
                                           datatype=INTEGER4_TYPE)
    real_kind = symbol_table.new_symbol(root_name="RKIND",
                                        symbol_type=DataSymbol,
                                        datatype=INTEGER_TYPE,
                                        constant_value=8)
    routine_symbol = RoutineSymbol("my_sub")

    # Array using precision defined by another symbol
    scalar_type = ScalarType(ScalarType.Intrinsic.REAL, real_kind)
    array = symbol_table.new_symbol(root_name="a",
                                    symbol_type=DataSymbol,
                                    datatype=ArrayType(scalar_type, [10]))

    # Make generators for nodes which do not have other Nodes as children,
    # with some predefined scalar datatypes
    def zero():
        return Literal("0.0", REAL_TYPE)

    def one():
        return Literal("1.0", REAL4_TYPE)

    def two():
        return Literal("2.0", scalar_type)

    def int_zero():
        return Literal("0", INTEGER_SINGLE_TYPE)

    def int_one():
        return Literal("1", INTEGER8_TYPE)

    def tmp1():
        return Reference(arg1)

    def tmp2():
        return Reference(tmp_symbol)

    # Unary Operation
    oper = UnaryOperation.Operator.SIN
    unaryoperation = UnaryOperation.create(oper, tmp2())

    # Binary Operation
    oper = BinaryOperation.Operator.ADD
    binaryoperation = BinaryOperation.create(oper, one(), unaryoperation)

    # Nary Operation
    oper = NaryOperation.Operator.MAX
    naryoperation = NaryOperation.create(oper, [tmp1(), tmp2(), one()])

    # Array reference using a range
    lbound = BinaryOperation.create(BinaryOperation.Operator.LBOUND,
                                    Reference(array), int_one())
    ubound = BinaryOperation.create(BinaryOperation.Operator.UBOUND,
                                    Reference(array), int_one())
    my_range = Range.create(lbound, ubound)
    tmparray = ArrayReference.create(array, [my_range])

    # Assignments
    assign1 = Assignment.create(tmp1(), zero())
    assign2 = Assignment.create(tmp2(), zero())
    assign3 = Assignment.create(tmp2(), binaryoperation)
    assign4 = Assignment.create(tmp1(), tmp2())
    assign5 = Assignment.create(tmp1(), naryoperation)
    assign6 = Assignment.create(tmparray, two())

    # Call
    call = Call.create(routine_symbol, [tmp1(), binaryoperation.copy()])

    # If statement
    if_condition = BinaryOperation.create(BinaryOperation.Operator.GT, tmp1(),
                                          zero())
    ifblock = IfBlock.create(if_condition, [assign3, assign4])

    # Loop
    loop = Loop.create(index_symbol, int_zero(), int_one(), int_one(),
                       [ifblock])

    # KernelSchedule
    kernel_schedule = KernelSchedule.create(
        "work", symbol_table, [assign1, call, assign2, loop, assign5, assign6])

    # Container
    container_symbol_table = SymbolTable()
    container = Container.create("CONTAINER", container_symbol_table,
                                 [kernel_schedule])

    # Import data from another container
    external_container = ContainerSymbol("some_mod")
    container_symbol_table.add(external_container)
    external_var = DataSymbol("some_var",
                              INTEGER_TYPE,
                              interface=GlobalInterface(external_container))
    container_symbol_table.add(external_var)
    routine_symbol.interface = GlobalInterface(external_container)
    container_symbol_table.add(routine_symbol)
    return container
    def apply(self, node, options=None):
        '''Apply this transformation to the supplied node.

        :param node: the node to transform.
        :type node: :py:class:`psyclone.gocean1p0.GOKern`
        :param options: a dictionary with options for transformations.
        :type options: dict of string:values or None

        :returns: 2-tuple of new schedule and memento of transform.
        :rtype: (:py:class:`psyclone.gocean1p0.GOInvokeSchedule`, \
                 :py:class:`psyclone.undoredo.Memento`)

        '''
        self.validate(node, options)

        # Get useful references
        invoke_st = node.ancestor(InvokeSchedule).symbol_table
        inner_loop = node.ancestor(Loop)
        outer_loop = inner_loop.ancestor(Loop)
        cursor = outer_loop.position

        # Make sure the boundary symbols in the PSylayer exist
        inv_xstart = invoke_st.symbol_from_tag("xstart_" + node.name,
                                               root_name="xstart",
                                               symbol_type=DataSymbol,
                                               datatype=INTEGER_TYPE)
        inv_xstop = invoke_st.symbol_from_tag("xstop_" + node.name,
                                              root_name="xstop",
                                              symbol_type=DataSymbol,
                                              datatype=INTEGER_TYPE)
        inv_ystart = invoke_st.symbol_from_tag("ystart_" + node.name,
                                               root_name="ystart",
                                               symbol_type=DataSymbol,
                                               datatype=INTEGER_TYPE)
        inv_ystop = invoke_st.symbol_from_tag("ystop_" + node.name,
                                              root_name="ystop",
                                              symbol_type=DataSymbol,
                                              datatype=INTEGER_TYPE)

        # If the kernel acts on the whole iteration space, the boundary values
        # are not needed. This also avoids adding duplicated arguments if this
        # transformation is applied more than once to the same kernel. But the
        # declaration and initialisation above still needs to exist because the
        # boundary variables are expected to exist by the generation code.
        if (inner_loop.field_space == "go_every"
                and outer_loop.field_space == "go_every"
                and inner_loop.iteration_space == "go_all_pts"
                and outer_loop.iteration_space == "go_all_pts"):
            return node.root, None

        # Initialise the boundary values provided by the Loop construct
        assign1 = Assignment.create(Reference(inv_xstart),
                                    inner_loop.lower_bound().copy())
        outer_loop.parent.children.insert(cursor, assign1)
        cursor = cursor + 1
        assign2 = Assignment.create(Reference(inv_xstop),
                                    inner_loop.upper_bound().copy())
        outer_loop.parent.children.insert(cursor, assign2)
        cursor = cursor + 1
        assign3 = Assignment.create(Reference(inv_ystart),
                                    outer_loop.lower_bound().copy())
        outer_loop.parent.children.insert(cursor, assign3)
        cursor = cursor + 1
        assign4 = Assignment.create(Reference(inv_ystop),
                                    outer_loop.upper_bound().copy())
        outer_loop.parent.children.insert(cursor, assign4)

        # Update Kernel Call argument list
        for symbol in [inv_xstart, inv_xstop, inv_ystart, inv_ystop]:
            node.arguments.append(symbol.name, "go_i_scalar")

        # Now that the boundaries are inside the kernel, the looping should go
        # through all the field points
        inner_loop.field_space = "go_every"
        outer_loop.field_space = "go_every"
        inner_loop.iteration_space = "go_all_pts"
        outer_loop.iteration_space = "go_all_pts"

        # Update Kernel
        kschedule = node.get_kernel_schedule()
        kernel_st = kschedule.symbol_table
        iteration_indices = kernel_st.iteration_indices
        data_arguments = kernel_st.data_arguments

        # Create new symbols and insert them as kernel arguments at the end of
        # the kernel argument list
        xstart_symbol = kernel_st.new_symbol(
            "xstart",
            symbol_type=DataSymbol,
            datatype=INTEGER_TYPE,
            interface=ArgumentInterface(ArgumentInterface.Access.READ))
        xstop_symbol = kernel_st.new_symbol("xstop",
                                            symbol_type=DataSymbol,
                                            datatype=INTEGER_TYPE,
                                            interface=ArgumentInterface(
                                                ArgumentInterface.Access.READ))
        ystart_symbol = kernel_st.new_symbol(
            "ystart",
            symbol_type=DataSymbol,
            datatype=INTEGER_TYPE,
            interface=ArgumentInterface(ArgumentInterface.Access.READ))
        ystop_symbol = kernel_st.new_symbol("ystop",
                                            symbol_type=DataSymbol,
                                            datatype=INTEGER_TYPE,
                                            interface=ArgumentInterface(
                                                ArgumentInterface.Access.READ))
        kernel_st.specify_argument_list(
            iteration_indices + data_arguments +
            [xstart_symbol, xstop_symbol, ystart_symbol, ystop_symbol])

        # Create boundary masking conditions
        condition1 = BinaryOperation.create(BinaryOperation.Operator.LT,
                                            Reference(iteration_indices[0]),
                                            Reference(xstart_symbol))
        condition2 = BinaryOperation.create(BinaryOperation.Operator.GT,
                                            Reference(iteration_indices[0]),
                                            Reference(xstop_symbol))
        condition3 = BinaryOperation.create(BinaryOperation.Operator.LT,
                                            Reference(iteration_indices[1]),
                                            Reference(ystart_symbol))
        condition4 = BinaryOperation.create(BinaryOperation.Operator.GT,
                                            Reference(iteration_indices[1]),
                                            Reference(ystop_symbol))

        condition = BinaryOperation.create(
            BinaryOperation.Operator.OR,
            BinaryOperation.create(BinaryOperation.Operator.OR, condition1,
                                   condition2),
            BinaryOperation.create(BinaryOperation.Operator.OR, condition3,
                                   condition4))

        # Insert the conditional mask as the first statement of the kernel
        if_statement = IfBlock.create(condition, [Return()])
        kschedule.children.insert(0, if_statement)

        return node.root, None
Beispiel #25
0
def test_oclw_kernelschedule():
    '''Check the OpenCLWriter class kernelschedule_node visitor produces
    the expected OpenCL code.

    '''

    # The kernelschedule OpenCL Backend relies on abstract methods that
    # need to be implemented by the APIs. A generic kernelschedule will
    # produce a NotImplementedError.
    oclwriter = OpenCLWriter()
    kschedule = KernelSchedule("kname")
    with pytest.raises(NotImplementedError) as error:
        _ = oclwriter(kschedule)
    assert "Abstract property. Which symbols are data arguments is " \
        "API-specific." in str(error.value)

    # Mock abstract properties. (pytest monkeypatch does not work
    # with properties, used sub-class instead)
    class MockSymbolTable(SymbolTable):
        ''' Mock needed abstract methods of the Symbol Table '''
        @property
        def iteration_indices(self):
            return self.argument_list[:2]

        @property
        def data_arguments(self):
            return self.argument_list[2:]

    kschedule.symbol_table.__class__ = MockSymbolTable

    # Create a sample symbol table and kernel schedule
    interface = ArgumentInterface(ArgumentInterface.Access.UNKNOWN)
    i = DataSymbol('i', INTEGER_TYPE, interface=interface)
    j = DataSymbol('j', INTEGER_TYPE, interface=interface)
    array_type = ArrayType(REAL_TYPE, [10, 10])
    data1 = DataSymbol('data1', array_type, interface=interface)
    data2 = DataSymbol('data2', array_type, interface=interface)
    kschedule.symbol_table.add(i)
    kschedule.symbol_table.add(j)
    kschedule.symbol_table.add(data1)
    kschedule.symbol_table.add(data2)
    kschedule.symbol_table.specify_argument_list([i, j, data1, data2])
    kschedule.addchild(Return(parent=kschedule))

    result = oclwriter(kschedule)
    assert result == "" \
        "__kernel void kname(\n" \
        "  __global double * restrict data1,\n" \
        "  __global double * restrict data2\n" \
        "  ){\n" \
        "  int data1LEN1 = get_global_size(0);\n" \
        "  int data1LEN2 = get_global_size(1);\n" \
        "  int data2LEN1 = get_global_size(0);\n" \
        "  int data2LEN2 = get_global_size(1);\n" \
        "  int i = get_global_id(0);\n" \
        "  int j = get_global_id(1);\n" \
        "  return;\n" \
        "}\n\n"

    # Set a local_size value different to 1 into the KernelSchedule
    oclwriter = OpenCLWriter(kernels_local_size=4)
    result = oclwriter(kschedule)

    assert result == "" \
        "__attribute__((reqd_work_group_size(4, 1, 1)))\n" \
        "__kernel void kname(\n" \
        "  __global double * restrict data1,\n" \
        "  __global double * restrict data2\n" \
        "  ){\n" \
        "  int data1LEN1 = get_global_size(0);\n" \
        "  int data1LEN2 = get_global_size(1);\n" \
        "  int data2LEN1 = get_global_size(0);\n" \
        "  int data2LEN2 = get_global_size(1);\n" \
        "  int i = get_global_id(0);\n" \
        "  int j = get_global_id(1);\n" \
        "  return;\n" \
        "}\n\n"

    # Add a symbol with a deferred interface and check that this raises the
    # expected error
    array_type = ArrayType(REAL_TYPE, [10, 10])
    kschedule.symbol_table.add(
        DataSymbol('broken', array_type, interface=UnresolvedInterface()))
    with pytest.raises(VisitorError) as err:
        _ = oclwriter(kschedule)
    assert ("symbol table contains unresolved data entries (i.e. that have no "
            "defined Interface) which are not used purely to define the "
            "precision of other symbols: 'broken'" in str(err.value))
Beispiel #26
0
def test_argumentinterface_str():
    '''Test that an ArgumentInterface instance can be stringified'''

    argument_interface = ArgumentInterface()
    assert str(argument_interface) == "Argument(pass-by-value=False)"
Beispiel #27
0
def test_operator():
    '''Test the KernelInterface class operator method. We want to check
    that the correct symbol is referenced for the dimension of the
    operator symbol so the simplest solution is to use one of the Fortran
    test examples.

    '''
    kernel_interface = KernelInterface(None)
    _, invoke_info = parse(os.path.join(BASE_PATH, "10_operator.f90"),
                           api="dynamo0.3")
    psy = PSyFactory("dynamo0.3", distributed_memory=False).create(invoke_info)
    schedule = psy.invokes.invoke_list[0].schedule
    kernel = schedule[0].loop_body[0]
    operator_arg = kernel.args[0]
    kernel_interface.operator(operator_arg)

    # fs_from symbol declared
    fs_from_name = operator_arg.function_space_from.orig_name
    fs_from_tag = "ndf_{0}".format(fs_from_name)
    fs_from_symbol = kernel_interface._symbol_table.lookup(fs_from_tag)
    assert isinstance(fs_from_symbol, lfric_psyir.NumberOfDofsDataSymbol)
    assert fs_from_symbol.fs == fs_from_name
    assert isinstance(fs_from_symbol.interface, ArgumentInterface)
    assert (fs_from_symbol.interface.access ==
            kernel_interface._read_access.access)

    # fs_to symbol declared
    fs_to_name = operator_arg.function_space_from.orig_name
    fs_to_tag = "ndf_{0}".format(fs_to_name)
    fs_to_symbol = kernel_interface._symbol_table.lookup(fs_to_tag)
    assert isinstance(fs_to_symbol, lfric_psyir.NumberOfDofsDataSymbol)
    assert fs_to_symbol.fs == fs_to_name
    assert isinstance(fs_to_symbol.interface, ArgumentInterface)
    assert (
        fs_to_symbol.interface.access == kernel_interface._read_access.access)

    # ncells symbol declared
    ncells_symbol = kernel_interface._symbol_table.lookup("ncell_3d")
    assert isinstance(ncells_symbol, lfric_psyir.NumberOfCellsDataSymbol)
    assert isinstance(ncells_symbol.interface, ArgumentInterface)
    assert (
        ncells_symbol.interface.access == kernel_interface._read_access.access)

    # operator declared, added to argument list, correct function
    # spaces specified and dimensioned correctly
    tag = operator_arg.name
    symbol = kernel_interface._symbol_table.lookup(tag)
    assert isinstance(symbol, lfric_psyir.OperatorDataSymbol)
    assert isinstance(symbol.interface, ArgumentInterface)
    assert (symbol.interface.access == ArgumentInterface(
        INTENT_MAPPING[operator_arg.intent]).access)
    assert kernel_interface._arglist[-1] is symbol
    assert symbol.fs_from == operator_arg.function_space_from.orig_name
    assert symbol.fs_to == operator_arg.function_space_to.orig_name
    assert len(symbol.shape) == 3
    assert isinstance(symbol.shape[0], Reference)
    assert symbol.shape[0].symbol is fs_from_symbol
    assert isinstance(symbol.shape[1], Reference)
    assert symbol.shape[1].symbol is fs_to_symbol
    assert isinstance(symbol.shape[2], Reference)
    assert symbol.shape[2].symbol is ncells_symbol
Beispiel #28
0
def create_psyir_tree():
    ''' Create an example PSyIR Tree.

    :returns: an example PSyIR tree.
    :rtype: :py:class:`psyclone.psyir.nodes.Container`

    '''
    # Symbol table, symbols and scalar datatypes
    symbol_table = SymbolTable()
    arg1 = symbol_table.new_symbol(symbol_type=DataSymbol,
                                   datatype=REAL_TYPE,
                                   interface=ArgumentInterface(
                                       ArgumentInterface.Access.READWRITE))
    symbol_table.specify_argument_list([arg1])
    tmp_symbol = symbol_table.new_symbol(symbol_type=DataSymbol,
                                         datatype=REAL_DOUBLE_TYPE)
    index_symbol = symbol_table.new_symbol(root_name="i",
                                           symbol_type=DataSymbol,
                                           datatype=INTEGER4_TYPE)
    real_kind = symbol_table.new_symbol(root_name="RKIND",
                                        symbol_type=DataSymbol,
                                        datatype=INTEGER_TYPE,
                                        constant_value=8)
    routine_symbol = RoutineSymbol("my_sub")

    # Array using precision defined by another symbol
    scalar_type = ScalarType(ScalarType.Intrinsic.REAL, real_kind)
    array = symbol_table.new_symbol(root_name="a",
                                    symbol_type=DataSymbol,
                                    datatype=ArrayType(scalar_type, [10]))

    # Nodes which do not have Nodes as children and (some) predefined
    # scalar datatypes
    # TODO: Issue #1136 looks at how to avoid all of the _x versions
    zero_1 = Literal("0.0", REAL_TYPE)
    zero_2 = Literal("0.0", REAL_TYPE)
    zero_3 = Literal("0.0", REAL_TYPE)
    one_1 = Literal("1.0", REAL4_TYPE)
    one_2 = Literal("1.0", REAL4_TYPE)
    one_3 = Literal("1.0", REAL4_TYPE)
    two = Literal("2.0", scalar_type)
    int_zero = Literal("0", INTEGER_SINGLE_TYPE)
    int_one_1 = Literal("1", INTEGER8_TYPE)
    int_one_2 = Literal("1", INTEGER8_TYPE)
    int_one_3 = Literal("1", INTEGER8_TYPE)
    int_one_4 = Literal("1", INTEGER8_TYPE)
    tmp1_1 = Reference(arg1)
    tmp1_2 = Reference(arg1)
    tmp1_3 = Reference(arg1)
    tmp1_4 = Reference(arg1)
    tmp1_5 = Reference(arg1)
    tmp1_6 = Reference(arg1)
    tmp2_1 = Reference(tmp_symbol)
    tmp2_2 = Reference(tmp_symbol)
    tmp2_3 = Reference(tmp_symbol)
    tmp2_4 = Reference(tmp_symbol)
    tmp2_5 = Reference(tmp_symbol)
    tmp2_6 = Reference(tmp_symbol)

    # Unary Operation
    oper = UnaryOperation.Operator.SIN
    unaryoperation_1 = UnaryOperation.create(oper, tmp2_1)
    unaryoperation_2 = UnaryOperation.create(oper, tmp2_2)

    # Binary Operation
    oper = BinaryOperation.Operator.ADD
    binaryoperation_1 = BinaryOperation.create(oper, one_1, unaryoperation_1)
    binaryoperation_2 = BinaryOperation.create(oper, one_2, unaryoperation_2)

    # Nary Operation
    oper = NaryOperation.Operator.MAX
    naryoperation = NaryOperation.create(oper, [tmp1_1, tmp2_3, one_3])

    # Array reference using a range
    lbound = BinaryOperation.create(BinaryOperation.Operator.LBOUND,
                                    Reference(array), int_one_1)
    ubound = BinaryOperation.create(BinaryOperation.Operator.UBOUND,
                                    Reference(array), int_one_2)
    my_range = Range.create(lbound, ubound)
    tmparray = ArrayReference.create(array, [my_range])

    # Assignments
    assign1 = Assignment.create(tmp1_2, zero_1)
    assign2 = Assignment.create(tmp2_4, zero_2)
    assign3 = Assignment.create(tmp2_5, binaryoperation_1)
    assign4 = Assignment.create(tmp1_3, tmp2_6)
    assign5 = Assignment.create(tmp1_4, naryoperation)
    assign6 = Assignment.create(tmparray, two)

    # Call
    call = Call.create(routine_symbol, [tmp1_5, binaryoperation_2])

    # If statement
    if_condition = BinaryOperation.create(BinaryOperation.Operator.GT, tmp1_6,
                                          zero_3)
    ifblock = IfBlock.create(if_condition, [assign3, assign4])

    # Loop
    loop = Loop.create(index_symbol, int_zero, int_one_3, int_one_4, [ifblock])

    # KernelSchedule
    kernel_schedule = KernelSchedule.create(
        "work", symbol_table, [assign1, call, assign2, loop, assign5, assign6])

    # Container
    container_symbol_table = SymbolTable()
    container = Container.create("CONTAINER", container_symbol_table,
                                 [kernel_schedule])

    # Import data from another container
    external_container = ContainerSymbol("some_mod")
    container_symbol_table.add(external_container)
    external_var = DataSymbol("some_var",
                              INTEGER_TYPE,
                              interface=GlobalInterface(external_container))
    container_symbol_table.add(external_var)
    routine_symbol.interface = GlobalInterface(external_container)
    container_symbol_table.add(routine_symbol)
    return container
Beispiel #29
0
def test_validate_kernel_code_arg(monkeypatch):
    '''Test that this method returns successfully if its two arguments
    have identical content, otherwise test that the expected
    exceptions are raised.

    '''
    kernel = DynKern()
    # Kernel name needs to be set when testing exceptions.
    kernel._name = "dummy"
    read_access = ArgumentInterface(ArgumentInterface.Access.READ)

    real_scalar_symbol = DataSymbol("generic_real_scalar",
                                    REAL_TYPE,
                                    interface=read_access)
    int_scalar_symbol = DataSymbol("generic_int_scalar",
                                   INTEGER_TYPE,
                                   interface=read_access)
    real_scalar_rw_symbol = DataSymbol("generic_scalar_rw",
                                       REAL_TYPE,
                                       interface=ArgumentInterface(
                                           ArgumentInterface.Access.READWRITE))
    lfric_real_scalar_symbol = LfricRealScalarDataSymbol("scalar",
                                                         interface=read_access)
    lfric_int_scalar_symbol = LfricIntegerScalarDataSymbol(
        "scalar", interface=read_access)
    lfric_real_field_symbol = RealFieldDataDataSymbol("field",
                                                      dims=[1],
                                                      fs="w0",
                                                      interface=read_access)

    kernel._validate_kernel_code_arg(lfric_real_scalar_symbol,
                                     lfric_real_scalar_symbol)

    with pytest.raises(GenerationError) as info:
        kernel._validate_kernel_code_arg(lfric_real_scalar_symbol,
                                         lfric_int_scalar_symbol)
    assert ("Kernel argument 'scalar' has datatype 'Intrinsic.REAL' in kernel "
            "'dummy' but the LFRic API expects 'Intrinsic.INTEGER'"
            in str(info.value))

    with pytest.raises(GenerationError) as info:
        kernel._validate_kernel_code_arg(real_scalar_symbol,
                                         lfric_real_scalar_symbol)
    assert ("Kernel argument 'generic_real_scalar' has precision 'UNDEFINED' "
            "in kernel 'dummy' but the LFRic API expects 'r_def'."
            in str(info.value))

    with pytest.raises(GenerationError) as info:
        kernel._validate_kernel_code_arg(real_scalar_symbol,
                                         real_scalar_rw_symbol)
    assert ("Kernel argument 'generic_real_scalar' has intent 'READ' in "
            "kernel 'dummy' but the LFRic API expects intent "
            "'READWRITE'." in str(info.value))

    with pytest.raises(GenerationError) as info:
        kernel._validate_kernel_code_arg(lfric_real_field_symbol,
                                         lfric_real_scalar_symbol)
    assert ("Argument 'field' to kernel 'dummy' should be a scalar "
            "according to the LFRic API, but it is not." in str(info.value))

    with pytest.raises(GenerationError) as info:
        kernel._validate_kernel_code_arg(lfric_real_scalar_symbol,
                                         lfric_real_field_symbol)
    assert ("Argument 'scalar' to kernel 'dummy' should be an array "
            "according to the LFRic API, but it is not." in str(info.value))

    undf = NumberOfUniqueDofsDataSymbol("undf", fs="w0", interface=read_access)
    lfric_real_field_symbol2 = RealFieldDataDataSymbol("field",
                                                       dims=[Reference(undf)],
                                                       fs="w0",
                                                       interface=read_access)
    # if one of the dimensions is not a datasymbol then the arguments
    # are not checked.
    kernel._validate_kernel_code_arg(lfric_real_field_symbol,
                                     lfric_real_field_symbol2)
    kernel._validate_kernel_code_arg(lfric_real_field_symbol2,
                                     lfric_real_field_symbol)

    lfric_real_field_symbol3 = RealFieldDataDataSymbol("field",
                                                       dims=[Reference(undf)],
                                                       fs="w0",
                                                       interface=read_access)
    monkeypatch.setattr(lfric_real_field_symbol3.datatype, "_shape",
                        [Reference(undf), Reference(undf)])
    with pytest.raises(GenerationError) as info:
        kernel._validate_kernel_code_arg(lfric_real_field_symbol2,
                                         lfric_real_field_symbol3)
    assert ("Argument 'field' to kernel 'dummy' should be an array with 2 "
            "dimension(s) according to the LFRic API, but found 1."
            in str(info.value))

    lfric_real_field_symbol4 = RealFieldDataDataSymbol(
        "field",
        dims=[Reference(int_scalar_symbol)],
        fs="w0",
        interface=read_access)
    with pytest.raises(GenerationError) as info:
        kernel._validate_kernel_code_arg(lfric_real_field_symbol4,
                                         lfric_real_field_symbol2)
    assert ("For dimension 1 in array argument 'field' to kernel 'dummy' the "
            "following error was found: Kernel argument 'generic_int_scalar' "
            "has precision 'UNDEFINED' in kernel 'dummy' but the LFRic API "
            "expects 'i_def'" in str(info.value))

    # monkeypatch lfric_real_scalar_symbol to return that it is not a
    # scalar in order to force the required exception. We do this by
    # changing the ScalarType as it is used when determining whether
    # the symbol is a scalar.
    monkeypatch.setattr(psyclone.psyir.symbols, "ScalarType", str)
    with pytest.raises(InternalError) as info:
        kernel._validate_kernel_code_arg(lfric_real_scalar_symbol,
                                         lfric_real_scalar_symbol)
    assert ("unexpected argument type found for 'scalar' in kernel 'dummy'. "
            "Expecting a scalar or an array." in str(info.value))
Beispiel #30
0
class KernelInterface(ArgOrdering):
    '''Create the kernel arguments for the supplied kernel as specified by
    the associated kernel metadata and the kernel ordering rules
    encoded in the ArgOrdering base class as method callbacks.

    A PSyIR symbol table is created containing appropriate LFRic PSyIR
    symbols to specify the arguments. If an argument is an array with
    one or more dimension sizes specified by another argument, then
    the associated array symbol is created so that it references the
    appropriate symbol.

    Related arguments - e.g. a field has an associated dofmap - are
    not directly connected, they must be inferred from the function
    space names. It is not yet clear whether this would be useful or
    not.

    TBD: This class should replace the current kernel stub generation
    code when all of its methods are implemented, see issue #928.

    :param kern: the kernel for which to create arguments.
    :type kern: :py:class:`psyclone.dynamo0p3.DynKern`

    '''
    #: Mapping from a generic PSyIR datatype to the equivalent
    #: LFRic-specific field datasymbol.
    field_mapping = {
        "integer": lfric_psyir.IntegerFieldDataDataSymbol,
        "real": lfric_psyir.RealFieldDataDataSymbol,
        "logical": lfric_psyir.LogicalFieldDataDataSymbol
    }
    #: Mapping from a generic PSyIR datatype to the equivalent
    #: LFRic-specific vector field datasymbol.
    vector_field_mapping = {
        "integer": lfric_psyir.IntegerVectorFieldDataDataSymbol,
        "real": lfric_psyir.RealVectorFieldDataDataSymbol,
        "logical": lfric_psyir.LogicalVectorFieldDataDataSymbol
    }
    #: Mapping from the LFRic metadata description of quadrature to the
    #: associated LFRic-specific basis function datasymbol.
    basis_mapping = {
        "gh_quadrature_xyoz": lfric_psyir.BasisFunctionQrXyozDataSymbol,
        "gh_quadrature_face": lfric_psyir.BasisFunctionQrFaceDataSymbol,
        "gh_quadrature_edge": lfric_psyir.BasisFunctionQrEdgeDataSymbol
    }
    #: Mapping from the LFRic metadata description of quadrature to the
    #: associated LFRic-specific differential basis function datasymbol.
    diff_basis_mapping = {
        "gh_quadrature_xyoz": lfric_psyir.DiffBasisFunctionQrXyozDataSymbol,
        "gh_quadrature_face": lfric_psyir.DiffBasisFunctionQrFaceDataSymbol,
        "gh_quadrature_edge": lfric_psyir.DiffBasisFunctionQrEdgeDataSymbol
    }
    _read_access = ArgumentInterface(ArgumentInterface.Access.READ)

    def __init__(self, kern):
        super(KernelInterface, self).__init__(kern)
        self._symbol_table = SymbolTable()
        self._arglist = []

    def generate(self, var_accesses=None):
        '''Call the generate base class then add the argument list as it can't
        be appended as we go along.

        :param var_accesses: an unused optional argument that stores \
            information about variable accesses.
        :type var_accesses: :\
            py:class:`psyclone.core.access_info.VariablesAccessInfo`

        '''
        super(KernelInterface, self).generate(var_accesses=var_accesses)
        # Set the argument list for the symbol table. This is done at
        # the end after incrementally adding symbols to the _args
        # list, as it is not possible to incrementally add symbols to
        # the symbol table argument list.
        self._symbol_table.specify_argument_list(self._arglist)
        # While data dependence analysis does not use the symbol
        # table, see #845, we have to provide variable information
        # separately. This is done by using the base class append()
        # method. However, this method also adds the variable names to the
        # internal _arglist list which we do not want as we have
        # already added our symbols there. Therefore we need to remove
        # them afterwards.
        # Map from symbol table accesses to dependence analysis accesses.
        mapping = {
            ArgumentInterface.Access.READ: AccessType.READ,
            ArgumentInterface.Access.READWRITE: AccessType.READWRITE,
            ArgumentInterface.Access.WRITE: AccessType.WRITE
        }
        len_arglist = len(self._arglist)
        for symbol in self._symbol_table.symbols:
            self.append(symbol.name,
                        var_accesses,
                        mode=mapping[symbol.interface.access])
        self._arglist = self._arglist[:len_arglist]

    def cell_position(self, var_accesses=None):
        '''Create an LFRic cell-position object and add it to the symbol table
        and argument list.

        :param var_accesses: an unused optional argument that stores \
            information about variable accesses.
        :type var_accesses: :\
            py:class:`psyclone.core.access_info.VariablesAccessInfo`

        '''
        symbol = self._symbol_table.symbol_from_tag(
            "cell",
            symbol_type=lfric_psyir.CellPositionDataSymbol,
            interface=self._read_access)
        self._arglist.append(symbol)

    def mesh_height(self, var_accesses=None):
        '''Create an LFRic mesh height object and add it to the symbol table
        and argument list.

        :param var_accesses: an unused optional argument that stores \
            information about variable accesses.
        :type var_accesses: :\
            py:class:`psyclone.core.access_info.VariablesAccessInfo`

        '''
        symbol = self._symbol_table.symbol_from_tag(
            "nlayers",
            symbol_type=lfric_psyir.MeshHeightDataSymbol,
            interface=self._read_access)
        self._arglist.append(symbol)

    def mesh_ncell2d(self, var_accesses=None):
        '''Not implemented.

        :param var_accesses: an unused optional argument that stores \
            information about variable accesses.
        :type var_accesses: :\
            py:class:`psyclone.core.access_info.VariablesAccessInfo`

        :raises NotImplementedError: as this method is not implemented.

        '''
        raise NotImplementedError("mesh_ncell2d not implemented")

    def cell_map(self, var_accesses=None):
        '''Not implemented.

        :param var_accesses: an unused optional argument that stores \
            information about variable accesses.
        :type var_accesses: :\
            py:class:`psyclone.core.access_info.VariablesAccessInfo`

        :raises NotImplementedError: as this method is not implemented.

        '''
        raise NotImplementedError("cell_map not implemented")

    def field_vector(self, argvect, var_accesses=None):
        '''Create LFRic field vector arguments and add them to the symbol
        table and argument list. Also declare the associated "undf"
        symbol if it has not already been declared so that it can be
        used to dimension the field vector arguments.

        :param argvect: the field vector to add.
        :type argvect: :py:class:`psyclone.dynamo0p3.DynKernelArgument`
        :param var_accesses: an unused optional argument that stores \
            information about variable accesses.
        :type var_accesses: :\
            py:class:`psyclone.core.access_info.VariablesAccessInfo`

        :raises NotImplementedError: if the datatype of the vector \
            field is not supported.

        '''
        fs_name = argvect.function_space.orig_name
        undf_symbol = self._symbol_table.symbol_from_tag(
            "undf_{0}".format(fs_name),
            fs=fs_name,
            symbol_type=lfric_psyir.NumberOfUniqueDofsDataSymbol,
            interface=self._read_access)

        interface = ArgumentInterface(INTENT_MAPPING[argvect.intent])
        try:
            field_class = self.vector_field_mapping[argvect.intrinsic_type]
        except KeyError as info:
            message = ("kernel interface does not support a vector field of "
                       "type '{0}'.".format(argvect.intrinsic_type))
            six.raise_from(NotImplementedError(message), info)
        for idx in range(argvect.vector_size):
            tag = "{0}_v{1}".format(argvect.name, idx)
            field_data_symbol = self._symbol_table.symbol_from_tag(
                tag,
                symbol_type=field_class,
                dims=[Reference(undf_symbol)],
                fs=fs_name,
                interface=interface)
            self._arglist.append(field_data_symbol)

    def field(self, arg, var_accesses=None):
        '''Create an LFRic field argument and add it to the symbol table and
        argument list. Also declare the associated "undf" symbol if it
        has not already been declared so that it can be used to
        dimension the field argument.

        :param arg: the field to add.
        :type arg: :py:class:`psyclone.dynamo0p3.DynKernelArgument`
        :param var_accesses: an unused optional argument that stores \
            information about variable accesses.
        :type var_accesses: :\
            py:class:`psyclone.core.access_info.VariablesAccessInfo`

        :raises NotImplementedError: if the datatype of the field is \
            not supported.

        '''
        fs_name = arg.function_space.orig_name
        undf_symbol = self._symbol_table.symbol_from_tag(
            "undf_{0}".format(fs_name),
            symbol_type=lfric_psyir.NumberOfUniqueDofsDataSymbol,
            fs=fs_name,
            interface=self._read_access)

        try:
            field_class = self.field_mapping[arg.intrinsic_type]
        except KeyError as info:
            message = ("kernel interface does not support a field of type "
                       "'{0}'.".format(arg.intrinsic_type))
            six.raise_from(NotImplementedError(message), info)
        field_data_symbol = self._symbol_table.symbol_from_tag(
            arg.name,
            interface=ArgumentInterface(INTENT_MAPPING[arg.intent]),
            symbol_type=field_class,
            dims=[Reference(undf_symbol)],
            fs=fs_name)
        self._arglist.append(field_data_symbol)

    def stencil_unknown_extent(self, arg, var_accesses=None):
        '''Not implemented.

        :param arg: the kernel argument with which the stencil is associated.
        :type arg: :py:class:`psyclone.dynamo0p3.DynKernelArgument`
        :param var_accesses: an unused optional argument that stores \
            information about variable accesses.
        :type var_accesses: :\
            py:class:`psyclone.core.access_info.VariablesAccessInfo`

        :raises NotImplementedError: as this method is not implemented.

        '''
        raise NotImplementedError("stencil_unknown_extent not implemented")

    def stencil_unknown_direction(self, arg, var_accesses=None):
        '''Not implemented.

        :param arg: the kernel argument with which the stencil is associated.
        :type arg: :py:class:`psyclone.dynamo0p3.DynKernelArgument`
        :param var_accesses: an unused optional argument that stores \
            information about variable accesses.
        :type var_accesses: :\
            py:class:`psyclone.core.access_info.VariablesAccessInfo`

        :raises NotImplementedError: as this method is not implemented.

        '''
        raise NotImplementedError("stencil_unknown_direction not implemented")

    def stencil(self, arg, var_accesses=None):
        '''Not implemented.

        :param arg: the kernel argument with which the stencil is associated.
        :type arg: :py:class:`psyclone.dynamo0p3.DynKernelArgument`
        :param var_accesses: an unused optional argument that stores \
            information about variable accesses.
        :type var_accesses: :\
            py:class:`psyclone.core.access_info.VariablesAccessInfo`

        :raises NotImplementedError: as this method is not implemented.

        '''
        raise NotImplementedError("stencil not implemented")

    def operator(self, arg, var_accesses=None):
        '''Create an LFRic operator argument and an ncells argument and add
        them to the symbol table and argument list. Also declare the
        associated 'fs_from', 'fs_to' symbols if they have not already
        been declared so that they can be used to dimension the
        operator symbol (as well as ncells).

        :param arg: the operator to add.
        :type arg: :py:class:`psyclone.dynamo0p3.DynKernelArgument`
        :param var_accesses: an unused optional argument that stores \
            information about variable accesses.
        :type var_accesses: :\
            py:class:`psyclone.core.access_info.VariablesAccessInfo`

        :raises NotImplementedError: if the datatype of the field is \
            not supported.

        '''
        fs_from_name = arg.function_space_from.orig_name
        ndf_symbol_from = self._symbol_table.symbol_from_tag(
            "ndf_{0}".format(fs_from_name),
            fs=fs_from_name,
            symbol_type=lfric_psyir.NumberOfDofsDataSymbol,
            interface=self._read_access)
        fs_to_name = arg.function_space_to.orig_name
        ndf_symbol_to = self._symbol_table.symbol_from_tag(
            "ndf_{0}".format(fs_to_name),
            fs=fs_to_name,
            symbol_type=lfric_psyir.NumberOfDofsDataSymbol,
            interface=self._read_access)

        ncells = lfric_psyir.NumberOfCellsDataSymbol(
            "ncell_3d", interface=self._read_access)
        self._symbol_table.add(ncells)
        self._arglist.append(ncells)

        op_arg_symbol = self._symbol_table.symbol_from_tag(
            arg.name,
            symbol_type=lfric_psyir.OperatorDataSymbol,
            dims=[
                Reference(ndf_symbol_from),
                Reference(ndf_symbol_to),
                Reference(ncells)
            ],
            fs_from=fs_from_name,
            fs_to=fs_to_name,
            interface=ArgumentInterface(INTENT_MAPPING[arg.intent]))
        self._arglist.append(op_arg_symbol)

    def cma_operator(self, arg, var_accesses=None):
        '''Not implemented.

        :param arg: the CMA operator argument.
        :type arg: :py:class:`psyclone.dynamo0p3.DynKernelArgument`
        :param var_accesses: an unused optional argument that stores \
            information about variable accesses.
        :type var_accesses: :\
            py:class:`psyclone.core.access_info.VariablesAccessInfo`

        :raises NotImplementedError: as this method is not implemented.

        '''
        raise NotImplementedError("cma_operator not implemented")

    def scalar(self, scalar_arg, var_accesses=None):
        '''Create an LFRic scalar argument and add it to the symbol table and
        argument list.

        :param scalar_arg: the scalar to add.
        :type scalar_arg: :py:class:`psyclone.dynamo0p3.DynKernelArgument`
        :param var_accesses: an unused optional argument that stores \
            information about variable accesses.
        :type var_accesses: :\
            py:class:`psyclone.core.access_info.VariablesAccessInfo`

        :raises NotImplementedError: if the datatype of the scalar is \
            not supported.

        '''
        mapping = {
            "integer": lfric_psyir.LfricIntegerScalarDataSymbol,
            "real": lfric_psyir.LfricRealScalarDataSymbol,
            "logical": lfric_psyir.LfricLogicalScalarDataSymbol
        }
        try:
            symbol = self._symbol_table.symbol_from_tag(
                scalar_arg.name,
                symbol_type=mapping[scalar_arg.intrinsic_type],
                interface=ArgumentInterface(INTENT_MAPPING[scalar_arg.intent]))
        except KeyError as info:
            message = (
                "scalar of type '{0}' not implemented in KernelInterface "
                "class.".format(scalar_arg.intrinsic_type))
            six.raise_from(NotImplementedError(message), info)
        self._arglist.append(symbol)

    def fs_common(self, function_space, var_accesses=None):
        '''Create any arguments that are common to a particular function
        space. At this time the only common argument is the number of
        degrees of freedom. Create the associated LFRic symbol, and
        add it to the symbol table and argument list.

        :param function_space: the function space for any common arguments.
        :type function_space: :py:class:`psyclone.domain.lfric.FunctionSpace`
        :param var_accesses: an unused optional argument that stores \
            information about variable accesses.
        :type var_accesses: :\
            py:class:`psyclone.core.access_info.VariablesAccessInfo`

        '''
        fs_name = function_space.orig_name
        ndf_symbol = self._symbol_table.symbol_from_tag(
            "ndf_{0}".format(fs_name),
            fs=fs_name,
            symbol_type=lfric_psyir.NumberOfDofsDataSymbol,
            interface=self._read_access)
        self._arglist.append(ndf_symbol)

    def fs_intergrid(self, function_space, var_accesses=None):
        '''Not implemented.

        :param arg: the CMA operator argument.
        :type arg: :py:class:`psyclone.dynamo0p3.DynKernelArgument`
        :param var_accesses: an unused optional argument that stores \
            information about variable accesses.
        :type var_accesses: :\
            py:class:`psyclone.core.access_info.VariablesAccessInfo`

        :raises NotImplementedError: as this method is not implemented.

        '''
        raise NotImplementedError("fs_intergrid not implemented")

    def fs_compulsory_field(self, function_space, var_accesses=None):
        '''Create any arguments that are compulsory for a field on a
        particular function space. At this time the compulsory
        arguments are the unique number of degrees of freedom and the
        dofmap. Create the associated LFRic symbol, and add it to the
        symbol table and argument list. Also declare the number of
        degrees of freedom and add to the symbol table if one has not
        yet been added.

        :param function_space: the function space for any compulsory arguments.
        :type function_space: :py:class:`psyclone.domain.lfric.FunctionSpace`
        :param var_accesses: an unused optional argument that stores \
            information about variable accesses.
        :type var_accesses: :\
            py:class:`psyclone.core.access_info.VariablesAccessInfo`

        '''
        fs_name = function_space.orig_name
        undf_symbol = self._symbol_table.symbol_from_tag(
            "undf_{0}".format(fs_name),
            fs=fs_name,
            symbol_type=lfric_psyir.NumberOfUniqueDofsDataSymbol,
            interface=self._read_access)
        self._arglist.append(undf_symbol)

        fs_name = function_space.orig_name
        ndf_symbol = self._symbol_table.symbol_from_tag(
            "ndf_{0}".format(fs_name),
            fs=fs_name,
            symbol_type=lfric_psyir.NumberOfDofsDataSymbol,
            interface=self._read_access)

        dofmap_symbol = self._symbol_table.symbol_from_tag(
            "dofmap_{0}".format(fs_name),
            fs=fs_name,
            symbol_type=lfric_psyir.DofMapDataSymbol,
            dims=[Reference(ndf_symbol)],
            interface=self._read_access)
        self._arglist.append(dofmap_symbol)

    def banded_dofmap(self, function_space, var_accesses=None):
        '''Not implemented.

        :param function_space: the function space for this dofmap.
        :type function_space: :py:class:`psyclone.domain.lfric.FunctionSpace`
        :param var_accesses: an unused optional argument that stores \
            information about variable accesses.
        :type var_accesses: :\
            py:class:`psyclone.core.access_info.VariablesAccessInfo`

        :raises NotImplementedError: as this method is not implemented.

        '''
        raise NotImplementedError("banded_dofmap not implemented")

    def indirection_dofmap(self,
                           function_space,
                           operator=None,
                           var_accesses=None):
        '''Not implemented.

        :param function_space: the function space for this dofmap.
        :type function_space: :py:class:`psyclone.domain.lfric.FunctionSpace`
        :param operator: the CMA operator.
        :type operator: :py:class:`psyclone.dynamo0p3.DynKernelArgument`
        :param var_accesses: an unused optional argument that stores \
            information about variable accesses.
        :type var_accesses: :\
            py:class:`psyclone.core.access_info.VariablesAccessInfo`

        :raises NotImplementedError: as this method is not implemented.

        '''
        raise NotImplementedError("indirection_dofmap not implemented")

    def basis(self, function_space, var_accesses=None):
        '''Create an LFRic basis function argument and add it to the symbol
        table and argument list.

        :param function_space: the function space for this basis function.
        :type function_space: :py:class:`psyclone.domain.lfric.FunctionSpace`
        :param var_accesses: an unused optional argument that stores \
            information about variable accesses.
        :type var_accesses: :\
            py:class:`psyclone.core.access_info.VariablesAccessInfo`

        '''
        basis_name_func = function_space.get_basis_name
        # This import must be placed here to avoid circular dependencies
        # pylint: disable=import-outside-toplevel
        from psyclone.dynamo0p3 import DynBasisFunctions
        first_dim_value_func = DynBasisFunctions.basis_first_dim_value
        self._create_basis(function_space, self.basis_mapping, basis_name_func,
                           first_dim_value_func)

    def diff_basis(self, function_space, var_accesses=None):
        '''Create an LFRic differential basis function argument and add it to
        the symbol table and argument list.

        :param function_space: the function space for this \
            differential basis function.
        :type function_space: :py:class:`psyclone.domain.lfric.FunctionSpace`
        :param var_accesses: an unused optional argument that stores \
            information about variable accesses.
        :type var_accesses: :\
            py:class:`psyclone.core.access_info.VariablesAccessInfo`

        '''
        basis_name_func = function_space.get_diff_basis_name
        # This import must be placed here to avoid circular dependencies
        # pylint: disable=import-outside-toplevel
        from psyclone.dynamo0p3 import DynBasisFunctions
        first_dim_value_func = DynBasisFunctions.diff_basis_first_dim_value
        self._create_basis(function_space, self.diff_basis_mapping,
                           basis_name_func, first_dim_value_func)

    def field_bcs_kernel(self, function_space, var_accesses=None):
        '''Not implemented.

        :param function_space: the function space for this boundary condition.
        :type function_space: :py:class:`psyclone.domain.lfric.FunctionSpace`
        :param var_accesses: an unused optional argument that stores \
            information about variable accesses.
        :type var_accesses: :\
            py:class:`psyclone.core.access_info.VariablesAccessInfo`

        :raises NotImplementedError: as this method is not implemented.

        '''
        raise NotImplementedError("field_bcs_kernel not implemented")

    def operator_bcs_kernel(self, function_space, var_accesses=None):
        '''Not implemented.

        :param function_space: the function space for this bcs kernel
        :type function_space: :py:class:`psyclone.domain.lfric.FunctionSpace`
        :param var_accesses: an unused optional argument that stores \
            information about variable accesses.
        :type var_accesses: :\
            py:class:`psyclone.core.access_info.VariablesAccessInfo`

        :raises NotImplementedError: as this method is not implemented.

        '''
        raise NotImplementedError("operator_bcs_kernel not implemented")

    def ref_element_properties(self, var_accesses=None):
        ''' Properties associated with the reference element

        :param var_accesses: an unused optional argument that stores \
            information about variable accesses.
        :type var_accesses: :\
            py:class:`psyclone.core.access_info.VariablesAccessInfo`

        '''
        # This callback does not contribute any kernel arguments

    def mesh_properties(self, var_accesses=None):
        ''' Properties associated with the mesh

        :param var_accesses: an unused optional argument that stores \
            information about variable accesses.
        :type var_accesses: :\
            py:class:`psyclone.core.access_info.VariablesAccessInfo`

        '''
        # This callback does not contribute any kernel arguments

    def quad_rule(self, var_accesses=None):
        '''Create LFRic arguments associated with the required quadrature, if
        they do not already exist, and add them to the symbol table
        and argument list. The arguments depend on the type of
        quadrature requested.

        :param var_accesses: an unused optional argument that stores \
            information about variable accesses.
        :type var_accesses: :\
            py:class:`psyclone.core.access_info.VariablesAccessInfo`

        :raises InternalError: if an unsupported quadrature shape is \
            found.

        '''
        # The kernel captures all the required quadrature shapes
        for shape in self._kern.qr_rules:
            if shape == "gh_quadrature_xyoz":
                nqp_xy = self._symbol_table.symbol_from_tag(
                    "nqp_xy",
                    symbol_type=lfric_psyir.NumberOfQrPointsInXyDataSymbol,
                    interface=self._read_access)
                nqp_z = self._symbol_table.symbol_from_tag(
                    "nqp_z",
                    symbol_type=lfric_psyir.NumberOfQrPointsInZDataSymbol,
                    interface=self._read_access)
                weights_xy = self._symbol_table.symbol_from_tag(
                    "weights_xy",
                    symbol_type=lfric_psyir.QrWeightsInXyDataSymbol,
                    dims=[Reference(nqp_xy)],
                    interface=self._read_access)
                weights_z = self._symbol_table.symbol_from_tag(
                    "weights_z",
                    symbol_type=lfric_psyir.QrWeightsInZDataSymbol,
                    dims=[Reference(nqp_z)],
                    interface=self._read_access)
                self._arglist.extend([nqp_xy, nqp_z, weights_xy, weights_z])
            elif shape == "gh_quadrature_face":
                nfaces = self._symbol_table.symbol_from_tag(
                    "nfaces",
                    symbol_type=lfric_psyir.NumberOfFacesDataSymbol,
                    interface=self._read_access)
                nqp = self._symbol_table.symbol_from_tag(
                    "nqp_faces",
                    symbol_type=lfric_psyir.NumberOfQrPointsInFacesDataSymbol,
                    interface=self._read_access)
                weights = self._symbol_table.symbol_from_tag(
                    "weights_faces",
                    symbol_type=lfric_psyir.QrWeightsInFacesDataSymbol,
                    dims=[Reference(nqp)],
                    interface=self._read_access)
                self._arglist.extend([nfaces, nqp, weights])
            elif shape == "gh_quadrature_edge":
                nedges = self._symbol_table.symbol_from_tag(
                    "nedges",
                    symbol_type=lfric_psyir.NumberOfEdgesDataSymbol,
                    interface=self._read_access)
                nqp = self._symbol_table.symbol_from_tag(
                    "nqp_edges",
                    symbol_type=lfric_psyir.NumberOfQrPointsInEdgesDataSymbol,
                    interface=self._read_access)
                weights = self._symbol_table.symbol_from_tag(
                    "weights_edges",
                    symbol_type=lfric_psyir.QrWeightsInEdgesDataSymbol,
                    dims=[Reference(nqp)],
                    interface=self._read_access)
                self._arglist.extend([nedges, nqp, weights])
            else:
                raise InternalError("Unsupported quadrature shape '{0}' "
                                    "found in kernel_interface.".format(shape))

    def _create_basis(self, function_space, mapping, basis_name_func,
                      first_dim_value_func):
        '''Internal utility to create an LFRic basis or differential basis
        function argument specific to the particular quadrature that
        is being used and add it to the symbol table and argument
        list. Also declare the associated "ndf" symbol and any
        quadrature-specific symbols if they have not already been
        declared so that they can be used to dimension the basis or
        differential basis symbol.

        This utility function is used to avoid code replication as the
        structure of a basis function is very similar to the structure
        of a differential basis function.

        :param function_space: the function space that this basis or \
            differential basis function is on.
        :type function_space: :py:class:`psyclone.domain.lfric.FunctionSpace`
        :param dict mapping: a mapping from quadrature type to basis \
            or differential basis class name.
        :param method basis_name_func: a method that returns the name \
            of the basis or differential basis function for the \
            current function space.
        :param function first_dim_value_func: a function that returns \
            the size of the first dimension of the basis or \
            differential basis function for the current function \
            space.

        :raises NotImplementedError: if an evaluator shape is found \
            that is not a quadrature shape (currently just \
            'gh_evaluator').
        :raises InternalError: if the supplied evaluator shape is not \
            recognised.

        '''
        # pylint: disable=too-many-locals
        # This import must be placed here to avoid circular dependencies
        # pylint: disable=import-outside-toplevel
        from psyclone.dynamo0p3 import VALID_EVALUATOR_SHAPES
        for shape in self._kern.eval_shapes:
            fs_name = function_space.orig_name
            ndf_symbol = self._symbol_table.symbol_from_tag(
                "ndf_{0}".format(fs_name),
                symbol_type=lfric_psyir.NumberOfDofsDataSymbol,
                fs=fs_name,
                interface=self._read_access)

            # Create the qr tag by appending the last part of the shape
            # name to "qr_".
            quad_name = shape.split("_")[-1]
            basis_tag = basis_name_func(qr_var="qr_" + quad_name)
            if shape == "gh_quadrature_xyoz":
                nqp_xy = self._symbol_table.symbol_from_tag(
                    "nqp_xy",
                    symbol_type=lfric_psyir.NumberOfQrPointsInXyDataSymbol,
                    interface=self._read_access)
                nqp_z = self._symbol_table.symbol_from_tag(
                    "nqp_z",
                    symbol_type=lfric_psyir.NumberOfQrPointsInZDataSymbol,
                    interface=self._read_access)
                arg = mapping["gh_quadrature_xyoz"](
                    basis_tag, [
                        int(first_dim_value_func(function_space)),
                        Reference(ndf_symbol),
                        Reference(nqp_xy),
                        Reference(nqp_z)
                    ],
                    fs_name,
                    interface=self._read_access)
            elif shape == "gh_quadrature_face":
                nfaces = self._symbol_table.symbol_from_tag(
                    "nfaces",
                    symbol_type=lfric_psyir.NumberOfFacesDataSymbol,
                    interface=self._read_access)
                nqp = self._symbol_table.symbol_from_tag(
                    "nqp_faces",
                    symbol_type=lfric_psyir.NumberOfQrPointsInFacesDataSymbol,
                    interface=self._read_access)
                arg = mapping["gh_quadrature_face"](
                    basis_tag, [
                        int(first_dim_value_func(function_space)),
                        Reference(ndf_symbol),
                        Reference(nqp),
                        Reference(nfaces)
                    ],
                    fs_name,
                    interface=self._read_access)
            elif shape == "gh_quadrature_edge":
                nedges = self._symbol_table.symbol_from_tag(
                    "nedges",
                    symbol_type=lfric_psyir.NumberOfEdgesDataSymbol,
                    interface=self._read_access)
                nqp = self._symbol_table.symbol_from_tag(
                    "nqp_edges",
                    symbol_type=lfric_psyir.NumberOfQrPointsInEdgesDataSymbol,
                    interface=self._read_access)
                arg = mapping["gh_quadrature_edge"](
                    basis_tag, [
                        int(first_dim_value_func(function_space)),
                        Reference(ndf_symbol),
                        Reference(nqp),
                        Reference(nedges)
                    ],
                    fs_name,
                    interface=self._read_access)
            elif shape in VALID_EVALUATOR_SHAPES:
                # Need a (diff) basis array for each target space upon
                # which the basis functions have been
                # evaluated. _kern.eval_targets is a dict where the
                # values are 2-tuples of (FunctionSpace, argument).
                for _, _ in self._kern.eval_targets.items():
                    raise NotImplementedError(
                        "Evaluator shapes not implemented in kernel_interface "
                        "class.")
            else:
                raise InternalError(
                    "Unrecognised quadrature or evaluator shape '{0}'. "
                    "Expected one of: {1}.".format(shape,
                                                   VALID_EVALUATOR_SHAPES))
            self._symbol_table.add(arg)
            self._arglist.append(arg)