[OpenACC] declare directive

Hi!

     This patch adds the processing of OpenACC declare directive in C
     and C++. (Note: Support in Fortran is already in trunk.)
     Commentary on the changes is included as an attachment (NOTES).

     All of the code is in the gomp-4_0-branch.

     Regtested on x86_64-linux.

     Thanks!
     Jim
Background
        The declare directive is used to allocate device memory for the
        entire scope of a variable / array within a program, function,
        or subroutine. Consider the following example.

            int a[10];
            #pragma acc declare create (a)

            void func (int *a)
            {
                int b[10];
                #pragma acc declare create (b)

                #pragma acc parallel present (a, b)
                {
                  int i;

                  for (i = 0; i < 10; i++)
                  {
                    b[i] = a[i];
                    a[i] = b[i] + i;
                  }

                }

                return;
            }

            int main (int argc, char **argv)
            {
                func (&a[0]);

                return 0;
            }

        In the example, array 'a' will be allocated on the device at the
        outset of device activity and be available for the duration.
        Whereas, array 'b', will only be available when 'func' is executing.
        In other words, array 'b' will be allocated at the outset of
        execution of 'func' and deallocated at the return from 'func'. In
        some instances, the clause may require that the host copy of a
        variable / array be updated prior to a return from a function or
        subroutine or exiting of the program.

    C and C++ front-ends

        Definitions for use in C and C++ were added to identify the
        declare directive pragma and its' valid clauses. After the
        clauses have been validated, if the declare directive is for a
	global variable, then an attribute is created and chained.
        These attributes will be used during gimplification.

        Once the user-specified clauses have been parsed, the clauses
        have to be examined and potentially altered and/or added to.
        As mentioned in the previous section, with some clauses, e.g.,
        e.g, copy, movement of data has to occur at the entry to 
        something like a function as well as at exit. Hence the need
        to examine/modify/add to the clauses so as to effect the
        correct data movement.

        For all instances of the declare directive, there is at least
        one set of 'entry' clauses. If the clauses pertain to global
        variables, a constructor is created. This constructor will
        'register' the variable(s) / arrays so that at beginning of
        OpenACC runtime the variable / arrays will be allocated and
        be made available throughout program execution.

        If on the other hand, the 'entry' clauses are not found to be
        of a global type, then a node is created and the clauses are
        associated with it. Also note that the 'return' clauses are
        also associated with the node. Notice that there are 'return'
        clauses only for non-global variables / arrays. The clauses
        available for global variables / arrays only allow for data
        movement at the initiation of program execution.

	Middle-end

        The OACC_DECLARE node is handled much the same as other OpenACC
        nodes that represent directives. However, there is one thing
        unique to declare, and that is the handling of the 'return'
        clauses. The 'return' clauses are scanned and then a gimple
        statment is created, but is not added. However, it is saved to
        be added after the body has been gimplified.

        The intent of this last-minute addition is to allow this statement
        to be executed prior to returning from a function. JAKUB: While
        this has been working, I'm not completely sure this is the proper
        means by which to do this in order to guarantee this statement
        is the last one executed. Please advise otherwise.

	Callgraph

        The make_offload functionality has been refactored to handle 
        OpenACC variables / arrays that are not external and and have
        'oacc declare' attributes associated with them. Once 'offloaded',
        the attributes are removed.

	libgomp

        A function has been added to handle the 'registration' of
        global variables that will be allocated on the device. This
        function is called by the constructor which was generated as
        part of handling the declare directive by the compiler.

        An additional function has been added to allocate device memory
        from the list of registered variables / arrays. Another function
        has been added to deallocate that device. These functions are
        called at the appropriate places in the runtime.

        Finally, a function has been added to handle the declare builtin
        which is emitted by the compiler.

    Testing

        New compile and runtime tests have been added. (NOTE: The numbering
        for the runtime tests has a gap in it. These tests use both the
        declare and routine directive. The support for the routine directive
        has yet to be added to trunk, so these tests will appear once
        the support has been committed.)
2015-10-27  James Norris  <jnorris@codesourcery.com>

	gcc/
	* builtin-types.def (BT_FN_VOID_PTR_INT_UINT): New type.
	* c-family/c-common.c (c_common_attribute_table): New oacc_declare.
	* c-family/c-pragma.c (oacc_pragmas): Add entry for declare directive. 
	* c-family/c-pragma.h (enum pragma_kind): Add PRAGMA_OACC_DECLARE.
	(enum pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT.
	* c/c-parser.c (c_parser_pragma): Handle PRAGMA_OACC_DECLARE.
	(c_parser_omp_clause_name): Handle 'device_resident' clause.
	(c_parser_oacc_data_clause): Handle PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT
	and PRAGMA_OMP_CLAUSE_LINK.
	(c_parser_oacc_all_clauses): Handle PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT
	and PRAGMA_OACC_CLAUSE_LINK.
	(OACC_DECLARE_CLAUSE_MASK): New definition.
	(c_parser_oacc_declare): New function.
	* cp/parser.c (cp_parser_omp_clause_name): Handle 'device_resident'
	clause.
	(cp_parser_oacc_data_clause): Handle PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT
	and PRAGMA_OMP_CLAUSE_LINK.
	(OACC_DECLARE_CLAUSE_MASK): New definition.
	(cp_parser_oacc_declare): New function.
	(cp_paser_oacc_all_clauses): Handle PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT
	and PRAGMA_OMP_CLAUSE_LINK.
	(cp_parser_pragma): Handle PRAGMA_OACC_DECLARE.
	* cp/pt.c (tsubst_expr): Handle OACC_DECLARE.
	* fortran/types.def (BT_FN_VOID_PTR_INT_UINT): New type.
	* gimple.h (enum gf_mask): Add GF_OMP_TARGET_KIND_OACC_DECLARE.
	(is_gomple_omp_oacc): Handle GF_OMP_TARGET_KIND_OACC_DECLARE.
	* gimplify.c (struct gimplify_omp_ctx): New field.
	(new_omp_context): Initialize new field.
	(omp_default_clause): Handle device resident variable.
	(gimplify_oacc_declare): New function.
	(device_resident_p): New function.
	(gimplify_expr): Handle OACC_DECLARE.
	(gimplify_body): Handle updating of declare'd variables.
	* omp-builtins.def (BUILT_IN_GOACC_STATIC, BUILT_IN_GOACC_DECLARE):
	New builtins.
	* omp-low.c (expand_omp_target): Handle
	GF_OMP_TARGET_KIND_OACC_DECLARE and BUILTIN_GOACC_DECLARE.
	(lower_omp_target): Handle GF_OMP_TARGET_KIND_OACC_DECLARE,
	GOMP_MAP_DEVICE_RESIDENT and GOMP_MAP_LINK.
	(make_gimple_omp_edges): Handle GF_OMP_TARGET_KIND_OACC_DECLARE.
	* testsuite/c-c++-common/goacc/declare-1.c: New test.
	* testsuite/c-c++-common/goacc/declare-2.c: Likewise.
	* tree.def (OACC_DECLARE): Update operands.
	* tree.h (OACC_DECLARE_RETURN_CLAUSES): New definition.
	* varpool.c (make_offloadable_1, make_offloadable): New functions.
	(get_create): Refactor offload functionality.

	include/
	* gomp-constants.h (enum gomp_map_kind): Add GOMP_MAP_DEVICE_RESIDENT
	and GOMP_MAP_LINK.

	libgomp/
	* libgomp.map (GOACC_2.0): Export GOACC_declare.
	* oacc-init.c (acc_shutdown_1): Add call.
	* oacc-int.h (goacc_allocate_static, goacc_deallocate_static): New
	declarations.
	* oacc-parallel.c (struct oacc_static): New struture.
	(goacc_allocate_static, goacc_deallocate_static, GOACC_register_static,
	GOACC_declare): New functions.
	* testsuite/libgomp.oacc-c-c++-common/declare-1.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/declare-5.c: Likewise.

[OpenACC] declare directive

Commit Message

Comments

Patch