[OpenACC] declare directive

On 10/27/2015 03:18 PM, James Norris wrote:
> 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

     Ping!

     I've revised the patch since I originally submitted it for review
     (https://gcc.gnu.org/ml/gcc-patches/2015-10/msg02967.html). The
     revision is due to Jakub and et al OpenMP 4.5 work in the area of
     'omp declare target'. I now exploit that functionality and have
     revised the patch accordingly.

     Updated ChangeLog, patch, and commentary (NOTES) are attached.

     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. A variable is an OpenACC declare'd
        is not known at the time the varpool node is created. This
        requires that a check of the offloadable bit to determine
        whether make_offloadable should be called to assert the bit if
        the right conditions are met.

	libgomp

        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.
	* tree.def (OACC_DECLARE): Update operands.
	* tree.h (OACC_DECLARE_RETURN_CLAUSES): New definition.
	* varpool.c (make_offloadable): New function.
	(get_create): Refactor offload functionality.

	gcc/testsuite
	* c-c++-common/goacc/declare-1.c: New test.
	* c-c++-common/goacc/declare-2.c: Likewise.

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

	libgomp/

	* libgomp.map (GOACC_2.0.1): Export GOACC_declare.
	* oacc-parallel.c (GOACC_declare): New function.
	* 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