commit c21b0aea867745ea98d9d91cf43c3b2cdf26dd98
Author: Bernd Schmidt <bernds@codesourcery.com>
Date: Tue Nov 18 22:22:26 2014 +0100
OpenACC middle end.
@@ -873,7 +873,8 @@ FIXED_VALUE_H = fixed-value.h $(MACHMODE_H) double-int.h
RTL_H = $(RTL_BASE_H) $(FLAGS_H) genrtl.h
READ_MD_H = $(OBSTACK_H) $(HASHTAB_H) read-md.h
PARAMS_H = params.h params.def
-BUILTINS_DEF = builtins.def sync-builtins.def omp-builtins.def \
+BUILTINS_DEF = builtins.def sync-builtins.def \
+ oacc-builtins.def omp-builtins.def \
gtm-builtins.def sanitizer.def cilkplus.def cilk-builtins.def
INTERNAL_FN_DEF = internal-fn.def
INTERNAL_FN_H = internal-fn.h $(INTERNAL_FN_DEF)
@@ -5341,6 +5341,12 @@ enum c_builtin_type
#define DEF_FUNCTION_TYPE_VAR_4(NAME, RETURN, ARG1, ARG2, ARG3, ARG4) NAME,
#define DEF_FUNCTION_TYPE_VAR_5(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5) \
NAME,
+#define DEF_FUNCTION_TYPE_VAR_8(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
+ ARG6, ARG7, ARG8) \
+ NAME,
+#define DEF_FUNCTION_TYPE_VAR_12(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
+ ARG6, ARG7, ARG8, ARG9, ARG10, ARG11, ARG12) \
+ NAME,
#define DEF_POINTER_TYPE(NAME, TYPE) NAME,
#include "builtin-types.def"
#undef DEF_PRIMITIVE_TYPE
@@ -5359,6 +5365,8 @@ enum c_builtin_type
#undef DEF_FUNCTION_TYPE_VAR_3
#undef DEF_FUNCTION_TYPE_VAR_4
#undef DEF_FUNCTION_TYPE_VAR_5
+#undef DEF_FUNCTION_TYPE_VAR_8
+#undef DEF_FUNCTION_TYPE_VAR_12
#undef DEF_POINTER_TYPE
BT_LAST
};
@@ -5464,6 +5472,14 @@ install_builtin_function_types (void)
def_fn_type (ENUM, RETURN, 1, 4, ARG1, ARG2, ARG3, ARG4);
#define DEF_FUNCTION_TYPE_VAR_5(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5) \
def_fn_type (ENUM, RETURN, 1, 5, ARG1, ARG2, ARG3, ARG4, ARG5);
+#define DEF_FUNCTION_TYPE_VAR_8(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
+ ARG6, ARG7, ARG8) \
+ def_fn_type (ENUM, RETURN, 1, 5, ARG1, ARG2, ARG3, ARG4, ARG5, ARG6, \
+ ARG7, ARG8);
+#define DEF_FUNCTION_TYPE_VAR_12(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
+ ARG6, ARG7, ARG8, ARG9, ARG10, ARG11, ARG12) \
+ def_fn_type (ENUM, RETURN, 1, 5, ARG1, ARG2, ARG3, ARG4, ARG5, ARG6, \
+ ARG7, ARG8, ARG9, ARG10, ARG11, ARG12);
#define DEF_POINTER_TYPE(ENUM, TYPE) \
builtin_types[(int) ENUM] = build_pointer_type (builtin_types[(int) TYPE]);
@@ -5485,6 +5501,8 @@ install_builtin_function_types (void)
#undef DEF_FUNCTION_TYPE_VAR_3
#undef DEF_FUNCTION_TYPE_VAR_4
#undef DEF_FUNCTION_TYPE_VAR_5
+#undef DEF_FUNCTION_TYPE_VAR_8
+#undef DEF_FUNCTION_TYPE_VAR_12
#undef DEF_POINTER_TYPE
builtin_types[(int) BT_LAST] = NULL_TREE;
}
@@ -568,6 +568,8 @@ DEF_FUNCTION_TYPE_VAR_2 (BT_FN_INT_INT_CONST_STRING_VAR,
BT_INT, BT_INT, BT_CONST_STRING)
DEF_FUNCTION_TYPE_VAR_2 (BT_FN_PTR_CONST_PTR_SIZE_VAR, BT_PTR,
BT_CONST_PTR, BT_SIZE)
+DEF_FUNCTION_TYPE_VAR_2 (BT_FN_VOID_INT_INT_VAR, BT_VOID,
+ BT_INT, BT_INT)
DEF_FUNCTION_TYPE_VAR_3 (BT_FN_INT_STRING_SIZE_CONST_STRING_VAR,
BT_INT, BT_STRING, BT_SIZE, BT_CONST_STRING)
@@ -586,6 +588,14 @@ DEF_FUNCTION_TYPE_VAR_5 (BT_FN_INT_STRING_SIZE_INT_SIZE_CONST_STRING_VAR,
DEF_FUNCTION_TYPE_VAR_5 (BT_FN_INT_INT_INT_INT_INT_INT_VAR,
BT_INT, BT_INT, BT_INT, BT_INT, BT_INT, BT_INT)
+DEF_FUNCTION_TYPE_VAR_8 (BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR_INT_INT_VAR,
+ BT_VOID, BT_INT, BT_PTR, BT_SIZE, BT_PTR, BT_PTR,
+ BT_PTR, BT_INT, BT_INT)
+
+DEF_FUNCTION_TYPE_VAR_12 (BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR_INT_INT_INT_INT_INT_VAR,
+ BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_PTR, BT_SIZE, BT_PTR, BT_PTR,
+ BT_PTR, BT_INT, BT_INT, BT_INT, BT_INT, BT_INT)
+
DEF_POINTER_TYPE (BT_PTR_FN_VOID_VAR, BT_FN_VOID_VAR)
DEF_FUNCTION_TYPE_3 (BT_FN_PTR_PTR_FN_VOID_VAR_PTR_SIZE,
BT_PTR, BT_PTR_FN_VOID_VAR, BT_PTR, BT_SIZE)
@@ -5885,6 +5885,49 @@ expand_stack_save (void)
return ret;
}
+
+/* Expand OpenACC acc_on_device.
+
+ This has to happen late (that is, not in early folding; expand_builtin_*,
+ rather than fold_builtin_*), as we have to act differently for host and
+ acceleration device (ACCEL_COMPILER conditional). */
+
+static rtx
+expand_builtin_acc_on_device (tree exp, rtx target ATTRIBUTE_UNUSED)
+{
+ if (!validate_arglist (exp, INTEGER_TYPE, VOID_TYPE))
+ return NULL_RTX;
+
+ tree arg, v1, v2, ret;
+ location_t loc;
+
+ arg = CALL_EXPR_ARG (exp, 0);
+ arg = builtin_save_expr (arg);
+ loc = EXPR_LOCATION (exp);
+
+ /* Build: (arg == v1 || arg == v2) ? 1 : 0. */
+
+#ifdef ACCEL_COMPILER
+ v1 = build_int_cst (TREE_TYPE (arg), /* TODO: acc_device_not_host */ 3);
+ v2 = build_int_cst (TREE_TYPE (arg), ACCEL_COMPILER_acc_device);
+#else
+ v1 = build_int_cst (TREE_TYPE (arg), /* TODO: acc_device_none */ 0);
+ v2 = build_int_cst (TREE_TYPE (arg), /* TODO: acc_device_host */ 2);
+#endif
+
+ v1 = fold_build2_loc (loc, EQ_EXPR, integer_type_node, arg, v1);
+ v2 = fold_build2_loc (loc, EQ_EXPR, integer_type_node, arg, v2);
+
+ /* Can't use TRUTH_ORIF_EXPR, as that is not supported by
+ expand_expr_real*. */
+ ret = fold_build3_loc (loc, COND_EXPR, integer_type_node, v1, v1, v2);
+ ret = fold_build3_loc (loc, COND_EXPR, integer_type_node,
+ ret, integer_one_node, integer_zero_node);
+
+ return expand_normal (ret);
+}
+
+
/* Expand an expression EXP that calls a built-in function,
with result going to TARGET if that's convenient
(and in mode MODE if that's convenient).
@@ -7023,6 +7066,12 @@ expand_builtin (tree exp, rtx target, rtx subtarget, machine_mode mode,
error ("Your target platform does not support -fcheck-pointer-bounds");
break;
+ case BUILT_IN_ACC_ON_DEVICE:
+ target = expand_builtin_acc_on_device (exp, target);
+ if (target)
+ return target;
+ break;
+
default: /* just do library call, if unknown builtin */
break;
}
@@ -13018,6 +13067,7 @@ is_inexpensive_builtin (tree decl)
case BUILT_IN_LABS:
case BUILT_IN_LLABS:
case BUILT_IN_PREFETCH:
+ case BUILT_IN_ACC_ON_DEVICE:
return true;
default:
@@ -146,6 +146,17 @@ along with GCC; see the file COPYING3. If not see
DEF_BUILTIN (ENUM, NAME, BUILT_IN_NORMAL, BT_LAST, BT_LAST, false, false, \
false, ATTR_LAST, false, false)
+/* Builtin used by the implementation of GNU OpenACC. Few of these are
+ actually implemented in the compiler; most are in libgomp. */
+#undef DEF_GOACC_BUILTIN
+#define DEF_GOACC_BUILTIN(ENUM, NAME, TYPE, ATTRS) \
+ DEF_BUILTIN (ENUM, "__builtin_" NAME, BUILT_IN_NORMAL, TYPE, TYPE, \
+ false, true, true, ATTRS, false, flag_openacc)
+#undef DEF_GOACC_BUILTIN_COMPILER
+#define DEF_GOACC_BUILTIN_COMPILER(ENUM, NAME, TYPE, ATTRS) \
+ DEF_BUILTIN (ENUM, "__builtin_" NAME, BUILT_IN_NORMAL, TYPE, TYPE, \
+ flag_openacc, true, true, ATTRS, false, true)
+
/* Builtin used by the implementation of GNU OpenMP. None of these are
actually implemented in the compiler; they're all in libgomp. */
#undef DEF_GOMP_BUILTIN
@@ -894,6 +905,9 @@ DEF_GCC_BUILTIN (BUILT_IN_LINE, "LINE", BT_FN_INT, ATTR_NOTHROW_LEAF_LIST)
/* Synchronization Primitives. */
#include "sync-builtins.def"
+/* OpenACC builtins. */
+#include "oacc-builtins.def"
+
/* OpenMP builtins. */
#include "omp-builtins.def"
@@ -5213,6 +5213,11 @@ enum c_builtin_type
#define DEF_FUNCTION_TYPE_VAR_4(NAME, RETURN, ARG1, ARG2, ARG3, ARG4) NAME,
#define DEF_FUNCTION_TYPE_VAR_5(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5) \
NAME,
+#define DEF_FUNCTION_TYPE_VAR_8(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
+ ARG6, ARG7, ARG8) NAME,
+#define DEF_FUNCTION_TYPE_VAR_12(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
+ ARG6, ARG7, ARG8, ARG9, ARG10, ARG11, \
+ ARG12) NAME,
#define DEF_POINTER_TYPE(NAME, TYPE) NAME,
#include "builtin-types.def"
#undef DEF_PRIMITIVE_TYPE
@@ -5231,6 +5236,8 @@ enum c_builtin_type
#undef DEF_FUNCTION_TYPE_VAR_3
#undef DEF_FUNCTION_TYPE_VAR_4
#undef DEF_FUNCTION_TYPE_VAR_5
+#undef DEF_FUNCTION_TYPE_VAR_8
+#undef DEF_FUNCTION_TYPE_VAR_12
#undef DEF_POINTER_TYPE
BT_LAST
};
@@ -5323,6 +5330,14 @@ c_define_builtins (tree va_list_ref_type_node, tree va_list_arg_type_node)
def_fn_type (ENUM, RETURN, 1, 4, ARG1, ARG2, ARG3, ARG4);
#define DEF_FUNCTION_TYPE_VAR_5(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5) \
def_fn_type (ENUM, RETURN, 1, 5, ARG1, ARG2, ARG3, ARG4, ARG5);
+#define DEF_FUNCTION_TYPE_VAR_8(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
+ ARG6, ARG7, ARG8) \
+ def_fn_type (ENUM, RETURN, 1, 8, ARG1, ARG2, ARG3, ARG4, ARG5, ARG6, \
+ ARG7, ARG8);
+#define DEF_FUNCTION_TYPE_VAR_12(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
+ ARG6, ARG7, ARG8, ARG9, ARG10, ARG11, ARG12) \
+ def_fn_type (ENUM, RETURN, 1, 12, ARG1, ARG2, ARG3, ARG4, ARG5, ARG6, \
+ ARG7, ARG8, ARG9, ARG10, ARG11, ARG12);
#define DEF_POINTER_TYPE(ENUM, TYPE) \
builtin_types[(int) ENUM] = build_pointer_type (builtin_types[(int) TYPE]);
@@ -5344,6 +5359,8 @@ c_define_builtins (tree va_list_ref_type_node, tree va_list_arg_type_node)
#undef DEF_FUNCTION_TYPE_VAR_3
#undef DEF_FUNCTION_TYPE_VAR_4
#undef DEF_FUNCTION_TYPE_VAR_5
+#undef DEF_FUNCTION_TYPE_VAR_8
+#undef DEF_FUNCTION_TYPE_VAR_12
#undef DEF_POINTER_TYPE
builtin_types[(int) BT_LAST] = NULL_TREE;
@@ -173,7 +173,7 @@ along with GCC; see the file COPYING3. If not see
%(linker) %l " LINK_PIE_SPEC "%X %{o*} %{A} %{d} %{e*} %{m} %{N} %{n} %{r}\
%{s} %{t} %{u*} %{x} %{z} %{Z} %{!A:%{!nostdlib:%{!nostartfiles:%S}}}\
%{static:} %{L*} %(mfwrap) %(link_libgcc) %o\
- %{fopenmp|ftree-parallelize-loops=*:%:include(libgomp.spec)%(link_gomp)}\
+ %{fopenacc|fopenmp|ftree-parallelize-loops=*:%:include(libgomp.spec)%(link_gomp)}\
%(mflib)\
%{fprofile-arcs|fprofile-generate|coverage:-lgcov}\
%{!nostdlib:%{!nodefaultlibs:%(link_ssp) %(link_gcc_c_sequence)}}\
@@ -177,7 +177,7 @@ extern GTY(()) int darwin_ms_struct;
%{o*}%{!o:-o a.out} \
%{!nostdlib:%{!nostartfiles:%S}} \
%{L*} %(link_libgcc) %o %{fprofile-arcs|fprofile-generate*|coverage:-lgcov} \
- %{fopenmp|ftree-parallelize-loops=*: \
+ %{fopenacc|fopenmp|ftree-parallelize-loops=*: \
%{static|static-libgcc|static-libstdc++|static-libgfortran: libgomp.a%s; : -lgomp } } \
%{fgnu-tm: \
%{static|static-libgcc|static-libstdc++|static-libgfortran: libitm.a%s; : -litm } } \
@@ -199,7 +199,7 @@ do { \
/* mingw32 uses the -mthreads option to enable thread support. */
#undef GOMP_SELF_SPECS
-#define GOMP_SELF_SPECS "%{fopenmp|ftree-parallelize-loops=*: " \
+#define GOMP_SELF_SPECS "%{fopenacc|fopenmp|ftree-parallelize-loops=*: " \
"-mthreads -pthread}"
#undef GTM_SELF_SPECS
#define GTM_SELF_SPECS "%{fgnu-tm:-mthreads -pthread}"
@@ -92,7 +92,7 @@ do { \
#undef LIB_SPEC
#define LIB_SPEC \
"%{!shared: \
- %{mt|pthread:%{fopenmp|ftree-parallelize-loops=*:-lrt} -lpthread} \
+ %{mt|pthread:%{fopenacc|fopenmp|ftree-parallelize-loops=*:-lrt} -lpthread} \
%{p:%{!mlp64:-L/usr/lib/hpux32/libp} \
%{mlp64:-L/usr/lib/hpux64/libp} -lprof} \
%{pg:%{!mlp64:-L/usr/lib/hpux32/libp} \
@@ -353,4 +353,10 @@ struct GTY(()) machine_function
#define FUNCTION_MODE QImode
#define HAS_INIT_SECTION 1
+/* Support for acc_on_device. */
+
+#include "gomp-constants.h"
+
+#define ACCEL_COMPILER_acc_device GOMP_TARGET_NVIDIA_PTX
+
#endif /* GCC_NVPTX_H */
@@ -122,8 +122,8 @@ along with GCC; see the file COPYING3. If not see
#undef LIB_SPEC
#define LIB_SPEC \
"%{!shared:\
- %{fopenmp|ftree-parallelize-loops=*:%{static:-a archive_shared} -lrt\
- %{static:-a archive}}\
+ %{fopenacc|fopenmp|ftree-parallelize-loops=*:\
+ %{static:-a archive_shared} -lrt %{static:-a archive}}\
%{mt|pthread:-lpthread} -lc\
%{static:%{!nolibdld:-a archive_shared -ldld -a archive -lc}\
%{!mt:%{!pthread:-a shared -lc -a archive}}}}\
@@ -58,22 +58,22 @@ along with GCC; see the file COPYING3. If not see
#if ((TARGET_DEFAULT | TARGET_CPU_DEFAULT) & MASK_GNU_LD)
#define LIB_SPEC \
"%{!shared:\
- %{!p:%{!pg:%{fopenmp|ftree-parallelize-loops=*:%{static:-a shared} -lrt\
- %{static:-a archive}}\
+ %{!p:%{!pg:%{fopenacc|fopenmp|ftree-parallelize-loops=*:\
+ %{static:-a shared} -lrt %{static:-a archive}}\
%{mt|pthread:-lpthread} -lc\
%{static:%{!nolibdld:-a shared -ldld -a archive -lc}\
%{!mt:%{!pthread:-a shared -lc -a archive}}}}}\
%{p:%{!pg:%{static:%{!mhp-ld:-a shared}%{mhp-ld:-a archive_shared}}\
-lprof %{static:-a archive}\
- %{fopenmp|ftree-parallelize-loops=*:%{static:-a shared} -lrt\
- %{static:-a archive}}\
+ %{fopenacc|fopenmp|ftree-parallelize-loops=*:\
+ %{static:-a shared} -lrt %{static:-a archive}}\
%{mt|pthread:-lpthread} -lc\
%{static:%{!nolibdld:-a shared -ldld -a archive -lc}\
%{!mt:%{!pthread:-a shared -lc -a archive}}}}}\
%{pg:%{static:%{!mhp-ld:-a shared}%{mhp-ld:-a archive_shared}}\
-lgprof %{static:-a archive}\
- %{fopenmp|ftree-parallelize-loops=*:%{static:-a shared} -lrt\
- %{static:-a archive}}\
+ %{fopenacc|fopenmp|ftree-parallelize-loops=*:\
+ %{static:-a shared} -lrt %{static:-a archive}}\
%{mt|pthread:-lpthread} -lc\
%{static:%{!nolibdld:-a shared -ldld -a archive -lc}\
%{!mt:%{!pthread:-a shared -lc -a archive}}}}}\
@@ -81,22 +81,22 @@ along with GCC; see the file COPYING3. If not see
#else
#define LIB_SPEC \
"%{!shared:\
- %{!p:%{!pg:%{fopenmp|ftree-parallelize-loops=*:%{static:-a shared} -lrt\
- %{static:-a archive}}\
+ %{!p:%{!pg:%{fopenacc|fopenmp|ftree-parallelize-loops=*:\
+ %{static:-a shared} -lrt %{static:-a archive}}\
%{mt|pthread:-lpthread} -lc\
%{static:%{!nolibdld:-a shared -ldld -a archive -lc}\
%{!mt:%{!pthread:-a shared -lc -a archive}}}}}\
%{p:%{!pg:%{static:%{mgnu-ld:-a shared}%{!mgnu-ld:-a archive_shared}}\
-lprof %{static:-a archive}\
- %{fopenmp|ftree-parallelize-loops=*:%{static:-a shared} -lrt\
- %{static:-a archive}}\
+ %{fopenacc|fopenmp|ftree-parallelize-loops=*:\
+ %{static:-a shared} -lrt %{static:-a archive}}\
%{mt|pthread:-lpthread} -lc\
%{static:%{!nolibdld:-a shared -ldld -a archive -lc}\
%{!mt:%{!pthread:-a shared -lc -a archive}}}}}\
%{pg:%{static:%{mgnu-ld:-a shared}%{!mgnu-ld:-a archive_shared}}\
-lgprof %{static:-a archive}\
- %{fopenmp|ftree-parallelize-loops=*:%{static:-a shared} -lrt\
- %{static:-a archive}}\
+ %{fopenacc|fopenmp|ftree-parallelize-loops=*:\
+ %{static:-a shared} -lrt %{static:-a archive}}\
%{mt|pthread:-lpthread} -lc\
%{static:%{!nolibdld:-a shared -ldld -a archive -lc}\
%{!mt:%{!pthread:-a shared -lc -a archive}}}}}\
@@ -1818,7 +1818,7 @@ There are also several varieties of complex statements.
* Empty Statements::
* Jumps::
* Cleanups::
-* OpenMP::
+* OpenACC and OpenMP::
@end menu
@node Basic Statements
@@ -2047,8 +2047,18 @@ EH lowering pass which runs before most of the optimization passes
eliminates these expressions by explicitly adding the cleanup to each
edge. Rethrowing the exception is represented using @code{RESX_EXPR}.
-@node OpenMP
-@subsection OpenMP
+@node OpenACC and OpenMP
+@subsection OpenACC and OpenMP
+@tindex OACC_CACHE
+@tindex OACC_DATA
+@tindex OACC_DECLARE
+@tindex OACC_ENTER_DATA
+@tindex OACC_EXIT_DATA
+@tindex OACC_HOST_DATA
+@tindex OACC_KERNELS
+@tindex OACC_LOOP
+@tindex OACC_PARALLEL
+@tindex OACC_UPDATE
@tindex OMP_PARALLEL
@tindex OMP_FOR
@tindex OMP_SECTIONS
@@ -2062,10 +2072,54 @@ edge. Rethrowing the exception is represented using @code{RESX_EXPR}.
@tindex OMP_ATOMIC
@tindex OMP_CLAUSE
-All the statements starting with @code{OMP_} represent directives and
-clauses used by the OpenMP API @w{@uref{http://www.openmp.org/}}.
+All the statements starting with @code{OACC_}, and @code{OMP_}
+represent directives and clauses used by the OpenACC API
+@w{@uref{http://www.openacc.org/}}, and OpenMP API
+@w{@uref{http://www.openmp.org/}}, respectively.
@table @code
+@item OACC_CACHE
+
+Represents @code{#pragma acc cache (var @dots{})}.
+
+@item OACC_DATA
+
+Represents @code{#pragma acc data [clause1 @dots{} clauseN]}.
+
+@item OACC_DECLARE
+
+Represents @code{#pragma acc declare [clause1 @dots{} clauseN]}.
+
+@item OACC_ENTER_DATA
+
+Represents @code{#pragma acc enter data [clause1 @dots{} clauseN]}.
+
+@item OACC_EXIT_DATA
+
+Represents @code{#pragma acc exit data [clause1 @dots{} clauseN]}.
+
+@item OACC_HOST_DATA
+
+Represents @code{#pragma acc host_data [clause1 @dots{} clauseN]}.
+
+@item OACC_KERNELS
+
+Represents @code{#pragma acc kernels [clause1 @dots{} clauseN]}.
+
+@item OACC_LOOP
+
+Represents @code{#pragma acc loop [clause1 @dots{} clauseN]}.
+
+See the description of the @code{OMP_FOR} code.
+
+@item OACC_PARALLEL
+
+Represents @code{#pragma acc parallel [clause1 @dots{} clauseN]}.
+
+@item OACC_UPDATE
+
+Represents @code{#pragma acc update [clause1 @dots{} clauseN]}.
+
@item OMP_PARALLEL
Represents @code{#pragma omp parallel [clause1 @dots{} clauseN]}. It
@@ -2093,8 +2147,8 @@ variables.
@item OMP_FOR
-Represents @code{#pragma omp for [clause1 @dots{} clauseN]}. It
-has 5 operands:
+Represents @code{#pragma omp for [clause1 @dots{} clauseN]}. It has
+six operands:
Operand @code{OMP_FOR_BODY} contains the loop body.
@@ -2184,10 +2238,10 @@ building code (@code{omp-low.c}).
@item OMP_CONTINUE
Similarly, this instruction does not represent an OpenMP
-directive, it is used by @code{OMP_FOR} and
+directive, it is used by @code{OMP_FOR} (and similar codes, such as
+@code{OACC_LOOP}) as well as
@code{OMP_SECTIONS} to mark the place where the code needs to
-loop to the next iteration (in the case of @code{OMP_FOR}) or
-the next section (in the case of @code{OMP_SECTIONS}).
+loop to the next iteration, or the next section, respectively.
In some cases, @code{OMP_CONTINUE} is placed right before
@code{OMP_RETURN}. But if there are cleanups that need to
@@ -454,6 +454,8 @@ The following table briefly describes the GIMPLE instruction set.
@item @code{GIMPLE_GOTO} @tab x @tab x
@item @code{GIMPLE_LABEL} @tab x @tab x
@item @code{GIMPLE_NOP} @tab x @tab x
+@item @code{GIMPLE_OACC_KERNELS} @tab x @tab x
+@item @code{GIMPLE_OACC_PARALLEL} @tab x @tab x
@item @code{GIMPLE_OMP_ATOMIC_LOAD} @tab x @tab x
@item @code{GIMPLE_OMP_ATOMIC_STORE} @tab x @tab x
@item @code{GIMPLE_OMP_CONTINUE} @tab x @tab x
@@ -1021,6 +1023,8 @@ Return a deep copy of statement @code{STMT}.
* @code{GIMPLE_EH_FILTER}::
* @code{GIMPLE_LABEL}::
* @code{GIMPLE_NOP}::
+* @code{GIMPLE_OACC_KERNELS}::
+* @code{GIMPLE_OACC_PARALLEL}::
* @code{GIMPLE_OMP_ATOMIC_LOAD}::
* @code{GIMPLE_OMP_ATOMIC_STORE}::
* @code{GIMPLE_OMP_CONTINUE}::
@@ -1666,6 +1670,17 @@ Build a @code{GIMPLE_NOP} statement.
Returns @code{TRUE} if statement @code{G} is a @code{GIMPLE_NOP}.
@end deftypefn
+
+@node @code{GIMPLE_OACC_KERNELS}
+@subsection @code{GIMPLE_OACC_KERNELS}
+@cindex @code{GIMPLE_OACC_KERNELS}
+
+
+@node @code{GIMPLE_OACC_PARALLEL}
+@subsection @code{GIMPLE_OACC_PARALLEL}
+@cindex @code{GIMPLE_OACC_PARALLEL}
+
+
@node @code{GIMPLE_OMP_ATOMIC_LOAD}
@subsection @code{GIMPLE_OMP_ATOMIC_LOAD}
@cindex @code{GIMPLE_OMP_ATOMIC_LOAD}
@@ -1780,9 +1795,8 @@ Set @code{NAME} to be the name associated with @code{OMP} critical statement @co
tree clauses, tree index, tree initial, tree final, tree incr, @
gimple_seq pre_body, enum tree_code omp_for_cond)
Build a @code{GIMPLE_OMP_FOR} statement. @code{BODY} is sequence of statements
-inside the for loop. @code{CLAUSES}, are any of the @code{OMP} loop
-construct's clauses: private, firstprivate, lastprivate,
-reductions, ordered, schedule, and nowait. @code{PRE_BODY} is the
+inside the for loop. @code{CLAUSES}, are any of the loop
+construct's clauses. @code{PRE_BODY} is the
sequence of statements that are loop invariant. @code{INDEX} is the
index variable. @code{INITIAL} is the initial value of @code{INDEX}. @code{FINAL} is
final value of @code{INDEX}. OMP_FOR_COND is the predicate used to
@@ -168,8 +168,8 @@ in the following sections.
@gccoptlist{-ansi -std=@var{standard} -fgnu89-inline @gol
-aux-info @var{filename} -fallow-parameterless-variadic-functions @gol
-fno-asm -fno-builtin -fno-builtin-@var{function} @gol
--fhosted -ffreestanding -fopenmp -fopenmp-simd -fms-extensions @gol
--fplan9-extensions -trigraphs -traditional -traditional-cpp @gol
+-fhosted -ffreestanding -fopenacc -fopenmp -fopenmp-simd @gol
+-fms-extensions -fplan9-extensions -trigraphs -traditional -traditional-cpp @gol
-fallow-single-precision -fcond-mismatch -flax-vector-conversions @gol
-fsigned-bitfields -fsigned-char @gol
-funsigned-bitfields -funsigned-char}
@@ -1872,6 +1872,16 @@ This is equivalent to @option{-fno-hosted}.
@xref{Standards,,Language Standards Supported by GCC}, for details of
freestanding and hosted environments.
+@item -fopenacc
+@opindex fopenacc
+@cindex OpenACC accelerator programming
+Enable handling of OpenACC directives @code{#pragma acc} in C/C++ and
+@code{!$acc} in Fortran. When @option{-fopenacc} is specified, the
+compiler generates accelerated code according to the OpenACC Application
+Programming Interface v2.0 @w{@uref{http://www.openacc.org/}}. This option
+implies @option{-pthread}, and thus is only supported on targets that
+have support for @option{-pthread}.
+
@item -fopenmp
@opindex fopenmp
@cindex OpenMP parallel
@@ -1827,6 +1827,9 @@ Target supports Graphite optimizations.
@item fixed_point
Target supports fixed-point extension to C.
+@item fopenacc
+Target supports OpenACC via @option{-fopenacc}.
+
@item fopenmp
Target supports OpenMP via @option{-fopenmp}.
@@ -666,6 +666,11 @@ gfc_init_builtin_functions (void)
#define DEF_FUNCTION_TYPE_8(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
ARG6, ARG7, ARG8) NAME,
#define DEF_FUNCTION_TYPE_VAR_0(NAME, RETURN) NAME,
+#define DEF_FUNCTION_TYPE_VAR_2(NAME, RETURN, ARG1, ARG2) NAME,
+#define DEF_FUNCTION_TYPE_VAR_8(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
+ ARG6, ARG7, ARG8) NAME,
+#define DEF_FUNCTION_TYPE_VAR_12(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
+ ARG6, ARG7, ARG8, ARG9, ARG10, ARG11, ARG12) NAME,
#define DEF_POINTER_TYPE(NAME, TYPE) NAME,
#include "types.def"
#undef DEF_PRIMITIVE_TYPE
@@ -679,6 +684,9 @@ gfc_init_builtin_functions (void)
#undef DEF_FUNCTION_TYPE_7
#undef DEF_FUNCTION_TYPE_8
#undef DEF_FUNCTION_TYPE_VAR_0
+#undef DEF_FUNCTION_TYPE_VAR_2
+#undef DEF_FUNCTION_TYPE_VAR_8
+#undef DEF_FUNCTION_TYPE_VAR_12
#undef DEF_POINTER_TYPE
BT_LAST
};
@@ -1113,6 +1121,42 @@ gfc_init_builtin_functions (void)
builtin_types[(int) ENUM] \
= build_varargs_function_type_list (builtin_types[(int) RETURN], \
NULL_TREE);
+#define DEF_FUNCTION_TYPE_VAR_2(ENUM, RETURN, ARG1, ARG2) \
+ builtin_types[(int) ENUM] \
+ = build_varargs_function_type_list (builtin_types[(int) RETURN], \
+ builtin_types[(int) ARG1], \
+ builtin_types[(int) ARG2], \
+ NULL_TREE);
+#define DEF_FUNCTION_TYPE_VAR_8(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
+ ARG6, ARG7, ARG8) \
+ builtin_types[(int) ENUM] \
+ = build_varargs_function_type_list (builtin_types[(int) RETURN], \
+ builtin_types[(int) ARG1], \
+ builtin_types[(int) ARG2], \
+ builtin_types[(int) ARG3], \
+ builtin_types[(int) ARG4], \
+ builtin_types[(int) ARG5], \
+ builtin_types[(int) ARG6], \
+ builtin_types[(int) ARG7], \
+ builtin_types[(int) ARG8], \
+ NULL_TREE);
+#define DEF_FUNCTION_TYPE_VAR_12(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
+ ARG6, ARG7, ARG8, ARG9, ARG10, ARG11, ARG12) \
+ builtin_types[(int) ENUM] \
+ = build_varargs_function_type_list (builtin_types[(int) RETURN], \
+ builtin_types[(int) ARG1], \
+ builtin_types[(int) ARG2], \
+ builtin_types[(int) ARG3], \
+ builtin_types[(int) ARG4], \
+ builtin_types[(int) ARG5], \
+ builtin_types[(int) ARG6], \
+ builtin_types[(int) ARG7], \
+ builtin_types[(int) ARG8], \
+ builtin_types[(int) ARG9], \
+ builtin_types[(int) ARG10], \
+ builtin_types[(int) ARG11], \
+ builtin_types[(int) ARG12], \
+ NULL_TREE);
#define DEF_POINTER_TYPE(ENUM, TYPE) \
builtin_types[(int) ENUM] \
= build_pointer_type (builtin_types[(int) TYPE]);
@@ -1128,6 +1172,9 @@ gfc_init_builtin_functions (void)
#undef DEF_FUNCTION_TYPE_7
#undef DEF_FUNCTION_TYPE_8
#undef DEF_FUNCTION_TYPE_VAR_0
+#undef DEF_FUNCTION_TYPE_VAR_2
+#undef DEF_FUNCTION_TYPE_VAR_8
+#undef DEF_FUNCTION_TYPE_VAR_12
#undef DEF_POINTER_TYPE
builtin_types[(int) BT_LAST] = NULL_TREE;
@@ -82,6 +82,7 @@ DEF_FUNCTION_TYPE_0 (BT_FN_VOID, BT_VOID)
DEF_FUNCTION_TYPE_1 (BT_FN_VOID_PTR, BT_VOID, BT_PTR)
DEF_FUNCTION_TYPE_1 (BT_FN_VOID_PTRPTR, BT_VOID, BT_PTR_PTR)
DEF_FUNCTION_TYPE_1 (BT_FN_VOID_VPTR, BT_VOID, BT_VOLATILE_PTR)
+DEF_FUNCTION_TYPE_1 (BT_FN_INT_INT, BT_INT, BT_INT)
DEF_FUNCTION_TYPE_1 (BT_FN_UINT_UINT, BT_UINT, BT_UINT)
DEF_FUNCTION_TYPE_1 (BT_FN_PTR_PTR, BT_PTR, BT_PTR)
DEF_FUNCTION_TYPE_1 (BT_FN_VOID_INT, BT_VOID, BT_INT)
@@ -209,3 +210,13 @@ DEF_FUNCTION_TYPE_8 (BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR,
BT_BOOL, BT_UINT, BT_PTR)
DEF_FUNCTION_TYPE_VAR_0 (BT_FN_VOID_VAR, BT_VOID)
+
+DEF_FUNCTION_TYPE_VAR_2 (BT_FN_VOID_INT_INT_VAR, BT_VOID, BT_INT, BT_INT)
+
+DEF_FUNCTION_TYPE_VAR_8 (BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR_INT_INT_VAR,
+ BT_VOID, BT_INT, BT_PTR, BT_SIZE, BT_PTR, BT_PTR,
+ BT_PTR, BT_INT, BT_INT)
+
+DEF_FUNCTION_TYPE_VAR_12 (BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR_INT_INT_INT_INT_INT_VAR,
+ BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_PTR, BT_SIZE, BT_PTR, BT_PTR,
+ BT_PTR, BT_INT, BT_INT, BT_INT, BT_INT, BT_INT)
@@ -829,7 +829,7 @@ proper position among the other output files. */
"%X %{o*} %{e*} %{N} %{n} %{r}\
%{s} %{t} %{u*} %{z} %{Z} %{!nostdlib:%{!nostartfiles:%S}} " VTABLE_VERIFICATION_SPEC " \
%{static:} %{L*} %(mfwrap) %(link_libgcc) " SANITIZER_EARLY_SPEC " %o\
- %{fopenmp|ftree-parallelize-loops=*:%:include(libgomp.spec)%(link_gomp)}\
+ %{fopenacc|fopenmp|ftree-parallelize-loops=*:%:include(libgomp.spec)%(link_gomp)}\
%{fcilkplus:%:include(libcilkrts.spec)%(link_cilkrts)}\
%{fgnu-tm:%:include(libitm.spec)%(link_itm)}\
%(mflib) " STACK_SPLIT_SPEC "\
@@ -990,7 +990,8 @@ static const char *const multilib_defaults_raw[] = MULTILIB_DEFAULTS;
/* Linking to libgomp implies pthreads. This is particularly important
for targets that use different start files and suchlike. */
#ifndef GOMP_SELF_SPECS
-#define GOMP_SELF_SPECS "%{fopenmp|ftree-parallelize-loops=*: -pthread}"
+#define GOMP_SELF_SPECS "%{fopenacc|fopenmp|ftree-parallelize-loops=*: " \
+ "-pthread}"
#endif
/* Likewise for -fgnu-tm. */
@@ -371,6 +371,8 @@ lower_stmt (gimple_stmt_iterator *gsi, struct lower_data *data)
}
break;
+ case GIMPLE_OACC_KERNELS:
+ case GIMPLE_OACC_PARALLEL:
case GIMPLE_OMP_PARALLEL:
case GIMPLE_OMP_TASK:
case GIMPLE_OMP_TARGET:
@@ -1145,18 +1145,21 @@ dump_gimple_omp_for (pretty_printer *buffer, gomp_for *gs, int spc, int flags)
case GF_OMP_FOR_KIND_FOR:
kind = "";
break;
- case GF_OMP_FOR_KIND_SIMD:
- kind = " simd";
- break;
- case GF_OMP_FOR_KIND_CILKSIMD:
- kind = " cilksimd";
- break;
case GF_OMP_FOR_KIND_DISTRIBUTE:
kind = " distribute";
break;
case GF_OMP_FOR_KIND_CILKFOR:
kind = " _Cilk_for";
break;
+ case GF_OMP_FOR_KIND_OACC_LOOP:
+ kind = " oacc_loop";
+ break;
+ case GF_OMP_FOR_KIND_SIMD:
+ kind = " simd";
+ break;
+ case GF_OMP_FOR_KIND_CILKSIMD:
+ kind = " cilksimd";
+ break;
default:
gcc_unreachable ();
}
@@ -1182,17 +1185,20 @@ dump_gimple_omp_for (pretty_printer *buffer, gomp_for *gs, int spc, int flags)
case GF_OMP_FOR_KIND_FOR:
pp_string (buffer, "#pragma omp for");
break;
+ case GF_OMP_FOR_KIND_DISTRIBUTE:
+ pp_string (buffer, "#pragma omp distribute");
+ break;
+ case GF_OMP_FOR_KIND_CILKFOR:
+ break;
+ case GF_OMP_FOR_KIND_OACC_LOOP:
+ pp_string (buffer, "#pragma acc loop");
+ break;
case GF_OMP_FOR_KIND_SIMD:
pp_string (buffer, "#pragma omp simd");
break;
case GF_OMP_FOR_KIND_CILKSIMD:
pp_string (buffer, "#pragma simd");
break;
- case GF_OMP_FOR_KIND_DISTRIBUTE:
- pp_string (buffer, "#pragma omp distribute");
- break;
- case GF_OMP_FOR_KIND_CILKFOR:
- break;
default:
gcc_unreachable ();
}
@@ -1338,6 +1344,15 @@ dump_gimple_omp_target (pretty_printer *buffer, gomp_target *gs,
case GF_OMP_TARGET_KIND_UPDATE:
kind = " update";
break;
+ case GF_OMP_TARGET_KIND_OACC_DATA:
+ kind = " oacc_data";
+ break;
+ case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ kind = " oacc_enter_exit_data";
+ break;
+ case GF_OMP_TARGET_KIND_OACC_UPDATE:
+ kind = " oacc_update";
+ break;
default:
gcc_unreachable ();
}
@@ -1871,6 +1886,81 @@ dump_gimple_phi (pretty_printer *buffer, gphi *phi, int spc, bool comment,
}
+/* Dump an OpenACC offload tuple on the pretty_printer BUFFER, SPC spaces
+ of indent. FLAGS specifies details to show in the dump (see TDF_* in
+ dumpfile.h). */
+
+static void
+dump_gimple_oacc_offload (pretty_printer *buffer, gimple gs, int spc,
+ int flags)
+{
+ tree (*gimple_omp_clauses) (const_gimple);
+ tree (*gimple_omp_child_fn) (const_gimple);
+ tree (*gimple_omp_data_arg) (const_gimple);
+ const char *kind;
+ switch (gimple_code (gs))
+ {
+ case GIMPLE_OACC_KERNELS:
+ gimple_omp_clauses = gimple_oacc_kernels_clauses;
+ gimple_omp_child_fn = gimple_oacc_kernels_child_fn;
+ gimple_omp_data_arg = gimple_oacc_kernels_data_arg;
+ kind = "kernels";
+ break;
+ case GIMPLE_OACC_PARALLEL:
+ gimple_omp_clauses = gimple_oacc_parallel_clauses;
+ gimple_omp_child_fn = gimple_oacc_parallel_child_fn;
+ gimple_omp_data_arg = gimple_oacc_parallel_data_arg;
+ kind = "parallel";
+ break;
+ default:
+ gcc_unreachable ();
+ }
+ if (flags & TDF_RAW)
+ {
+ dump_gimple_fmt (buffer, spc, flags, "%G <%+BODY <%S>%nCLAUSES <", gs,
+ gimple_omp_body (gs));
+ dump_omp_clauses (buffer, gimple_omp_clauses (gs), spc, flags);
+ dump_gimple_fmt (buffer, spc, flags, " >, %T, %T%n>",
+ gimple_omp_child_fn (gs), gimple_omp_data_arg (gs));
+ }
+ else
+ {
+ gimple_seq body;
+ pp_string (buffer, "#pragma acc ");
+ pp_string (buffer, kind);
+ dump_omp_clauses (buffer, gimple_omp_clauses (gs), spc, flags);
+ if (gimple_omp_child_fn (gs))
+ {
+ pp_string (buffer, " [child fn: ");
+ dump_generic_node (buffer, gimple_omp_child_fn (gs),
+ spc, flags, false);
+ pp_string (buffer, " (");
+ if (gimple_omp_data_arg (gs))
+ dump_generic_node (buffer, gimple_omp_data_arg (gs),
+ spc, flags, false);
+ else
+ pp_string (buffer, "???");
+ pp_string (buffer, ")]");
+ }
+ body = gimple_omp_body (gs);
+ if (body && gimple_code (gimple_seq_first_stmt (body)) != GIMPLE_BIND)
+ {
+ newline_and_indent (buffer, spc + 2);
+ pp_left_brace (buffer);
+ pp_newline (buffer);
+ dump_gimple_seq (buffer, body, spc + 4, flags);
+ newline_and_indent (buffer, spc + 2);
+ pp_right_brace (buffer);
+ }
+ else if (body)
+ {
+ pp_newline (buffer);
+ dump_gimple_seq (buffer, body, spc + 2, flags);
+ }
+ }
+}
+
+
/* Dump a GIMPLE_OMP_PARALLEL tuple on the pretty_printer BUFFER, SPC spaces
of indent. FLAGS specifies details to show in the dump (see TDF_* in
dumpfile.h). */
@@ -2155,6 +2245,11 @@ pp_gimple_stmt_1 (pretty_printer *buffer, gimple gs, int spc, int flags)
dump_gimple_phi (buffer, as_a <gphi *> (gs), spc, false, flags);
break;
+ case GIMPLE_OACC_KERNELS:
+ case GIMPLE_OACC_PARALLEL:
+ dump_gimple_oacc_offload (buffer, gs, spc, flags);
+ break;
+
case GIMPLE_OMP_PARALLEL:
dump_gimple_omp_parallel (buffer, as_a <gomp_parallel *> (gs), spc,
flags);
@@ -305,6 +305,36 @@ walk_gimple_op (gimple stmt, walk_tree_fn callback_op,
return ret;
break;
+ case GIMPLE_OACC_KERNELS:
+ ret = walk_tree (gimple_oacc_kernels_clauses_ptr (stmt), callback_op,
+ wi, pset);
+ if (ret)
+ return ret;
+ ret = walk_tree (gimple_oacc_kernels_child_fn_ptr (stmt), callback_op,
+ wi, pset);
+ if (ret)
+ return ret;
+ ret = walk_tree (gimple_oacc_kernels_data_arg_ptr (stmt), callback_op,
+ wi, pset);
+ if (ret)
+ return ret;
+ break;
+
+ case GIMPLE_OACC_PARALLEL:
+ ret = walk_tree (gimple_oacc_parallel_clauses_ptr (stmt), callback_op,
+ wi, pset);
+ if (ret)
+ return ret;
+ ret = walk_tree (gimple_oacc_parallel_child_fn_ptr (stmt), callback_op,
+ wi, pset);
+ if (ret)
+ return ret;
+ ret = walk_tree (gimple_oacc_parallel_data_arg_ptr (stmt), callback_op,
+ wi, pset);
+ if (ret)
+ return ret;
+ break;
+
case GIMPLE_OMP_CONTINUE:
{
gomp_continue *cont_stmt = as_a <gomp_continue *> (stmt);
@@ -616,6 +646,8 @@ walk_gimple_stmt (gimple_stmt_iterator *gsi, walk_stmt_fn callback_stmt,
return wi->callback_result;
/* FALL THROUGH. */
+ case GIMPLE_OACC_KERNELS:
+ case GIMPLE_OACC_PARALLEL:
case GIMPLE_OMP_CRITICAL:
case GIMPLE_OMP_MASTER:
case GIMPLE_OMP_TASKGROUP:
@@ -833,6 +833,40 @@ gimple_build_debug_source_bind_stat (tree var, tree value,
}
+/* Build a GIMPLE_OACC_KERNELS statement.
+
+ BODY is sequence of statements which are executed as kernels.
+ CLAUSES are the OpenACC kernels construct's clauses. */
+
+gimple
+gimple_build_oacc_kernels (gimple_seq body, tree clauses)
+{
+ gimple p = gimple_alloc (GIMPLE_OACC_KERNELS, 0);
+ if (body)
+ gimple_omp_set_body (p, body);
+ gimple_oacc_kernels_set_clauses (p, clauses);
+
+ return p;
+}
+
+
+/* Build a GIMPLE_OACC_PARALLEL statement.
+
+ BODY is sequence of statements which are executed in parallel.
+ CLAUSES are the OpenACC parallel construct's clauses. */
+
+gimple
+gimple_build_oacc_parallel (gimple_seq body, tree clauses)
+{
+ gimple p = gimple_alloc (GIMPLE_OACC_PARALLEL, 0);
+ if (body)
+ gimple_omp_set_body (p, body);
+ gimple_oacc_parallel_set_clauses (p, clauses);
+
+ return p;
+}
+
+
/* Build a GIMPLE_OMP_CRITICAL statement.
BODY is the sequence of statements for which only one thread can execute.
@@ -854,8 +888,7 @@ gimple_build_omp_critical (gimple_seq body, tree name)
BODY is sequence of statements inside the for loop.
KIND is the `for' variant.
- CLAUSES, are any of the OMP loop construct's clauses: private, firstprivate,
- lastprivate, reductions, ordered, schedule, and nowait.
+ CLAUSES, are any of the loop construct's clauses.
COLLAPSE is the collapse count.
PRE_BODY is the sequence of statements that are loop invariant. */
@@ -1070,7 +1103,8 @@ gimple_build_omp_single (gimple_seq body, tree clauses)
/* Build a GIMPLE_OMP_TARGET statement.
BODY is the sequence of statements that will be executed.
- CLAUSES are any of the OMP target construct's clauses. */
+ KIND is the kind of target region.
+ CLAUSES are any of the construct's clauses. */
gomp_target *
gimple_build_omp_target (gimple_seq body, int kind, tree clauses)
@@ -1737,7 +1771,12 @@ gimple_copy (gimple stmt)
}
break;
+ case GIMPLE_OACC_KERNELS:
+ case GIMPLE_OACC_PARALLEL:
+ gcc_unreachable ();
+
case GIMPLE_OMP_FOR:
+ gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
new_seq = gimple_seq_copy (gimple_omp_for_pre_body (stmt));
gimple_omp_for_set_pre_body (copy, new_seq);
t = unshare_expr (gimple_omp_for_clauses (stmt));
@@ -1811,6 +1850,7 @@ gimple_copy (gimple stmt)
case GIMPLE_OMP_TASKGROUP:
case GIMPLE_OMP_ORDERED:
copy_omp_body:
+ gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
new_seq = gimple_seq_copy (gimple_omp_body (stmt));
gimple_omp_set_body (copy, new_seq);
break;
@@ -205,10 +205,34 @@ DEFGSCODE(GIMPLE_NOP, "gimple_nop", GSS_BASE)
/* IMPORTANT.
- Do not rearrange any of the GIMPLE_OMP_* codes. This ordering is
- exposed by the range check in gimple_omp_subcode(). */
+ Do not rearrange any of the GIMPLE_OACC_* and GIMPLE_OMP_* codes. This
+ ordering is exposed by the range check in gimple_omp_subcode. */
+/* GIMPLE_OACC_KERNELS <BODY, CLAUSES, CHILD_FN, DATA_ARG> represents
+ #pragma acc kernels [CLAUSES]
+ BODY is the sequence of statements inside the kernels construct.
+ CLAUSES is an OMP_CLAUSE chain holding the associated clauses.
+ CHILD_FN is set when outlining the body of the kernels region.
+ All the statements in BODY are moved into this newly created
+ function when converting OMP constructs into low-GIMPLE.
+ DATA_ARG is a vec of 3 local variables in the parent function
+ containing data to be mapped to CHILD_FN. This is used to
+ implement the MAP clauses. */
+DEFGSCODE(GIMPLE_OACC_KERNELS, "gimple_oacc_kernels", GSS_OMP_PARALLEL_LAYOUT)
+
+/* GIMPLE_OACC_PARALLEL <BODY, CLAUSES, CHILD_FN, DATA_ARG> represents
+ #pragma acc parallel [CLAUSES]
+ BODY is the sequence of statements inside the parallel construct.
+ CLAUSES is an OMP_CLAUSE chain holding the associated clauses.
+ CHILD_FN is set when outlining the body of the parallel region.
+ All the statements in BODY are moved into this newly created
+ function when converting OMP constructs into low-GIMPLE.
+ DATA_ARG is a vec of 3 local variables in the parent function
+ containing data to be mapped to CHILD_FN. This is used to
+ implement the MAP clauses. */
+DEFGSCODE(GIMPLE_OACC_PARALLEL, "gimple_oacc_parallel", GSS_OMP_PARALLEL_LAYOUT)
+
/* Tuples used for lowering of OMP_ATOMIC. Although the form of the OMP_ATOMIC
expression is very simple (just in form mem op= expr), various implicit
conversions may cause the expression to become more complex, so that it does
@@ -243,6 +267,9 @@ DEFGSCODE(GIMPLE_OMP_CRITICAL, "gimple_omp_critical", GSS_OMP_CRITICAL)
for (INDEX = INITIAL; INDEX COND FINAL; INDEX {+=,-=} INCR)
BODY
+ Likewise for:
+ #pragma acc loop [clause1 ... clauseN]
+
BODY is the loop body.
CLAUSES is the list of clauses.
@@ -269,7 +296,7 @@ DEFGSCODE(GIMPLE_OMP_CRITICAL, "gimple_omp_critical", GSS_OMP_CRITICAL)
INITIAL, FINAL and INCR are required to be loop invariant integer
expressions that are evaluated without any synchronization.
The evaluation order, frequency of evaluation and side-effects are
- unspecified by the standard. */
+ unspecified by the standards. */
DEFGSCODE(GIMPLE_OMP_FOR, "gimple_omp_for", GSS_OMP_FOR)
/* GIMPLE_OMP_MASTER <BODY> represents #pragma omp master.
@@ -354,6 +381,7 @@ DEFGSCODE(GIMPLE_OMP_SECTIONS_SWITCH, "gimple_omp_sections_switch", GSS_BASE)
DEFGSCODE(GIMPLE_OMP_SINGLE, "gimple_omp_single", GSS_OMP_SINGLE_LAYOUT)
/* GIMPLE_OMP_TARGET <BODY, CLAUSES, CHILD_FN> represents
+ #pragma acc {data,enter data,exit data,update}
#pragma omp target {,data,update}
BODY is the sequence of statements inside the target construct
(NULL for target update).
@@ -89,20 +89,24 @@ enum gf_mask {
GF_CALL_CTRL_ALTERING = 1 << 7,
GF_CALL_WITH_BOUNDS = 1 << 8,
GF_OMP_PARALLEL_COMBINED = 1 << 0,
- GF_OMP_FOR_KIND_MASK = 7 << 0,
+ GF_OMP_FOR_KIND_MASK = (1 << 3) - 1,
GF_OMP_FOR_KIND_FOR = 0,
GF_OMP_FOR_KIND_DISTRIBUTE = 1,
GF_OMP_FOR_KIND_CILKFOR = 2,
+ GF_OMP_FOR_KIND_OACC_LOOP = 3,
/* Flag for SIMD variants of OMP_FOR kinds. */
GF_OMP_FOR_SIMD = 1 << 2,
GF_OMP_FOR_KIND_SIMD = GF_OMP_FOR_SIMD | 0,
GF_OMP_FOR_KIND_CILKSIMD = GF_OMP_FOR_SIMD | 1,
GF_OMP_FOR_COMBINED = 1 << 3,
GF_OMP_FOR_COMBINED_INTO = 1 << 4,
- GF_OMP_TARGET_KIND_MASK = (1 << 2) - 1,
+ GF_OMP_TARGET_KIND_MASK = (1 << 3) - 1,
GF_OMP_TARGET_KIND_REGION = 0,
GF_OMP_TARGET_KIND_DATA = 1,
GF_OMP_TARGET_KIND_UPDATE = 2,
+ GF_OMP_TARGET_KIND_OACC_DATA = 3,
+ GF_OMP_TARGET_KIND_OACC_UPDATE = 4,
+ GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA = 5,
/* True on an GIMPLE_OMP_RETURN statement if the return does not require
a thread synchronization via some sort of barrier. The exact barrier
@@ -552,7 +556,8 @@ struct GTY((tag("GSS_OMP_FOR")))
};
-/* GIMPLE_OMP_PARALLEL, GIMPLE_OMP_TARGET */
+/* GIMPLE_OACC_KERNELS, GIMPLE_OACC_PARALLEL, GIMPLE_OMP_PARALLEL,
+ GIMPLE_OMP_TARGET, GIMPLE_OMP_TASK */
struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
gimple_statement_omp_parallel_layout : public gimple_statement_omp
{
@@ -571,6 +576,22 @@ struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
tree data_arg;
};
+/* GIMPLE_OACC_KERNELS */
+struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
+ gimple_statement_oacc_kernels : public gimple_statement_omp_parallel_layout
+{
+ /* No extra fields; adds invariant:
+ stmt->code == GIMPLE_OACC_KERNELS. */
+};
+
+/* GIMPLE_OACC_PARALLEL */
+struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
+ gimple_statement_oacc_parallel : public gimple_statement_omp_parallel_layout
+{
+ /* No extra fields; adds invariant:
+ stmt->code == GIMPLE_OACC_PARALLEL. */
+};
+
/* GIMPLE_OMP_PARALLEL or GIMPLE_TASK */
struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
gimple_statement_omp_taskreg : public gimple_statement_omp_parallel_layout
@@ -580,7 +601,6 @@ struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
|| stmt->code == GIMPLE_OMP_TASK. */
};
-
/* GIMPLE_OMP_PARALLEL */
struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
gomp_parallel : public gimple_statement_omp_taskreg
@@ -589,6 +609,7 @@ struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
stmt->code == GIMPLE_OMP_PARALLEL. */
};
+/* GIMPLE_OMP_TARGET */
struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
gomp_target : public gimple_statement_omp_parallel_layout
{
@@ -998,6 +1019,22 @@ is_a_helper <gomp_for *>::test (gimple gs)
template <>
template <>
inline bool
+is_a_helper <gimple_statement_oacc_kernels *>::test (gimple gs)
+{
+ return gs->code == GIMPLE_OACC_KERNELS;
+}
+
+template <>
+template <>
+inline bool
+is_a_helper <gimple_statement_oacc_parallel *>::test (gimple gs)
+{
+ return gs->code == GIMPLE_OACC_PARALLEL;
+}
+
+template <>
+template <>
+inline bool
is_a_helper <gimple_statement_omp_taskreg *>::test (gimple gs)
{
return gs->code == GIMPLE_OMP_PARALLEL || gs->code == GIMPLE_OMP_TASK;
@@ -1206,6 +1243,22 @@ is_a_helper <const gomp_for *>::test (const_gimple gs)
template <>
template <>
inline bool
+is_a_helper <const gimple_statement_oacc_kernels *>::test (const_gimple gs)
+{
+ return gs->code == GIMPLE_OACC_KERNELS;
+}
+
+template <>
+template <>
+inline bool
+is_a_helper <const gimple_statement_oacc_parallel *>::test (const_gimple gs)
+{
+ return gs->code == GIMPLE_OACC_PARALLEL;
+}
+
+template <>
+template <>
+inline bool
is_a_helper <const gimple_statement_omp_taskreg *>::test (const_gimple gs)
{
return gs->code == GIMPLE_OMP_PARALLEL || gs->code == GIMPLE_OMP_TASK;
@@ -1336,6 +1389,8 @@ gomp_for *gimple_build_omp_for (gimple_seq, int, tree, size_t, gimple_seq);
gomp_parallel *gimple_build_omp_parallel (gimple_seq, tree, tree, tree);
gomp_task *gimple_build_omp_task (gimple_seq, tree, tree, tree, tree,
tree, tree);
+gimple gimple_build_oacc_kernels (gimple_seq, tree);
+gimple gimple_build_oacc_parallel (gimple_seq, tree);
gimple gimple_build_omp_section (gimple_seq);
gimple gimple_build_omp_master (gimple_seq);
gimple gimple_build_omp_taskgroup (gimple_seq);
@@ -1580,6 +1635,8 @@ gimple_has_substatements (gimple g)
case GIMPLE_EH_FILTER:
case GIMPLE_EH_ELSE:
case GIMPLE_TRY:
+ case GIMPLE_OACC_KERNELS:
+ case GIMPLE_OACC_PARALLEL:
case GIMPLE_OMP_FOR:
case GIMPLE_OMP_MASTER:
case GIMPLE_OMP_TASKGROUP:
@@ -4384,6 +4441,197 @@ gimple_omp_set_body (gimple gs, gimple_seq body)
}
+/* Return the clauses associated with OACC_KERNELS statement GS. */
+
+static inline tree
+gimple_oacc_kernels_clauses (const_gimple gs)
+{
+ const gimple_statement_oacc_kernels *oacc_kernels_stmt =
+ as_a <const gimple_statement_oacc_kernels *> (gs);
+ return oacc_kernels_stmt->clauses;
+}
+
+/* Return a pointer to the clauses associated with OACC_KERNELS statement GS. */
+
+static inline tree *
+gimple_oacc_kernels_clauses_ptr (gimple gs)
+{
+ gimple_statement_oacc_kernels *oacc_kernels_stmt =
+ as_a <gimple_statement_oacc_kernels *> (gs);
+ return &oacc_kernels_stmt->clauses;
+}
+
+/* Set CLAUSES to be the list of clauses associated with OACC_KERNELS statement
+ GS. */
+
+static inline void
+gimple_oacc_kernels_set_clauses (gimple gs, tree clauses)
+{
+ gimple_statement_oacc_kernels *oacc_kernels_stmt =
+ as_a <gimple_statement_oacc_kernels *> (gs);
+ oacc_kernels_stmt->clauses = clauses;
+}
+
+/* Return the child function used to hold the body of OACC_KERNELS statement
+ GS. */
+
+static inline tree
+gimple_oacc_kernels_child_fn (const_gimple gs)
+{
+ const gimple_statement_oacc_kernels *oacc_kernels_stmt =
+ as_a <const gimple_statement_oacc_kernels *> (gs);
+ return oacc_kernels_stmt->child_fn;
+}
+
+/* Return a pointer to the child function used to hold the body of OACC_KERNELS
+ statement GS. */
+
+static inline tree *
+gimple_oacc_kernels_child_fn_ptr (gimple gs)
+{
+ gimple_statement_oacc_kernels *oacc_kernels_stmt =
+ as_a <gimple_statement_oacc_kernels *> (gs);
+ return &oacc_kernels_stmt->child_fn;
+}
+
+/* Set CHILD_FN to be the child function for OACC_KERNELS statement GS. */
+
+static inline void
+gimple_oacc_kernels_set_child_fn (gimple gs, tree child_fn)
+{
+ gimple_statement_oacc_kernels *oacc_kernels_stmt =
+ as_a <gimple_statement_oacc_kernels *> (gs);
+ oacc_kernels_stmt->child_fn = child_fn;
+}
+
+/* Return the artificial argument used to send variables and values
+ from the parent to the children threads in OACC_KERNELS statement GS. */
+
+static inline tree
+gimple_oacc_kernels_data_arg (const_gimple gs)
+{
+ const gimple_statement_oacc_kernels *oacc_kernels_stmt =
+ as_a <const gimple_statement_oacc_kernels *> (gs);
+ return oacc_kernels_stmt->data_arg;
+}
+
+/* Return a pointer to the data argument for OACC_KERNELS statement GS. */
+
+static inline tree *
+gimple_oacc_kernels_data_arg_ptr (gimple gs)
+{
+ gimple_statement_oacc_kernels *oacc_kernels_stmt =
+ as_a <gimple_statement_oacc_kernels *> (gs);
+ return &oacc_kernels_stmt->data_arg;
+}
+
+/* Set DATA_ARG to be the data argument for OACC_KERNELS statement GS. */
+
+static inline void
+gimple_oacc_kernels_set_data_arg (gimple gs, tree data_arg)
+{
+ gimple_statement_oacc_kernels *oacc_kernels_stmt =
+ as_a <gimple_statement_oacc_kernels *> (gs);
+ oacc_kernels_stmt->data_arg = data_arg;
+}
+
+
+/* Return the clauses associated with OACC_PARALLEL statement GS. */
+
+static inline tree
+gimple_oacc_parallel_clauses (const_gimple gs)
+{
+ const gimple_statement_oacc_parallel *oacc_parallel_stmt =
+ as_a <const gimple_statement_oacc_parallel *> (gs);
+ return oacc_parallel_stmt->clauses;
+}
+
+/* Return a pointer to the clauses associated with OACC_PARALLEL statement
+ GS. */
+
+static inline tree *
+gimple_oacc_parallel_clauses_ptr (gimple gs)
+{
+ gimple_statement_oacc_parallel *oacc_parallel_stmt =
+ as_a <gimple_statement_oacc_parallel *> (gs);
+ return &oacc_parallel_stmt->clauses;
+}
+
+/* Set CLAUSES to be the list of clauses associated with OACC_PARALLEL
+ statement GS. */
+
+static inline void
+gimple_oacc_parallel_set_clauses (gimple gs, tree clauses)
+{
+ gimple_statement_oacc_parallel *oacc_parallel_stmt =
+ as_a <gimple_statement_oacc_parallel *> (gs);
+ oacc_parallel_stmt->clauses = clauses;
+}
+
+/* Return the child function used to hold the body of OACC_PARALLEL statement
+ GS. */
+
+static inline tree
+gimple_oacc_parallel_child_fn (const_gimple gs)
+{
+ const gimple_statement_oacc_parallel *oacc_parallel_stmt =
+ as_a <const gimple_statement_oacc_parallel *> (gs);
+ return oacc_parallel_stmt->child_fn;
+}
+
+/* Return a pointer to the child function used to hold the body of
+ OACC_PARALLEL statement GS. */
+
+static inline tree *
+gimple_oacc_parallel_child_fn_ptr (gimple gs)
+{
+ gimple_statement_oacc_parallel *oacc_parallel_stmt =
+ as_a <gimple_statement_oacc_parallel *> (gs);
+ return &oacc_parallel_stmt->child_fn;
+}
+
+/* Set CHILD_FN to be the child function for OACC_PARALLEL statement GS. */
+
+static inline void
+gimple_oacc_parallel_set_child_fn (gimple gs, tree child_fn)
+{
+ gimple_statement_oacc_parallel *oacc_parallel_stmt =
+ as_a <gimple_statement_oacc_parallel *> (gs);
+ oacc_parallel_stmt->child_fn = child_fn;
+}
+
+/* Return the artificial argument used to send variables and values
+ from the parent to the children threads in OACC_PARALLEL statement GS. */
+
+static inline tree
+gimple_oacc_parallel_data_arg (const_gimple gs)
+{
+ const gimple_statement_oacc_parallel *oacc_parallel_stmt =
+ as_a <const gimple_statement_oacc_parallel *> (gs);
+ return oacc_parallel_stmt->data_arg;
+}
+
+/* Return a pointer to the data argument for OACC_PARALLEL statement GS. */
+
+static inline tree *
+gimple_oacc_parallel_data_arg_ptr (gimple gs)
+{
+ gimple_statement_oacc_parallel *oacc_parallel_stmt =
+ as_a <gimple_statement_oacc_parallel *> (gs);
+ return &oacc_parallel_stmt->data_arg;
+}
+
+/* Set DATA_ARG to be the data argument for OACC_PARALLEL statement GS. */
+
+static inline void
+gimple_oacc_parallel_set_data_arg (gimple gs, tree data_arg)
+{
+ gimple_statement_oacc_parallel *oacc_parallel_stmt =
+ as_a <gimple_statement_oacc_parallel *> (gs);
+ oacc_parallel_stmt->data_arg = data_arg;
+}
+
+
/* Return the name associated with OMP_CRITICAL statement CRIT_STMT. */
static inline tree
@@ -4411,7 +4659,7 @@ gimple_omp_critical_set_name (gomp_critical *crit_stmt, tree name)
}
-/* Return the kind of OMP for statemement. */
+/* Return the kind of the OMP_FOR statemement G. */
static inline int
gimple_omp_for_kind (const_gimple g)
@@ -4421,7 +4669,7 @@ gimple_omp_for_kind (const_gimple g)
}
-/* Set the OMP for kind. */
+/* Set the kind of the OMP_FOR statement G. */
static inline void
gimple_omp_for_set_kind (gomp_for *g, int kind)
@@ -4431,7 +4679,7 @@ gimple_omp_for_set_kind (gomp_for *g, int kind)
}
-/* Return true if OMP for statement G has the
+/* Return true if OMP_FOR statement G has the
GF_OMP_FOR_COMBINED flag set. */
static inline bool
@@ -4442,8 +4690,8 @@ gimple_omp_for_combined_p (const_gimple g)
}
-/* Set the GF_OMP_FOR_COMBINED field in G depending on the boolean
- value of COMBINED_P. */
+/* Set the GF_OMP_FOR_COMBINED field in the OMP_FOR statement G depending on
+ the boolean value of COMBINED_P. */
static inline void
gimple_omp_for_set_combined_p (gomp_for *g, bool combined_p)
@@ -4455,7 +4703,7 @@ gimple_omp_for_set_combined_p (gomp_for *g, bool combined_p)
}
-/* Return true if OMP for statement G has the
+/* Return true if the OMP_FOR statement G has the
GF_OMP_FOR_COMBINED_INTO flag set. */
static inline bool
@@ -4466,8 +4714,8 @@ gimple_omp_for_combined_into_p (const_gimple g)
}
-/* Set the GF_OMP_FOR_COMBINED_INTO field in G depending on the boolean
- value of COMBINED_P. */
+/* Set the GF_OMP_FOR_COMBINED_INTO field in the OMP_FOR statement G depending
+ on the boolean value of COMBINED_P. */
static inline void
gimple_omp_for_set_combined_into_p (gomp_for *g, bool combined_p)
@@ -4479,7 +4727,7 @@ gimple_omp_for_set_combined_into_p (gomp_for *g, bool combined_p)
}
-/* Return the clauses associated with OMP_FOR GS. */
+/* Return the clauses associated with the OMP_FOR statement GS. */
static inline tree
gimple_omp_for_clauses (const_gimple gs)
@@ -4489,7 +4737,8 @@ gimple_omp_for_clauses (const_gimple gs)
}
-/* Return a pointer to the OMP_FOR GS. */
+/* Return a pointer to the clauses associated with the OMP_FOR statement
+ GS. */
static inline tree *
gimple_omp_for_clauses_ptr (gimple gs)
@@ -4499,7 +4748,8 @@ gimple_omp_for_clauses_ptr (gimple gs)
}
-/* Set CLAUSES to be the list of clauses associated with OMP_FOR GS. */
+/* Set CLAUSES to be the list of clauses associated with the OMP_FOR statement
+ GS. */
static inline void
gimple_omp_for_set_clauses (gimple gs, tree clauses)
@@ -4509,7 +4759,7 @@ gimple_omp_for_set_clauses (gimple gs, tree clauses)
}
-/* Get the collapse count of OMP_FOR GS. */
+/* Get the collapse count of the OMP_FOR statement GS. */
static inline size_t
gimple_omp_for_collapse (gimple gs)
@@ -4519,7 +4769,7 @@ gimple_omp_for_collapse (gimple gs)
}
-/* Return the index variable for OMP_FOR GS. */
+/* Return the index variable for the OMP_FOR statement GS. */
static inline tree
gimple_omp_for_index (const_gimple gs, size_t i)
@@ -4530,7 +4780,7 @@ gimple_omp_for_index (const_gimple gs, size_t i)
}
-/* Return a pointer to the index variable for OMP_FOR GS. */
+/* Return a pointer to the index variable for the OMP_FOR statement GS. */
static inline tree *
gimple_omp_for_index_ptr (gimple gs, size_t i)
@@ -4541,7 +4791,7 @@ gimple_omp_for_index_ptr (gimple gs, size_t i)
}
-/* Set INDEX to be the index variable for OMP_FOR GS. */
+/* Set INDEX to be the index variable for the OMP_FOR statement GS. */
static inline void
gimple_omp_for_set_index (gimple gs, size_t i, tree index)
@@ -4552,7 +4802,7 @@ gimple_omp_for_set_index (gimple gs, size_t i, tree index)
}
-/* Return the initial value for OMP_FOR GS. */
+/* Return the initial value for the OMP_FOR statement GS. */
static inline tree
gimple_omp_for_initial (const_gimple gs, size_t i)
@@ -4563,7 +4813,7 @@ gimple_omp_for_initial (const_gimple gs, size_t i)
}
-/* Return a pointer to the initial value for OMP_FOR GS. */
+/* Return a pointer to the initial value for the OMP_FOR statement GS. */
static inline tree *
gimple_omp_for_initial_ptr (gimple gs, size_t i)
@@ -4574,7 +4824,7 @@ gimple_omp_for_initial_ptr (gimple gs, size_t i)
}
-/* Set INITIAL to be the initial value for OMP_FOR GS. */
+/* Set INITIAL to be the initial value for the OMP_FOR statement GS. */
static inline void
gimple_omp_for_set_initial (gimple gs, size_t i, tree initial)
@@ -4585,7 +4835,7 @@ gimple_omp_for_set_initial (gimple gs, size_t i, tree initial)
}
-/* Return the final value for OMP_FOR GS. */
+/* Return the final value for the OMP_FOR statement GS. */
static inline tree
gimple_omp_for_final (const_gimple gs, size_t i)
@@ -4596,7 +4846,7 @@ gimple_omp_for_final (const_gimple gs, size_t i)
}
-/* Return a pointer to the final value for OMP_FOR GS. */
+/* Return a pointer to the final value for the OMP_FOR statement GS. */
static inline tree *
gimple_omp_for_final_ptr (gimple gs, size_t i)
@@ -4607,7 +4857,7 @@ gimple_omp_for_final_ptr (gimple gs, size_t i)
}
-/* Set FINAL to be the final value for OMP_FOR GS. */
+/* Set FINAL to be the final value for the OMP_FOR statement GS. */
static inline void
gimple_omp_for_set_final (gimple gs, size_t i, tree final)
@@ -4618,7 +4868,7 @@ gimple_omp_for_set_final (gimple gs, size_t i, tree final)
}
-/* Return the increment value for OMP_FOR GS. */
+/* Return the increment value for the OMP_FOR statement GS. */
static inline tree
gimple_omp_for_incr (const_gimple gs, size_t i)
@@ -4629,7 +4879,7 @@ gimple_omp_for_incr (const_gimple gs, size_t i)
}
-/* Return a pointer to the increment value for OMP_FOR GS. */
+/* Return a pointer to the increment value for the OMP_FOR statement GS. */
static inline tree *
gimple_omp_for_incr_ptr (gimple gs, size_t i)
@@ -4640,7 +4890,7 @@ gimple_omp_for_incr_ptr (gimple gs, size_t i)
}
-/* Set INCR to be the increment value for OMP_FOR GS. */
+/* Set INCR to be the increment value for the OMP_FOR statement GS. */
static inline void
gimple_omp_for_set_incr (gimple gs, size_t i, tree incr)
@@ -4715,8 +4965,9 @@ gimple_omp_parallel_set_clauses (gomp_parallel *omp_parallel_stmt,
/* Return the child function used to hold the body of OMP_PARALLEL_STMT. */
static inline tree
-gimple_omp_parallel_child_fn (const gomp_parallel *omp_parallel_stmt)
+gimple_omp_parallel_child_fn (const_gimple gs)
{
+ const gomp_parallel *omp_parallel_stmt = as_a <const gomp_parallel *> (gs);
return omp_parallel_stmt->child_fn;
}
@@ -5134,8 +5385,9 @@ gimple_omp_target_set_kind (gomp_target *g, int kind)
/* Return the child function used to hold the body of OMP_TARGET_STMT. */
static inline tree
-gimple_omp_target_child_fn (const gomp_target *omp_target_stmt)
+gimple_omp_target_child_fn (const_gimple gs)
{
+ const gomp_target *omp_target_stmt = as_a <const gomp_target *> (gs);
return omp_target_stmt->child_fn;
}
@@ -5152,9 +5404,10 @@ gimple_omp_target_child_fn_ptr (gomp_target *omp_target_stmt)
/* Set CHILD_FN to be the child function for OMP_TARGET_STMT. */
static inline void
-gimple_omp_target_set_child_fn (gomp_target *omp_target_stmt,
+gimple_omp_target_set_child_fn (gimple gs,
tree child_fn)
{
+ gomp_target *omp_target_stmt = as_a <gomp_target *> (gs);
omp_target_stmt->child_fn = child_fn;
}
@@ -5163,8 +5416,9 @@ gimple_omp_target_set_child_fn (gomp_target *omp_target_stmt,
from the parent to the children threads in OMP_TARGET_STMT. */
static inline tree
-gimple_omp_target_data_arg (const gomp_target *omp_target_stmt)
+gimple_omp_target_data_arg (const_gimple gs)
{
+ const gomp_target *omp_target_stmt = as_a <const gomp_target *> (gs);
return omp_target_stmt->data_arg;
}
@@ -5181,9 +5435,9 @@ gimple_omp_target_data_arg_ptr (gomp_target *omp_target_stmt)
/* Set DATA_ARG to be the data argument for OMP_TARGET_STMT. */
static inline void
-gimple_omp_target_set_data_arg (gomp_target *omp_target_stmt,
- tree data_arg)
+gimple_omp_target_set_data_arg (gimple gs, tree data_arg)
{
+ gomp_target *omp_target_stmt = as_a <gomp_target *> (gs);
omp_target_stmt->data_arg = data_arg;
}
@@ -5552,6 +5806,8 @@ gimple_return_set_retbnd (gimple gs, tree retval)
/* Returns true when the gimple statement STMT is any of the OpenMP types. */
#define CASE_GIMPLE_OMP \
+ case GIMPLE_OACC_KERNELS: \
+ case GIMPLE_OACC_PARALLEL: \
case GIMPLE_OMP_PARALLEL: \
case GIMPLE_OMP_TASK: \
case GIMPLE_OMP_FOR: \
@@ -5582,6 +5838,65 @@ is_gimple_omp (const_gimple stmt)
}
}
+/* Return true if STMT is any of the OpenACC types specifically. */
+
+static inline bool
+is_gimple_omp_oacc_specifically (const_gimple stmt)
+{
+ gcc_assert (is_gimple_omp (stmt));
+ switch (gimple_code (stmt))
+ {
+ case GIMPLE_OACC_KERNELS:
+ case GIMPLE_OACC_PARALLEL:
+ return true;
+ case GIMPLE_OMP_FOR:
+ switch (gimple_omp_for_kind (stmt))
+ {
+ case GF_OMP_FOR_KIND_OACC_LOOP:
+ return true;
+ default:
+ return false;
+ }
+ case GIMPLE_OMP_TARGET:
+ switch (gimple_omp_target_kind (stmt))
+ {
+ case GF_OMP_TARGET_KIND_OACC_DATA:
+ case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ case GF_OMP_TARGET_KIND_OACC_UPDATE:
+ return true;
+ default:
+ return false;
+ }
+ default:
+ return false;
+ }
+}
+
+
+/* Return true if OMP_* STMT is offloaded. */
+
+static inline bool
+is_gimple_omp_offloaded (const_gimple stmt)
+{
+ gcc_assert (is_gimple_omp (stmt));
+ switch (gimple_code (stmt))
+ {
+ case GIMPLE_OACC_KERNELS:
+ case GIMPLE_OACC_PARALLEL:
+ return true;
+ case GIMPLE_OMP_TARGET:
+ switch (gimple_omp_target_kind (stmt))
+ {
+ case GF_OMP_TARGET_KIND_REGION:
+ return true;
+ default:
+ return false;
+ }
+ default:
+ return false;
+ }
+}
+
/* Returns TRUE if statement G is a GIMPLE_NOP. */
@@ -90,7 +90,11 @@ enum gimplify_omp_var_data
GOVD_PRIVATE_OUTER_REF = 1024,
GOVD_LINEAR = 2048,
GOVD_ALIGNED = 4096,
+
+ /* Flags for GOVD_MAP. */
+ /* Don't copy back. */
GOVD_MAP_TO_ONLY = 8192,
+
GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
| GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
| GOVD_LOCAL)
@@ -99,15 +103,16 @@ enum gimplify_omp_var_data
enum omp_region_type
{
- ORT_WORKSHARE = 0,
- ORT_SIMD = 1,
- ORT_PARALLEL = 2,
- ORT_COMBINED_PARALLEL = 3,
- ORT_TASK = 4,
- ORT_UNTIED_TASK = 5,
- ORT_TEAMS = 8,
- ORT_TARGET_DATA = 16,
- ORT_TARGET = 32
+ /* An undefined region type. */
+ ORT_INVALID = 0,
+
+ ORT_WORKSHARE,
+ ORT_SIMD,
+ ORT_PARALLEL,
+ ORT_COMBINED_PARALLEL,
+ ORT_TASK,
+ ORT_TEAMS,
+ ORT_TARGET
};
/* Gimplify hashtable helper. */
@@ -149,6 +154,21 @@ struct gimplify_omp_ctx
location_t location;
enum omp_clause_default_kind default_kind;
enum omp_region_type region_type;
+ union
+ {
+ /* ORT_TASK. */
+ struct
+ {
+ /* Has an untied clause. */
+ unsigned untied : 1;
+ } task;
+ /* ORT_TARGET. */
+ struct
+ {
+ /* Prepare this region for offloading. */
+ unsigned offload : 1;
+ } target;
+ } region_type_flags;
bool combined_loop;
bool distribute;
};
@@ -356,7 +376,7 @@ splay_tree_compare_decl_uid (splay_tree_key xa, splay_tree_key xb)
/* Create a new omp construct that deals with variable remapping. */
static struct gimplify_omp_ctx *
-new_omp_context (enum omp_region_type region_type)
+new_omp_context (void)
{
struct gimplify_omp_ctx *c;
@@ -365,11 +385,8 @@ new_omp_context (enum omp_region_type region_type)
c->variables = splay_tree_new (splay_tree_compare_decl_uid, 0, 0);
c->privatized_types = new hash_set<tree>;
c->location = input_location;
- c->region_type = region_type;
- if ((region_type & ORT_TASK) == 0)
- c->default_kind = OMP_CLAUSE_DEFAULT_SHARED;
- else
- c->default_kind = OMP_CLAUSE_DEFAULT_UNSPECIFIED;
+ c->default_kind = OMP_CLAUSE_DEFAULT_SHARED;
+ c->region_type = ORT_INVALID;
return c;
}
@@ -1544,9 +1561,10 @@ gimplify_case_label_expr (tree *expr_p, gimple_seq *pre_p)
struct gimplify_ctx *ctxp;
glabel *label_stmt;
- /* Invalid OpenMP programs can play Duff's Device type games with
+ /* Invalid programs can play Duff's Device type games with, for example,
#pragma omp parallel. At least in the C front end, we don't
- detect such invalid branches until after gimplification. */
+ detect such invalid branches until after gimplification, in the
+ diagnose_omp_blocks pass. */
for (ctxp = gimplify_ctxp; ; ctxp = ctxp->prev_context)
if (ctxp->case_labels.exists ())
break;
@@ -2244,7 +2262,7 @@ gimplify_arg (tree *arg_p, gimple_seq *pre_p, location_t call_location)
return gimplify_expr (arg_p, pre_p, NULL, test, fb);
}
-/* Don't fold STMT inside ORT_TARGET, because it can break code by adding decl
+/* Don't fold inside offloading regsion: it can break code by adding decl
references that weren't in the source. We'll do it during omplower pass
instead. */
@@ -2253,7 +2271,8 @@ maybe_fold_stmt (gimple_stmt_iterator *gsi)
{
struct gimplify_omp_ctx *ctx;
for (ctx = gimplify_omp_ctxp; ctx; ctx = ctx->outer_context)
- if (ctx->region_type == ORT_TARGET)
+ if (ctx->region_type == ORT_TARGET
+ && ctx->region_type_flags.target.offload)
return false;
return fold_stmt (gsi);
}
@@ -4435,11 +4454,21 @@ is_gimple_stmt (tree t)
case CATCH_EXPR:
case ASM_EXPR:
case STATEMENT_LIST:
+ case OACC_PARALLEL:
+ case OACC_KERNELS:
+ case OACC_DATA:
+ case OACC_HOST_DATA:
+ case OACC_DECLARE:
+ case OACC_UPDATE:
+ case OACC_ENTER_DATA:
+ case OACC_EXIT_DATA:
+ case OACC_CACHE:
case OMP_PARALLEL:
case OMP_FOR:
case OMP_SIMD:
case CILK_SIMD:
case OMP_DISTRIBUTE:
+ case OACC_LOOP:
case OMP_SECTIONS:
case OMP_SECTION:
case OMP_SINGLE:
@@ -5499,10 +5528,12 @@ omp_firstprivatize_variable (struct gimplify_omp_ctx *ctx, tree decl)
return;
}
else if (ctx->region_type == ORT_TARGET)
- omp_add_variable (ctx, decl, GOVD_MAP | GOVD_MAP_TO_ONLY);
+ {
+ if (ctx->region_type_flags.target.offload)
+ omp_add_variable (ctx, decl, GOVD_MAP | GOVD_MAP_TO_ONLY);
+ }
else if (ctx->region_type != ORT_WORKSHARE
- && ctx->region_type != ORT_SIMD
- && ctx->region_type != ORT_TARGET_DATA)
+ && ctx->region_type != ORT_SIMD)
omp_add_variable (ctx, decl, GOVD_FIRSTPRIVATE);
ctx = ctx->outer_context;
@@ -5611,9 +5642,12 @@ omp_add_variable (struct gimplify_omp_ctx *ctx, tree decl, unsigned int flags)
copy into or out of the context. */
if (!(flags & GOVD_LOCAL))
{
- nflags = flags & GOVD_MAP
- ? GOVD_MAP | GOVD_MAP_TO_ONLY | GOVD_EXPLICIT
- : flags & GOVD_PRIVATE ? GOVD_PRIVATE : GOVD_FIRSTPRIVATE;
+ if (flags & GOVD_MAP)
+ nflags = GOVD_MAP | GOVD_MAP_TO_ONLY | GOVD_EXPLICIT;
+ else if (flags & GOVD_PRIVATE)
+ nflags = GOVD_PRIVATE;
+ else
+ nflags = GOVD_FIRSTPRIVATE;
nflags |= flags & GOVD_SEEN;
t = DECL_VALUE_EXPR (decl);
gcc_assert (TREE_CODE (t) == INDIRECT_REF);
@@ -5680,7 +5714,8 @@ omp_notice_threadprivate_variable (struct gimplify_omp_ctx *ctx, tree decl,
struct gimplify_omp_ctx *octx;
for (octx = ctx; octx; octx = octx->outer_context)
- if (octx->region_type == ORT_TARGET)
+ if (octx->region_type == ORT_TARGET
+ && octx->region_type_flags.target.offload)
{
n = splay_tree_lookup (octx->variables, (splay_tree_key)decl);
if (n == NULL)
@@ -5694,7 +5729,7 @@ omp_notice_threadprivate_variable (struct gimplify_omp_ctx *ctx, tree decl,
splay_tree_insert (octx->variables, (splay_tree_key)decl2, 0);
}
- if (ctx->region_type != ORT_UNTIED_TASK)
+ if (!(ctx->region_type == ORT_TASK && ctx->region_type_flags.task.untied))
return false;
n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
if (n == NULL)
@@ -5741,7 +5776,8 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
}
n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
- if (ctx->region_type == ORT_TARGET)
+ if (ctx->region_type == ORT_TARGET
+ && ctx->region_type_flags.target.offload)
{
ret = lang_hooks.decls.omp_disregard_value_expr (decl, true);
if (n == NULL)
@@ -5772,7 +5808,8 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
if (ctx->region_type == ORT_WORKSHARE
|| ctx->region_type == ORT_SIMD
- || ctx->region_type == ORT_TARGET_DATA)
+ || (ctx->region_type == ORT_TARGET
+ && !ctx->region_type_flags.target.offload))
goto do_outer;
/* ??? Some compiler-generated variables (like SAVE_EXPRs) could be
@@ -5786,13 +5823,14 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
switch (default_kind)
{
case OMP_CLAUSE_DEFAULT_NONE:
- if ((ctx->region_type & ORT_PARALLEL) != 0)
+ if (ctx->region_type == ORT_PARALLEL
+ || ctx->region_type == ORT_COMBINED_PARALLEL)
{
error ("%qE not specified in enclosing parallel",
DECL_NAME (lang_hooks.decls.omp_report_decl (decl)));
error_at (ctx->location, "enclosing parallel");
}
- else if ((ctx->region_type & ORT_TASK) != 0)
+ else if (ctx->region_type == ORT_TASK)
{
error ("%qE not specified in enclosing task",
DECL_NAME (lang_hooks.decls.omp_report_decl (decl)));
@@ -5818,14 +5856,14 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
break;
case OMP_CLAUSE_DEFAULT_UNSPECIFIED:
/* decl will be either GOVD_FIRSTPRIVATE or GOVD_SHARED. */
- gcc_assert ((ctx->region_type & ORT_TASK) != 0);
+ gcc_assert (ctx->region_type == ORT_TASK);
if (ctx->outer_context)
omp_notice_variable (ctx->outer_context, decl, in_code);
for (octx = ctx->outer_context; octx; octx = octx->outer_context)
{
splay_tree_node n2;
- if ((octx->region_type & (ORT_TARGET_DATA | ORT_TARGET)) != 0)
+ if (octx->region_type == ORT_TARGET)
continue;
n2 = splay_tree_lookup (octx->variables, (splay_tree_key) decl);
if (n2 && (n2->value & GOVD_DATA_SHARE_CLASS) != GOVD_SHARED)
@@ -5833,7 +5871,9 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
flags |= GOVD_FIRSTPRIVATE;
break;
}
- if ((octx->region_type & (ORT_PARALLEL | ORT_TEAMS)) != 0)
+ if (octx->region_type == ORT_PARALLEL
+ || octx->region_type == ORT_COMBINED_PARALLEL
+ || octx->region_type == ORT_TEAMS)
break;
}
if (flags & GOVD_FIRSTPRIVATE)
@@ -5978,7 +6018,7 @@ omp_check_private (struct gimplify_omp_ctx *ctx, tree decl, bool copyprivate)
|| (!copyprivate
&& lang_hooks.decls.omp_privatize_by_reference (decl)));
- if ((ctx->region_type & (ORT_TARGET | ORT_TARGET_DATA)) != 0)
+ if (ctx->region_type == ORT_TARGET)
continue;
n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
@@ -5990,19 +6030,17 @@ omp_check_private (struct gimplify_omp_ctx *ctx, tree decl, bool copyprivate)
return false;
}
-/* Scan the OpenMP clauses in *LIST_P, installing mappings into a new
- and previous omp contexts. */
+/* Scan the clauses in *LIST_P, installing mappings into CTX as well as outer
+ contexts, if applicable. Before returning, CTX will also be pushed to the
+ top of GIMPLIFY_OMP_CTXP. */
static void
gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
- enum omp_region_type region_type)
+ struct gimplify_omp_ctx *ctx)
{
- struct gimplify_omp_ctx *ctx, *outer_ctx;
+ struct gimplify_omp_ctx *outer_ctx = ctx->outer_context;
tree c;
- ctx = new_omp_context (region_type);
- outer_ctx = ctx->outer_context;
-
while ((c = *list_p) != NULL)
{
bool remove = false;
@@ -6101,6 +6139,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
+ case OMP_CLAUSE__CACHE_:
decl = OMP_CLAUSE_DECL (c);
if (error_operand_p (decl))
{
@@ -6244,7 +6283,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
if (outer_ctx)
omp_notice_variable (outer_ctx, decl, true);
if (check_non_private
- && region_type == ORT_WORKSHARE
+ && ctx->region_type == ORT_WORKSHARE
&& omp_check_private (ctx, decl, false))
{
error ("%s variable %qE is private in outer context",
@@ -6266,11 +6305,25 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
case OMP_CLAUSE_DIST_SCHEDULE:
case OMP_CLAUSE_DEVICE:
case OMP_CLAUSE__CILK_FOR_COUNT_:
+ case OMP_CLAUSE_NUM_GANGS:
+ case OMP_CLAUSE_NUM_WORKERS:
+ case OMP_CLAUSE_VECTOR_LENGTH:
if (gimplify_expr (&OMP_CLAUSE_OPERAND (c, 0), pre_p, NULL,
is_gimple_val, fb_rvalue) == GS_ERROR)
remove = true;
break;
+ case OMP_CLAUSE_DEVICE_RESIDENT:
+ case OMP_CLAUSE_USE_DEVICE:
+ case OMP_CLAUSE_GANG:
+ case OMP_CLAUSE_ASYNC:
+ case OMP_CLAUSE_WAIT:
+ case OMP_CLAUSE_INDEPENDENT:
+ case OMP_CLAUSE_WORKER:
+ case OMP_CLAUSE_VECTOR:
+ remove = true;
+ break;
+
case OMP_CLAUSE_NOWAIT:
case OMP_CLAUSE_ORDERED:
case OMP_CLAUSE_UNTIED:
@@ -6395,9 +6448,12 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
OMP_CLAUSE_PRIVATE_OUTER_REF (clause) = 1;
else if (code == OMP_CLAUSE_MAP)
{
- OMP_CLAUSE_MAP_KIND (clause) = flags & GOVD_MAP_TO_ONLY
- ? OMP_CLAUSE_MAP_TO
- : OMP_CLAUSE_MAP_TOFROM;
+ enum omp_clause_map_kind map_kind;
+ map_kind = (flags & GOVD_MAP_TO_ONLY
+ ? OMP_CLAUSE_MAP_TO
+ : OMP_CLAUSE_MAP_TOFROM);
+ OMP_CLAUSE_MAP_KIND (clause) = map_kind;
+
if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
{
@@ -6445,6 +6501,9 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
return 0;
}
+/* Apply adjustments to the clauses in *LIST_P. Before returning, the current
+ context will also be destroyed, and popped off of GIMPLIFY_OMP_CTXP. */
+
static void
gimplify_adjust_omp_clauses (gimple_seq *pre_p, tree *list_p)
{
@@ -6564,12 +6623,20 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, tree *list_p)
if (!DECL_P (decl))
break;
n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
- if (ctx->region_type == ORT_TARGET && !(n->value & GOVD_SEEN))
+ if (ctx->region_type == ORT_TARGET
+ && ctx->region_type_flags.target.offload
+ && !(n->value & GOVD_SEEN))
remove = true;
else if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST
&& OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_POINTER)
{
+ /* For OMP_CLAUSE_MAP_FORCE_DEVICEPTR, we'll never enter here,
+ because for these, TREE_CODE (DECL_SIZE (decl)) will always be
+ INTEGER_CST. */
+ gcc_assert (OMP_CLAUSE_MAP_KIND (c)
+ != OMP_CLAUSE_MAP_FORCE_DEVICEPTR);
+
tree decl2 = DECL_VALUE_EXPR (decl);
gcc_assert (TREE_CODE (decl2) == INDIRECT_REF);
decl2 = TREE_OPERAND (decl2, 0);
@@ -6598,6 +6665,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, tree *list_p)
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
+ case OMP_CLAUSE__CACHE_:
decl = OMP_CLAUSE_DECL (c);
if (!DECL_P (decl))
break;
@@ -6643,8 +6711,19 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, tree *list_p)
case OMP_CLAUSE_SAFELEN:
case OMP_CLAUSE_DEPEND:
case OMP_CLAUSE__CILK_FOR_COUNT_:
+ case OMP_CLAUSE_NUM_GANGS:
+ case OMP_CLAUSE_NUM_WORKERS:
+ case OMP_CLAUSE_VECTOR_LENGTH:
break;
+ case OMP_CLAUSE_DEVICE_RESIDENT:
+ case OMP_CLAUSE_USE_DEVICE:
+ case OMP_CLAUSE_GANG:
+ case OMP_CLAUSE_ASYNC:
+ case OMP_CLAUSE_WAIT:
+ case OMP_CLAUSE_INDEPENDENT:
+ case OMP_CLAUSE_WORKER:
+ case OMP_CLAUSE_VECTOR:
default:
gcc_unreachable ();
}
@@ -6665,6 +6744,24 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, tree *list_p)
delete_omp_context (ctx);
}
+/* Gimplify OACC_CACHE. */
+
+static void
+gimplify_oacc_cache (tree *expr_p, gimple_seq *pre_p)
+{
+ tree expr = *expr_p;
+ struct gimplify_omp_ctx *ctx = new_omp_context ();
+
+ ctx->region_type = ORT_WORKSHARE;
+
+ gimplify_scan_omp_clauses (&OACC_CACHE_CLAUSES (expr), pre_p, ctx);
+ gimplify_adjust_omp_clauses (pre_p, &OACC_CACHE_CLAUSES (expr));
+
+ /* TODO: Do something sensible with this information. */
+
+ *expr_p = NULL_TREE;
+}
+
/* Gimplify the contents of an OMP_PARALLEL statement. This involves
gimplification of the body, as well as scanning the body for used
variables. We need to do this scan now, because variable-sized
@@ -6676,11 +6773,12 @@ gimplify_omp_parallel (tree *expr_p, gimple_seq *pre_p)
tree expr = *expr_p;
gimple g;
gimple_seq body = NULL;
+ struct gimplify_omp_ctx *ctx = new_omp_context ();
+
+ ctx->region_type
+ = OMP_PARALLEL_COMBINED (expr) ? ORT_COMBINED_PARALLEL : ORT_PARALLEL;
- gimplify_scan_omp_clauses (&OMP_PARALLEL_CLAUSES (expr), pre_p,
- OMP_PARALLEL_COMBINED (expr)
- ? ORT_COMBINED_PARALLEL
- : ORT_PARALLEL);
+ gimplify_scan_omp_clauses (&OMP_PARALLEL_CLAUSES (expr), pre_p, ctx);
push_gimplify_context ();
@@ -6712,11 +6810,14 @@ gimplify_omp_task (tree *expr_p, gimple_seq *pre_p)
tree expr = *expr_p;
gimple g;
gimple_seq body = NULL;
+ struct gimplify_omp_ctx *ctx = new_omp_context ();
- gimplify_scan_omp_clauses (&OMP_TASK_CLAUSES (expr), pre_p,
- find_omp_clause (OMP_TASK_CLAUSES (expr),
- OMP_CLAUSE_UNTIED)
- ? ORT_UNTIED_TASK : ORT_TASK);
+ ctx->default_kind = OMP_CLAUSE_DEFAULT_UNSPECIFIED;
+ ctx->region_type = ORT_TASK;
+ if (find_omp_clause (OMP_TASK_CLAUSES (expr), OMP_CLAUSE_UNTIED))
+ ctx->region_type_flags.task.untied = true;
+
+ gimplify_scan_omp_clauses (&OMP_TASK_CLAUSES (expr), pre_p, ctx);
push_gimplify_context ();
@@ -6776,13 +6877,29 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
int i;
bool simd;
bitmap has_decl_expr = NULL;
+ struct gimplify_omp_ctx *ctx = new_omp_context ();
orig_for_stmt = for_stmt = *expr_p;
- simd = (TREE_CODE (for_stmt) == OMP_SIMD
- || TREE_CODE (for_stmt) == CILK_SIMD);
- gimplify_scan_omp_clauses (&OMP_FOR_CLAUSES (for_stmt), pre_p,
- simd ? ORT_SIMD : ORT_WORKSHARE);
+ switch (TREE_CODE (for_stmt))
+ {
+ case OMP_FOR:
+ case CILK_FOR:
+ case OMP_DISTRIBUTE:
+ case OACC_LOOP:
+ simd = false;
+ ctx->region_type = ORT_WORKSHARE;
+ break;
+ case OMP_SIMD:
+ case CILK_SIMD:
+ simd = true;
+ ctx->region_type = ORT_SIMD;
+ break;
+ default:
+ gcc_unreachable ();
+ }
+
+ gimplify_scan_omp_clauses (&OMP_FOR_CLAUSES (for_stmt), pre_p, ctx);
if (TREE_CODE (for_stmt) == OMP_DISTRIBUTE)
gimplify_omp_ctxp->distribute = true;
@@ -6816,6 +6933,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
if (OMP_FOR_INIT (for_stmt) == NULL_TREE)
{
+ gcc_assert (TREE_CODE (for_stmt) != OACC_LOOP);
for_stmt = walk_tree (&OMP_FOR_BODY (for_stmt), find_combined_omp_for,
NULL, NULL);
gcc_assert (for_stmt != NULL_TREE);
@@ -7117,6 +7235,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
case CILK_SIMD: kind = GF_OMP_FOR_KIND_CILKSIMD; break;
case CILK_FOR: kind = GF_OMP_FOR_KIND_CILKFOR; break;
case OMP_DISTRIBUTE: kind = GF_OMP_FOR_KIND_DISTRIBUTE; break;
+ case OACC_LOOP: kind = GF_OMP_FOR_KIND_OACC_LOOP; break;
default:
gcc_unreachable ();
}
@@ -7157,9 +7276,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
return GS_ALL_DONE;
}
-/* Gimplify the gross structure of other OpenMP constructs.
- In particular, OMP_SECTIONS, OMP_SINGLE, OMP_TARGET, OMP_TARGET_DATA
- and OMP_TEAMS. */
+/* Gimplify the gross structure of several OpenACC or OpenMP constructs. */
static void
gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
@@ -7167,27 +7284,34 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
tree expr = *expr_p;
gimple stmt;
gimple_seq body = NULL;
- enum omp_region_type ort = ORT_WORKSHARE;
+ struct gimplify_omp_ctx *ctx = new_omp_context ();
switch (TREE_CODE (expr))
{
case OMP_SECTIONS:
case OMP_SINGLE:
+ ctx->region_type = ORT_WORKSHARE;
break;
+ case OACC_KERNELS:
+ case OACC_PARALLEL:
case OMP_TARGET:
- ort = ORT_TARGET;
+ ctx->region_type = ORT_TARGET;
+ ctx->region_type_flags.target.offload = true;
break;
+ case OACC_DATA:
case OMP_TARGET_DATA:
- ort = ORT_TARGET_DATA;
+ ctx->region_type = ORT_TARGET;
break;
case OMP_TEAMS:
- ort = ORT_TEAMS;
+ ctx->region_type = ORT_TEAMS;
break;
default:
gcc_unreachable ();
}
- gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ort);
- if (ort == ORT_TARGET || ort == ORT_TARGET_DATA)
+
+ gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ctx);
+
+ if (ctx->region_type == ORT_TARGET)
{
push_gimplify_context ();
gimple g = gimplify_and_return_first (OMP_BODY (expr), &body);
@@ -7195,11 +7319,23 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
pop_gimplify_context (g);
else
pop_gimplify_context (NULL);
- if (ort == ORT_TARGET_DATA)
+ if (!ctx->region_type_flags.target.offload)
{
- gimple_seq cleanup = NULL;
- tree fn = builtin_decl_explicit (BUILT_IN_GOMP_TARGET_END_DATA);
+ enum built_in_function end_ix;
+ switch (TREE_CODE (expr))
+ {
+ case OACC_DATA:
+ end_ix = BUILT_IN_GOACC_DATA_END;
+ break;
+ case OMP_TARGET_DATA:
+ end_ix = BUILT_IN_GOMP_TARGET_END_DATA;
+ break;
+ default:
+ gcc_unreachable ();
+ }
+ tree fn = builtin_decl_explicit (end_ix);
g = gimple_build_call (fn, 0);
+ gimple_seq cleanup = NULL;
gimple_seq_add_stmt (&cleanup, g);
g = gimple_build_try (body, cleanup, GIMPLE_TRY_FINALLY);
body = NULL;
@@ -7212,6 +7348,16 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
switch (TREE_CODE (expr))
{
+ case OACC_DATA:
+ stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_DATA,
+ OACC_DATA_CLAUSES (expr));
+ break;
+ case OACC_KERNELS:
+ stmt = gimple_build_oacc_kernels (body, OACC_KERNELS_CLAUSES (expr));
+ break;
+ case OACC_PARALLEL:
+ stmt = gimple_build_oacc_parallel (body, OACC_PARALLEL_CLAUSES (expr));
+ break;
case OMP_SECTIONS:
stmt = gimple_build_omp_sections (body, OMP_CLAUSES (expr));
break;
@@ -7237,19 +7383,42 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
*expr_p = NULL_TREE;
}
-/* Gimplify the gross structure of OpenMP target update construct. */
+/* Gimplify the gross structure of OpenACC enter data and exit data, OpenACC
+ update, and OpenMP target update constructs. */
static void
gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
{
- tree expr = *expr_p;
- gomp_target *stmt;
+ tree expr = *expr_p, clauses;
+ int kind;
+ gimple stmt;
+ struct gimplify_omp_ctx *ctx = new_omp_context ();
- gimplify_scan_omp_clauses (&OMP_TARGET_UPDATE_CLAUSES (expr), pre_p,
- ORT_WORKSHARE);
- gimplify_adjust_omp_clauses (pre_p, &OMP_TARGET_UPDATE_CLAUSES (expr));
- stmt = gimple_build_omp_target (NULL, GF_OMP_TARGET_KIND_UPDATE,
- OMP_TARGET_UPDATE_CLAUSES (expr));
+ switch (TREE_CODE (expr))
+ {
+ case OACC_ENTER_DATA:
+ clauses = OACC_ENTER_DATA_CLAUSES (expr);
+ kind = GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA;
+ break;
+ case OACC_EXIT_DATA:
+ clauses = OACC_EXIT_DATA_CLAUSES (expr);
+ kind = GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA;
+ break;
+ case OACC_UPDATE:
+ clauses = OACC_UPDATE_CLAUSES (expr);
+ kind = GF_OMP_TARGET_KIND_OACC_UPDATE;
+ break;
+ case OMP_TARGET_UPDATE:
+ clauses = OMP_TARGET_UPDATE_CLAUSES (expr);
+ kind = GF_OMP_TARGET_KIND_UPDATE;
+ break;
+ default:
+ gcc_unreachable ();
+ }
+ ctx->region_type = ORT_WORKSHARE;
+ gimplify_scan_omp_clauses (&clauses, pre_p, ctx);
+ gimplify_adjust_omp_clauses (pre_p, &clauses);
+ stmt = gimple_build_omp_target (NULL, kind, clauses);
gimplify_seq_add_stmt (pre_p, stmt);
*expr_p = NULL_TREE;
@@ -8192,9 +8361,38 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
case CILK_SIMD:
case CILK_FOR:
case OMP_DISTRIBUTE:
+ case OACC_LOOP:
ret = gimplify_omp_for (expr_p, pre_p);
break;
+ case OACC_CACHE:
+ gimplify_oacc_cache (expr_p, pre_p);
+ ret = GS_ALL_DONE;
+ break;
+
+ case OACC_DECLARE:
+ case OACC_HOST_DATA:
+ sorry ("directive not yet implemented");
+ ret = GS_ALL_DONE;
+ break;
+
+ case OACC_KERNELS:
+ if (OACC_KERNELS_COMBINED (*expr_p))
+ sorry ("directive not yet implemented");
+ else
+ gimplify_omp_workshare (expr_p, pre_p);
+ ret = GS_ALL_DONE;
+ break;
+
+ case OACC_PARALLEL:
+ if (OACC_PARALLEL_COMBINED (*expr_p))
+ sorry ("directive not yet implemented");
+ else
+ gimplify_omp_workshare (expr_p, pre_p);
+ ret = GS_ALL_DONE;
+ break;
+
+ case OACC_DATA:
case OMP_SECTIONS:
case OMP_SINGLE:
case OMP_TARGET:
@@ -8204,6 +8402,9 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
ret = GS_ALL_DONE;
break;
+ case OACC_ENTER_DATA:
+ case OACC_EXIT_DATA:
+ case OACC_UPDATE:
case OMP_TARGET_UPDATE:
gimplify_omp_target_update (expr_p, pre_p);
ret = GS_ALL_DONE;
@@ -8585,8 +8786,18 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
&& code != LOOP_EXPR
&& code != SWITCH_EXPR
&& code != TRY_FINALLY_EXPR
+ && code != OACC_PARALLEL
+ && code != OACC_KERNELS
+ && code != OACC_DATA
+ && code != OACC_HOST_DATA
+ && code != OACC_DECLARE
+ && code != OACC_UPDATE
+ && code != OACC_ENTER_DATA
+ && code != OACC_EXIT_DATA
+ && code != OACC_CACHE
&& code != OMP_CRITICAL
&& code != OMP_FOR
+ && code != OACC_LOOP
&& code != OMP_MASTER
&& code != OMP_TASKGROUP
&& code != OMP_ORDERED
@@ -8813,11 +9024,17 @@ gimplify_body (tree fndecl, bool do_parms)
gcc_assert (gimplify_ctxp == NULL);
push_gimplify_context ();
- if (flag_openmp)
+ if (flag_openacc || flag_openmp)
{
gcc_assert (gimplify_omp_ctxp == NULL);
if (lookup_attribute ("omp declare target", DECL_ATTRIBUTES (fndecl)))
- gimplify_omp_ctxp = new_omp_context (ORT_TARGET);
+ {
+ struct gimplify_omp_ctx *ctx = new_omp_context ();
+ ctx->region_type = ORT_TARGET;
+ ctx->region_type_flags.target.offload = true;
+
+ gimplify_omp_ctxp = ctx;
+ }
}
/* Unshare most shared trees in the body and in that of any nested functions.
@@ -8897,7 +9114,8 @@ gimplify_body (tree fndecl, bool do_parms)
nonlocal_vlas = NULL;
}
- if ((flag_openmp || flag_openmp_simd) && gimplify_omp_ctxp)
+ if ((flag_openacc || flag_openmp || flag_openmp_simd)
+ && gimplify_omp_ctxp)
{
delete_omp_context (gimplify_omp_ctxp);
gimplify_omp_ctxp = NULL;
@@ -171,6 +171,11 @@ enum lto_builtin_type
#define DEF_FUNCTION_TYPE_VAR_4(NAME, RETURN, ARG1, ARG2, ARG3, ARG4) NAME,
#define DEF_FUNCTION_TYPE_VAR_5(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG6) \
NAME,
+#define DEF_FUNCTION_TYPE_VAR_8(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
+ ARG6, ARG7, ARG8) NAME,
+#define DEF_FUNCTION_TYPE_VAR_12(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
+ ARG6, ARG7, ARG8, ARG9, ARG10, ARG11, \
+ ARG12) NAME,
#define DEF_POINTER_TYPE(NAME, TYPE) NAME,
#include "builtin-types.def"
#undef DEF_PRIMITIVE_TYPE
@@ -189,6 +194,8 @@ enum lto_builtin_type
#undef DEF_FUNCTION_TYPE_VAR_3
#undef DEF_FUNCTION_TYPE_VAR_4
#undef DEF_FUNCTION_TYPE_VAR_5
+#undef DEF_FUNCTION_TYPE_VAR_8
+#undef DEF_FUNCTION_TYPE_VAR_12
#undef DEF_POINTER_TYPE
BT_LAST
};
@@ -673,6 +680,14 @@ lto_define_builtins (tree va_list_ref_type_node ATTRIBUTE_UNUSED,
def_fn_type (ENUM, RETURN, 1, 4, ARG1, ARG2, ARG3, ARG4);
#define DEF_FUNCTION_TYPE_VAR_5(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5) \
def_fn_type (ENUM, RETURN, 1, 5, ARG1, ARG2, ARG3, ARG4, ARG5);
+#define DEF_FUNCTION_TYPE_VAR_8(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
+ ARG6, ARG7, ARG8) \
+ def_fn_type (ENUM, RETURN, 1, 8, ARG1, ARG2, ARG3, ARG4, ARG5, ARG6, \
+ ARG7, ARG8);
+#define DEF_FUNCTION_TYPE_VAR_12(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
+ ARG6, ARG7, ARG8, ARG9, ARG10, ARG11, ARG12) \
+ def_fn_type (ENUM, RETURN, 1, 12, ARG1, ARG2, ARG3, ARG4, ARG5, ARG6, \
+ ARG7, ARG8, ARG9, ARG10, ARG11, ARG12);
#define DEF_POINTER_TYPE(ENUM, TYPE) \
builtin_types[(int) ENUM] = build_pointer_type (builtin_types[(int) TYPE]);
@@ -694,6 +709,8 @@ lto_define_builtins (tree va_list_ref_type_node ATTRIBUTE_UNUSED,
#undef DEF_FUNCTION_TYPE_VAR_3
#undef DEF_FUNCTION_TYPE_VAR_4
#undef DEF_FUNCTION_TYPE_VAR_5
+#undef DEF_FUNCTION_TYPE_VAR_8
+#undef DEF_FUNCTION_TYPE_VAR_12
#undef DEF_POINTER_TYPE
builtin_types[(int) BT_LAST] = NULL_TREE;
new file mode 100644
@@ -0,0 +1,56 @@
+/* This file contains the definitions and documentation for the
+ OpenACC builtins used in the GNU compiler.
+
+ Copyright (C) 2013-2014 Free Software Foundation, Inc.
+
+ Contributed by Thomas Schwinge <thomas@codesourcery.com>.
+
+This file is part of GCC.
+
+GCC is free software; you can redistribute it and/or modify it under
+the terms of the GNU General Public License as published by the Free
+Software Foundation; either version 3, or (at your option) any later
+version.
+
+GCC is distributed in the hope that it will be useful, but WITHOUT ANY
+WARRANTY; without even the implied warranty of MERCHANTABILITY or
+FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+for more details.
+
+You should have received a copy of the GNU General Public License
+along with GCC; see the file COPYING3. If not see
+<http://www.gnu.org/licenses/>. */
+
+/* Before including this file, you should define a macro:
+
+ DEF_GOACC_BUILTIN (ENUM, NAME, TYPE, ATTRS)
+
+ See builtins.def for details. */
+
+DEF_GOACC_BUILTIN (BUILT_IN_ACC_GET_DEVICE_TYPE, "acc_get_device_type",
+ BT_FN_INT, ATTR_NOTHROW_LIST)
+DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DATA_START, "GOACC_data_start",
+ BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
+DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DATA_END, "GOACC_data_end",
+ BT_FN_VOID, ATTR_NOTHROW_LIST)
+DEF_GOACC_BUILTIN (BUILT_IN_GOACC_ENTER_EXIT_DATA, "GOACC_enter_exit_data",
+ BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR_INT_INT_VAR,
+ ATTR_NOTHROW_LIST)
+DEF_GOACC_BUILTIN (BUILT_IN_GOACC_KERNELS, "GOACC_kernels",
+ BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR_INT_INT_INT_INT_INT_VAR,
+ ATTR_NOTHROW_LIST)
+DEF_GOACC_BUILTIN (BUILT_IN_GOACC_PARALLEL, "GOACC_parallel",
+ BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR_INT_INT_INT_INT_INT_VAR,
+ ATTR_NOTHROW_LIST)
+DEF_GOACC_BUILTIN (BUILT_IN_GOACC_UPDATE, "GOACC_update",
+ BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR_INT_INT_VAR,
+ ATTR_NOTHROW_LIST)
+DEF_GOACC_BUILTIN (BUILT_IN_GOACC_WAIT, "GOACC_wait",
+ BT_FN_VOID_INT_INT_VAR,
+ ATTR_NOTHROW_LIST)
+DEF_GOACC_BUILTIN_COMPILER (BUILT_IN_ACC_ON_DEVICE, "acc_on_device",
+ BT_FN_INT_INT, ATTR_CONST_NOTHROW_LEAF_LIST)
+DEF_GOACC_BUILTIN (BUILT_IN_GOACC_GET_THREAD_NUM, "GOACC_get_thread_num",
+ BT_FN_INT, ATTR_CONST_NOTHROW_LEAF_LIST)
+DEF_GOACC_BUILTIN (BUILT_IN_GOACC_GET_NUM_THREADS, "GOACC_get_num_threads",
+ BT_FN_INT, ATTR_CONST_NOTHROW_LEAF_LIST)
@@ -175,6 +175,11 @@ typedef struct omp_context
construct. In the case of a parallel, this is in the child function. */
tree block_vars;
+ /* A map of reduction pointer variables. For accelerators, each
+ reduction variable is replaced with an array. Each thread, in turn,
+ is assigned to a slot on that array. */
+ splay_tree reduction_map;
+
/* Label to which GOMP_cancel{,llation_point} and explicit and implicit
barriers should jump to during omplower pass. */
tree cancel_label;
@@ -195,6 +200,8 @@ typedef struct omp_context
bool cancellable;
} omp_context;
+/* A structure holding the elements of:
+ for (V = N1; V cond N2; V += STEP) [...] */
struct omp_for_data_loop
{
@@ -240,6 +247,87 @@ static tree scan_omp_1_op (tree *, int *, void *);
/* Holds offload tables with decls. */
vec<tree, va_gc> *offload_funcs, *offload_vars;
+/* Helper function to get the reduction array name */
+static const char *
+omp_get_id (tree node)
+{
+ const char *id = IDENTIFIER_POINTER (DECL_NAME (node));
+ int len = strlen ("omp$") + strlen (id);
+ char *temp_name = (char *)alloca (len+1);
+ snprintf (temp_name, len+1, "gfc$%s", id);
+ return IDENTIFIER_POINTER(get_identifier (temp_name));
+}
+
+/* Determine the number of threads OpenACC threads used to determine the
+ size of the array of partial reductions. Currently, this is num_gangs
+ * vector_length. This value may be different than GOACC_GET_NUM_THREADS,
+ because it is independed of the device used. */
+
+static tree
+oacc_max_threads (omp_context *ctx)
+{
+ tree nthreads, vector_length, gangs, clauses;
+
+ gangs = fold_convert (sizetype, integer_one_node);
+ vector_length = gangs;
+
+ /* The reduction clause may be nested inside a loop directive.
+ Scan for the innermost vector_length clause. */
+ for (omp_context *oc = ctx; oc; oc = oc->outer)
+ {
+ if (gimple_code (oc->stmt) != GIMPLE_OACC_PARALLEL)
+ continue;
+
+ clauses = gimple_oacc_parallel_clauses (oc->stmt);
+
+ vector_length = find_omp_clause (clauses, OMP_CLAUSE_VECTOR_LENGTH);
+ if (vector_length)
+ vector_length = fold_convert_loc (OMP_CLAUSE_LOCATION (vector_length),
+ sizetype,
+ OMP_CLAUSE_VECTOR_LENGTH_EXPR
+ (vector_length));
+ else
+ vector_length = fold_convert (sizetype, integer_one_node);
+
+ gangs = find_omp_clause (clauses, OMP_CLAUSE_NUM_GANGS);
+ if (gangs)
+ gangs = fold_convert_loc (OMP_CLAUSE_LOCATION (gangs), sizetype,
+ OMP_CLAUSE_NUM_GANGS_EXPR (gangs));
+ else
+ gangs = fold_convert (sizetype, integer_one_node);
+
+ break;
+ }
+
+ nthreads = fold_build2 (MULT_EXPR, sizetype, gangs, vector_length);
+
+ return nthreads;
+}
+
+/* Holds a decl for __OPENMP_TARGET__. */
+static GTY(()) tree offload_symbol_decl;
+
+/* Get the __OFFLOAD_TABLE__ symbol. */
+static tree
+get_offload_symbol_decl (void)
+{
+ if (!offload_symbol_decl)
+ {
+ tree decl = build_decl (UNKNOWN_LOCATION, VAR_DECL,
+ get_identifier ("__OFFLOAD_TABLE__"),
+ ptr_type_node);
+ TREE_ADDRESSABLE (decl) = 1;
+ TREE_PUBLIC (decl) = 1;
+ DECL_EXTERNAL (decl) = 1;
+ DECL_WEAK (decl) = 1;
+ DECL_ATTRIBUTES (decl)
+ = tree_cons (get_identifier ("weak"),
+ NULL_TREE, DECL_ATTRIBUTES (decl));
+ offload_symbol_decl = decl;
+ }
+ return offload_symbol_decl;
+}
+
/* Convenience function for calling scan_omp_1_op on tree operands. */
static inline tree
@@ -613,6 +701,15 @@ extract_omp_for_data (gomp_for *for_stmt, struct omp_for_data *fd,
fd->loop.step = build_int_cst (TREE_TYPE (fd->loop.v), 1);
fd->loop.cond_code = LT_EXPR;
}
+
+ /* For OpenACC loops, force a chunk size of one, as this avoids the default
+ scheduling where several subsequent iterations are being executed by the
+ same thread. */
+ if (gimple_omp_for_kind (for_stmt) == GF_OMP_FOR_KIND_OACC_LOOP)
+ {
+ gcc_assert (fd->chunk_size == NULL_TREE);
+ fd->chunk_size = build_int_cst (TREE_TYPE (fd->loop.v), 1);
+ }
}
@@ -841,7 +938,18 @@ is_reference (tree decl)
return lang_hooks.decls.omp_privatize_by_reference (decl);
}
-/* Lookup variables in the decl or field splay trees. The "maybe" form
+/* Return the type of a decl. If the decl is reference type,
+ return its base type. */
+static inline tree
+get_base_type (tree decl)
+{
+ tree type = TREE_TYPE (decl);
+ if (is_reference (decl))
+ type = TREE_TYPE (type);
+ return type;
+}
+
+/* Lookup variables. The "maybe" form
allows for the variable form to not have been entered, otherwise we
assert that the variable must have been entered. */
@@ -885,6 +993,25 @@ maybe_lookup_field (tree var, omp_context *ctx)
return n ? (tree) n->value : NULL_TREE;
}
+static inline tree
+lookup_reduction (const char *id, omp_context *ctx)
+{
+ gcc_assert (is_gimple_omp_oacc_specifically (ctx->stmt));
+
+ splay_tree_node n;
+ n = splay_tree_lookup (ctx->reduction_map, (splay_tree_key) id);
+ return (tree) n->value;
+}
+
+static inline tree
+maybe_lookup_reduction (tree var, omp_context *ctx)
+{
+ splay_tree_node n = NULL;
+ if (ctx->reduction_map)
+ n = splay_tree_lookup (ctx->reduction_map, (splay_tree_key) var);
+ return n ? (tree) n->value : NULL_TREE;
+}
+
/* Return true if DECL should be copied by pointer. SHARED_CTX is
the parallel context if DECL is to be shared. */
@@ -898,6 +1025,8 @@ use_pointer_for_field (tree decl, omp_context *shared_ctx)
when we know the value is not accessible from an outer scope. */
if (shared_ctx)
{
+ gcc_assert (!is_gimple_omp_oacc_specifically (shared_ctx->stmt));
+
/* ??? Trivially accessible from anywhere. But why would we even
be passing an address in this case? Should we simply assert
this to be false, or should we have a cleanup pass that removes
@@ -1102,6 +1231,8 @@ install_var_field (tree var, bool by_ref, int mask, omp_context *ctx)
|| !splay_tree_lookup (ctx->field_map, (splay_tree_key) var));
gcc_assert ((mask & 2) == 0 || !ctx->sfield_map
|| !splay_tree_lookup (ctx->sfield_map, (splay_tree_key) var));
+ gcc_assert ((mask & 3) == 3
+ || !is_gimple_omp_oacc_specifically (ctx->stmt));
type = TREE_TYPE (var);
if (mask & 4)
@@ -1378,6 +1509,7 @@ new_omp_context (gimple stmt, omp_context *outer_ctx)
ctx->cb = outer_ctx->cb;
ctx->cb.block = NULL;
ctx->depth = outer_ctx->depth + 1;
+ ctx->reduction_map = outer_ctx->reduction_map;
}
else
{
@@ -1448,6 +1580,11 @@ delete_omp_context (splay_tree_value value)
splay_tree_delete (ctx->field_map);
if (ctx->sfield_map)
splay_tree_delete (ctx->sfield_map);
+ if (ctx->reduction_map
+ /* Shared over several omp_contexts. */
+ && (ctx->outer == NULL
+ || ctx->reduction_map != ctx->outer->reduction_map))
+ splay_tree_delete (ctx->reduction_map);
/* We hijacked DECL_ABSTRACT_ORIGIN earlier. We need to clear it before
it produces corrupt debug information. */
@@ -1545,6 +1682,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
break;
case OMP_CLAUSE_SHARED:
+ gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt));
decl = OMP_CLAUSE_DECL (c);
/* Ignore shared directives in teams construct. */
if (gimple_code (ctx->stmt) == GIMPLE_OMP_TEAMS)
@@ -1579,6 +1717,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
goto do_private;
case OMP_CLAUSE_LASTPRIVATE:
+ gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt));
/* Let the corresponding firstprivate clause create
the variable. */
if (OMP_CLAUSE_LASTPRIVATE_FIRSTPRIVATE (c))
@@ -1586,8 +1725,16 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
/* FALLTHRU */
case OMP_CLAUSE_FIRSTPRIVATE:
- case OMP_CLAUSE_REDUCTION:
+ if (is_gimple_omp_oacc_specifically (ctx->stmt))
+ {
+ sorry ("clause not supported yet");
+ break;
+ }
+ /* FALLTHRU */
case OMP_CLAUSE_LINEAR:
+ gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt));
+ /* FALLTHRU */
+ case OMP_CLAUSE_REDUCTION:
decl = OMP_CLAUSE_DECL (c);
do_private:
if (is_variable_sized (decl))
@@ -1613,9 +1760,30 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
install_var_field (decl, by_ref, 3, ctx);
}
install_var_local (decl, ctx);
+ if (is_gimple_omp_oacc_specifically (ctx->stmt)
+ && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
+ {
+ /* Create a decl for the reduction array. */
+ tree var = OMP_CLAUSE_DECL (c);
+ tree type = get_base_type (var);
+ tree ptype = build_pointer_type (type);
+ tree array = create_tmp_var (ptype, omp_get_id (var));
+ omp_context *c = (ctx->field_map ? ctx : ctx->outer);
+ install_var_field (array, true, 3, c);
+ install_var_local (array, c);
+
+ /* Insert it into the current context. */
+ splay_tree_insert (ctx->reduction_map,
+ (splay_tree_key) omp_get_id(var),
+ (splay_tree_value) array);
+ splay_tree_insert (ctx->reduction_map,
+ (splay_tree_key) array,
+ (splay_tree_value) array);
+ }
break;
case OMP_CLAUSE__LOOPTEMP_:
+ gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt));
gcc_assert (is_parallel_ctx (ctx));
decl = OMP_CLAUSE_DECL (c);
install_var_field (decl, false, 3, ctx);
@@ -1624,17 +1792,18 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_COPYPRIVATE:
case OMP_CLAUSE_COPYIN:
+ gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt));
decl = OMP_CLAUSE_DECL (c);
by_ref = use_pointer_for_field (decl, NULL);
install_var_field (decl, by_ref, 3, ctx);
break;
case OMP_CLAUSE_DEFAULT:
+ gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt));
ctx->default_kind = OMP_CLAUSE_DEFAULT_KIND (c);
break;
case OMP_CLAUSE_FINAL:
- case OMP_CLAUSE_IF:
case OMP_CLAUSE_NUM_THREADS:
case OMP_CLAUSE_NUM_TEAMS:
case OMP_CLAUSE_THREAD_LIMIT:
@@ -1643,13 +1812,41 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_DIST_SCHEDULE:
case OMP_CLAUSE_DEPEND:
case OMP_CLAUSE__CILK_FOR_COUNT_:
+ gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt));
+ /* FALLTHRU */
+ case OMP_CLAUSE_IF:
+ case OMP_CLAUSE_NUM_GANGS:
+ case OMP_CLAUSE_NUM_WORKERS:
+ case OMP_CLAUSE_VECTOR_LENGTH:
if (ctx->outer)
scan_omp_op (&OMP_CLAUSE_OPERAND (c, 0), ctx->outer);
break;
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
+ gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt));
+ /* FALLTHRU */
case OMP_CLAUSE_MAP:
+ switch (OMP_CLAUSE_CODE (c))
+ {
+ case OMP_CLAUSE_TO:
+ case OMP_CLAUSE_FROM:
+ /* The to and from clauses are only ever seen with OpenMP target
+ update constructs. */
+ gcc_assert (gimple_code (ctx->stmt) == GIMPLE_OMP_TARGET
+ && (gimple_omp_target_kind (ctx->stmt)
+ == GF_OMP_TARGET_KIND_UPDATE));
+ break;
+ case OMP_CLAUSE_MAP:
+ /* The map clause is never seen with OpenMP target update
+ constructs. */
+ gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OMP_TARGET
+ || (gimple_omp_target_kind (ctx->stmt)
+ != GF_OMP_TARGET_KIND_UPDATE));
+ break;
+ default:
+ gcc_unreachable ();
+ }
if (ctx->outer)
scan_omp_op (&OMP_CLAUSE_SIZE (c), ctx->outer);
decl = OMP_CLAUSE_DECL (c);
@@ -1660,14 +1857,17 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
&& DECL_P (decl)
&& is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
&& varpool_node::get_create (decl)->offloadable)
- break;
+ {
+ gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt));
+ break;
+ }
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER)
{
/* Ignore OMP_CLAUSE_MAP_POINTER kind for arrays in
- #pragma omp target data, there is nothing to map for
+ data regions that are not offloaded; there is nothing to map for
those. */
- if (gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA
+ if (!is_gimple_omp_offloaded (ctx->stmt)
&& !POINTER_TYPE_P (TREE_TYPE (decl)))
break;
}
@@ -1686,6 +1886,10 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
}
else
{
+ gcc_assert (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+ || (OMP_CLAUSE_MAP_KIND (c)
+ != OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
+ || TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
@@ -1693,8 +1897,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
install_var_field (decl, true, 7, ctx);
else
install_var_field (decl, true, 3, ctx);
- if (gimple_omp_target_kind (ctx->stmt)
- == GF_OMP_TARGET_KIND_REGION)
+ if (is_gimple_omp_offloaded (ctx->stmt))
install_var_local (decl, ctx);
}
}
@@ -1734,20 +1937,35 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_NOWAIT:
case OMP_CLAUSE_ORDERED:
- case OMP_CLAUSE_COLLAPSE:
case OMP_CLAUSE_UNTIED:
case OMP_CLAUSE_MERGEABLE:
case OMP_CLAUSE_PROC_BIND:
case OMP_CLAUSE_SAFELEN:
+ gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt));
+ /* FALLTHRU */
+ case OMP_CLAUSE_COLLAPSE:
+ case OMP_CLAUSE_ASYNC:
+ case OMP_CLAUSE_WAIT:
break;
case OMP_CLAUSE_ALIGNED:
+ gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt));
decl = OMP_CLAUSE_DECL (c);
if (is_global_var (decl)
&& TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
install_var_local (decl, ctx);
break;
+ case OMP_CLAUSE_DEVICE_RESIDENT:
+ case OMP_CLAUSE_USE_DEVICE:
+ case OMP_CLAUSE_GANG:
+ case OMP_CLAUSE__CACHE_:
+ case OMP_CLAUSE_INDEPENDENT:
+ case OMP_CLAUSE_WORKER:
+ case OMP_CLAUSE_VECTOR:
+ sorry ("Clause not supported yet");
+ break;
+
default:
gcc_unreachable ();
}
@@ -1758,6 +1976,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
switch (OMP_CLAUSE_CODE (c))
{
case OMP_CLAUSE_LASTPRIVATE:
+ gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt));
/* Let the corresponding firstprivate clause create
the variable. */
if (OMP_CLAUSE_LASTPRIVATE_GIMPLE_SEQ (c))
@@ -1766,10 +1985,18 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
break;
/* FALLTHRU */
- case OMP_CLAUSE_PRIVATE:
case OMP_CLAUSE_FIRSTPRIVATE:
- case OMP_CLAUSE_REDUCTION:
+ if (is_gimple_omp_oacc_specifically (ctx->stmt))
+ {
+ sorry ("clause not supported yet");
+ break;
+ }
+ /* FALLTHRU */
case OMP_CLAUSE_LINEAR:
+ gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt));
+ /* FALLTHRU */
+ case OMP_CLAUSE_PRIVATE:
+ case OMP_CLAUSE_REDUCTION:
decl = OMP_CLAUSE_DECL (c);
if (is_variable_sized (decl))
install_var_local (decl, ctx);
@@ -1785,6 +2012,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
break;
case OMP_CLAUSE_SHARED:
+ gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt));
/* Ignore shared directives in teams construct. */
if (gimple_code (ctx->stmt) == GIMPLE_OMP_TEAMS)
break;
@@ -1794,15 +2022,26 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
break;
case OMP_CLAUSE_MAP:
- if (gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA)
+ /* The map clause is never seen with OpenMP target update
+ constructs. */
+ gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OMP_TARGET
+ || (gimple_omp_target_kind (ctx->stmt)
+ != GF_OMP_TARGET_KIND_UPDATE));
+ if (!is_gimple_omp_offloaded (ctx->stmt))
break;
decl = OMP_CLAUSE_DECL (c);
if (DECL_P (decl)
&& is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
&& varpool_node::get_create (decl)->offloadable)
- break;
+ {
+ gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt));
+ break;
+ }
if (DECL_P (decl))
{
+ gcc_assert ((OMP_CLAUSE_MAP_KIND (c)
+ != OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
+ || TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE);
if (OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
&& TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE
&& !COMPLETE_TYPE_P (TREE_TYPE (decl)))
@@ -1814,6 +2053,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
else if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
{
+ gcc_assert (OMP_CLAUSE_MAP_KIND (c)
+ != OMP_CLAUSE_MAP_FORCE_DEVICEPTR);
+
tree decl2 = DECL_VALUE_EXPR (decl);
gcc_assert (TREE_CODE (decl2) == INDIRECT_REF);
decl2 = TREE_OPERAND (decl2, 0);
@@ -1829,7 +2071,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_COPYPRIVATE:
case OMP_CLAUSE_COPYIN:
case OMP_CLAUSE_DEFAULT:
- case OMP_CLAUSE_IF:
case OMP_CLAUSE_NUM_THREADS:
case OMP_CLAUSE_NUM_TEAMS:
case OMP_CLAUSE_THREAD_LIMIT:
@@ -1838,7 +2079,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_DIST_SCHEDULE:
case OMP_CLAUSE_NOWAIT:
case OMP_CLAUSE_ORDERED:
- case OMP_CLAUSE_COLLAPSE:
case OMP_CLAUSE_UNTIED:
case OMP_CLAUSE_FINAL:
case OMP_CLAUSE_MERGEABLE:
@@ -1850,6 +2090,25 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
case OMP_CLAUSE__CILK_FOR_COUNT_:
+ gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt));
+ /* FALLTHRU */
+ case OMP_CLAUSE_ASYNC:
+ case OMP_CLAUSE_COLLAPSE:
+ case OMP_CLAUSE_IF:
+ case OMP_CLAUSE_NUM_GANGS:
+ case OMP_CLAUSE_NUM_WORKERS:
+ case OMP_CLAUSE_VECTOR_LENGTH:
+ case OMP_CLAUSE_WAIT:
+ break;
+
+ case OMP_CLAUSE_DEVICE_RESIDENT:
+ case OMP_CLAUSE_USE_DEVICE:
+ case OMP_CLAUSE_GANG:
+ case OMP_CLAUSE__CACHE_:
+ case OMP_CLAUSE_INDEPENDENT:
+ case OMP_CLAUSE_WORKER:
+ case OMP_CLAUSE_VECTOR:
+ sorry ("Clause not supported yet");
break;
default:
@@ -1857,6 +2116,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
}
}
+ gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt)
+ || !scan_array_reductions);
if (scan_array_reductions)
for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION
@@ -1939,6 +2200,8 @@ create_omp_child_function (omp_context *ctx, bool task_copy)
decl = build_decl (gimple_location (ctx->stmt), FUNCTION_DECL, name, type);
+ gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt)
+ || !task_copy);
if (!task_copy)
ctx->cb.dst_fn = decl;
else
@@ -1959,7 +2222,7 @@ create_omp_child_function (omp_context *ctx, bool task_copy)
{
omp_context *octx;
for (octx = ctx; octx; octx = octx->outer)
- if (is_targetreg_ctx (octx))
+ if (is_gimple_omp_offloaded (octx->stmt))
{
cgraph_node::get_create (decl)->offloadable = 1;
#ifdef ENABLE_OFFLOADING
@@ -2294,7 +2557,7 @@ finish_taskreg_scan (omp_context *ctx)
}
-/* Scan an OpenMP loop directive. */
+/* Scan a GIMPLE_OMP_FOR. */
static void
scan_omp_for (gomp_for *stmt, omp_context *outer_ctx)
@@ -2354,14 +2617,41 @@ scan_omp_single (gomp_single *stmt, omp_context *outer_ctx)
layout_type (ctx->record_type);
}
-/* Scan an OpenMP target{, data, update} directive. */
+/* Scan a GIMPLE_OMP_TARGET. */
static void
-scan_omp_target (gomp_target *stmt, omp_context *outer_ctx)
+scan_omp_target (gimple stmt, omp_context *outer_ctx)
{
omp_context *ctx;
tree name;
- int kind = gimple_omp_target_kind (stmt);
+ bool offloaded;
+ void (*gimple_omp_set_child_fn) (gimple, tree);
+ tree (*gimple_omp_clauses) (const_gimple);
+
+ offloaded = is_gimple_omp_offloaded (stmt);
+ switch (gimple_code (stmt))
+ {
+ case GIMPLE_OACC_KERNELS:
+ gimple_omp_set_child_fn = gimple_oacc_kernels_set_child_fn;
+ gimple_omp_clauses = gimple_oacc_kernels_clauses;
+ break;
+ case GIMPLE_OACC_PARALLEL:
+ gimple_omp_set_child_fn = gimple_oacc_parallel_set_child_fn;
+ gimple_omp_clauses = gimple_oacc_parallel_clauses;
+ break;
+ case GIMPLE_OMP_TARGET:
+ gimple_omp_set_child_fn = gimple_omp_target_set_child_fn;
+ gimple_omp_clauses = gimple_omp_target_clauses;
+ break;
+ default:
+ gcc_unreachable ();
+ }
+
+ if (is_gimple_omp_oacc_specifically (stmt))
+ {
+ gcc_assert (taskreg_nesting_level == 0);
+ gcc_assert (target_nesting_level == 0);
+ }
ctx = new_omp_context (stmt, outer_ctx);
ctx->field_map = splay_tree_new (splay_tree_compare_pointers, 0, 0);
@@ -2373,13 +2663,17 @@ scan_omp_target (gomp_target *stmt, omp_context *outer_ctx)
DECL_ARTIFICIAL (name) = 1;
DECL_NAMELESS (name) = 1;
TYPE_NAME (ctx->record_type) = name;
- if (kind == GF_OMP_TARGET_KIND_REGION)
+ if (offloaded)
{
+ if (is_gimple_omp_oacc_specifically (stmt))
+ ctx->reduction_map = splay_tree_new (splay_tree_compare_pointers,
+ 0, 0);
+
create_omp_child_function (ctx, false);
- gimple_omp_target_set_child_fn (stmt, ctx->cb.dst_fn);
+ gimple_omp_set_child_fn (stmt, ctx->cb.dst_fn);
}
- scan_sharing_clauses (gimple_omp_target_clauses (stmt), ctx);
+ scan_sharing_clauses (gimple_omp_clauses (stmt), ctx);
scan_omp (gimple_omp_body_ptr (stmt), ctx);
if (TYPE_FIELDS (ctx->record_type) == NULL)
@@ -2397,7 +2691,7 @@ scan_omp_target (gomp_target *stmt, omp_context *outer_ctx)
gcc_assert (DECL_ALIGN (field) == align);
#endif
layout_type (ctx->record_type);
- if (kind == GF_OMP_TARGET_KIND_REGION)
+ if (offloaded)
fixup_child_record_type (ctx);
}
}
@@ -2416,6 +2710,43 @@ scan_omp_teams (gomp_teams *stmt, omp_context *outer_ctx)
static bool
check_omp_nesting_restrictions (gimple stmt, omp_context *ctx)
{
+ /* While the OpenACC specification does allow for certain kinds of
+ nesting, we don't support many of these yet. */
+ if (is_gimple_omp (stmt)
+ && is_gimple_omp_oacc_specifically (stmt))
+ {
+ /* Regular handling of OpenACC loop constructs. */
+ if (gimple_code (stmt) == GIMPLE_OMP_FOR
+ && gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_OACC_LOOP)
+ goto cont;
+ /* No nesting of OpenACC STMT inside any OpenACC or OpenMP CTX different
+ from an OpenACC data construct. */
+ for (omp_context *ctx_ = ctx; ctx_ != NULL; ctx_ = ctx_->outer)
+ if (is_gimple_omp (ctx_->stmt)
+ && !(gimple_code (ctx_->stmt) == GIMPLE_OMP_TARGET
+ && (gimple_omp_target_kind (ctx_->stmt)
+ == GF_OMP_TARGET_KIND_OACC_DATA)))
+ {
+ error_at (gimple_location (stmt),
+ "may not be nested");
+ return false;
+ }
+ }
+ else
+ {
+ /* No nesting of non-OpenACC STMT (that is, an OpenMP one, or a GOMP
+ builtin) inside any OpenACC CTX. */
+ for (omp_context *ctx_ = ctx; ctx_ != NULL; ctx_ = ctx_->outer)
+ if (is_gimple_omp (ctx_->stmt)
+ && is_gimple_omp_oacc_specifically (ctx_->stmt))
+ {
+ error_at (gimple_location (stmt),
+ "may not be nested");
+ return false;
+ }
+ }
+ cont:
+
if (ctx != NULL)
{
if (gimple_code (ctx->stmt) == GIMPLE_OMP_FOR
@@ -2594,6 +2925,8 @@ check_omp_nesting_restrictions (gimple stmt, omp_context *ctx)
"of work-sharing, critical, ordered, master or explicit "
"task region");
return false;
+ case GIMPLE_OACC_KERNELS:
+ case GIMPLE_OACC_PARALLEL:
case GIMPLE_OMP_PARALLEL:
return true;
default:
@@ -2854,8 +3187,10 @@ scan_omp_1_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
scan_omp (gimple_omp_body_ptr (stmt), ctx);
break;
+ case GIMPLE_OACC_KERNELS:
+ case GIMPLE_OACC_PARALLEL:
case GIMPLE_OMP_TARGET:
- scan_omp_target (as_a <gomp_target *> (stmt), ctx);
+ scan_omp_target (stmt, ctx);
break;
case GIMPLE_OMP_TEAMS:
@@ -3150,6 +3485,8 @@ static bool
lower_rec_simd_input_clauses (tree new_var, omp_context *ctx, int &max_vf,
tree &idx, tree &lane, tree &ivar, tree &lvar)
{
+ gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt));
+
if (max_vf == 0)
{
max_vf = omp_max_vf ();
@@ -4069,6 +4406,57 @@ lower_lastprivate_clauses (tree clauses, tree predicate, gimple_seq *stmt_list,
gimple_seq_add_stmt (stmt_list, gimple_build_label (label));
}
+static void
+lower_reduction_var_helper (gimple_seq *stmt_seqp, omp_context *ctx, tree tid,
+ tree var, tree new_var)
+{
+ /* The atomic add at the end of the sum creates unnecessary
+ write contention on accelerators. To work around this,
+ create an array to store the partial reductions. Later, in
+ lower_omp_for (for openacc), the values of array will be
+ combined. */
+
+ tree t = NULL_TREE, array, x;
+ tree type = get_base_type (var);
+ gimple stmt;
+
+ /* Now insert the partial reductions into the array. */
+
+ /* Find the reduction array. */
+
+ tree ptype = build_pointer_type (type);
+
+ t = lookup_reduction (omp_get_id (var), ctx);
+ t = build_receiver_ref (t, false, ctx->outer);
+
+ array = create_tmp_var (ptype, NULL);
+ gimplify_assign (array, t, stmt_seqp);
+
+ tree ptr = create_tmp_var (TREE_TYPE (array), NULL);
+
+ /* Find the reduction array. */
+
+ /* testing a unary conversion. */
+ tree offset = create_tmp_var (sizetype, NULL);
+ gimplify_assign (offset, TYPE_SIZE_UNIT (type),
+ stmt_seqp);
+ t = create_tmp_var (sizetype, NULL);
+ gimplify_assign (t, unshare_expr (fold_build1 (NOP_EXPR, sizetype, tid)),
+ stmt_seqp);
+ stmt = gimple_build_assign_with_ops (MULT_EXPR, offset, offset, t);
+ gimple_seq_add_stmt (stmt_seqp, stmt);
+
+ /* Offset expression. Does the POINTER_PLUS_EXPR take care
+ of adding sizeof(var) to the array? */
+ ptr = create_tmp_var (ptype, NULL);
+ stmt = gimple_build_assign_with_ops (POINTER_PLUS_EXPR, unshare_expr(ptr),
+ array, offset);
+ gimple_seq_add_stmt (stmt_seqp, stmt);
+
+ /* Move the local sum to gfc$sum[i]. */
+ x = unshare_expr (build_simple_mem_ref (ptr));
+ stmt = gimplify_assign (x, new_var, stmt_seqp);
+}
/* Generate code to implement the REDUCTION clauses. */
@@ -4077,7 +4465,7 @@ lower_reduction_clauses (tree clauses, gimple_seq *stmt_seqp, omp_context *ctx)
{
gimple_seq sub_seq = NULL;
gimple stmt;
- tree x, c;
+ tree x, c, tid = NULL_TREE;
int count = 0;
/* SIMD reductions are handled in lower_rec_input_clauses. */
@@ -4102,6 +4490,17 @@ lower_reduction_clauses (tree clauses, gimple_seq *stmt_seqp, omp_context *ctx)
if (count == 0)
return;
+ /* Initialize thread info for OpenACC. */
+ if (is_gimple_omp_oacc_specifically (ctx->stmt))
+ {
+ /* Get the current thread id. */
+ tree call = builtin_decl_explicit (BUILT_IN_GOACC_GET_THREAD_NUM);
+ tid = create_tmp_var (TREE_TYPE (TREE_TYPE (call)), NULL);
+ gimple stmt = gimple_build_call (call, 0);
+ gimple_call_set_lhs (stmt, tid);
+ gimple_seq_add_stmt (stmt_seqp, stmt);
+ }
+
for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
{
tree var, ref, new_var;
@@ -4123,7 +4522,13 @@ lower_reduction_clauses (tree clauses, gimple_seq *stmt_seqp, omp_context *ctx)
if (code == MINUS_EXPR)
code = PLUS_EXPR;
- if (count == 1)
+ if (is_gimple_omp_oacc_specifically (ctx->stmt))
+ {
+ gcc_assert (!OMP_CLAUSE_REDUCTION_PLACEHOLDER (c));
+
+ lower_reduction_var_helper (stmt_seqp, ctx, tid, var, new_var);
+ }
+ else if (count == 1)
{
tree addr = build_fold_addr_expr_loc (clause_loc, ref);
@@ -4134,8 +4539,7 @@ lower_reduction_clauses (tree clauses, gimple_seq *stmt_seqp, omp_context *ctx)
gimplify_and_add (x, stmt_seqp);
return;
}
-
- if (OMP_CLAUSE_REDUCTION_PLACEHOLDER (c))
+ else if (OMP_CLAUSE_REDUCTION_PLACEHOLDER (c))
{
tree placeholder = OMP_CLAUSE_REDUCTION_PLACEHOLDER (c);
@@ -4158,6 +4562,9 @@ lower_reduction_clauses (tree clauses, gimple_seq *stmt_seqp, omp_context *ctx)
}
}
+ if (is_gimple_omp_oacc_specifically (ctx->stmt))
+ return;
+
stmt = gimple_build_call (builtin_decl_explicit (BUILT_IN_GOMP_ATOMIC_START),
0);
gimple_seq_add_stmt (stmt_seqp, stmt);
@@ -4176,6 +4583,8 @@ static void
lower_copyprivate_clauses (tree clauses, gimple_seq *slist, gimple_seq *rlist,
omp_context *ctx)
{
+ gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt));
+
tree c;
for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
@@ -4226,6 +4635,8 @@ static void
lower_send_clauses (tree clauses, gimple_seq *ilist, gimple_seq *olist,
omp_context *ctx)
{
+ gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt));
+
tree c;
for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
@@ -4317,6 +4728,8 @@ lower_send_clauses (tree clauses, gimple_seq *ilist, gimple_seq *olist,
static void
lower_send_shared_vars (gimple_seq *ilist, gimple_seq *olist, omp_context *ctx)
{
+ gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt));
+
tree var, ovar, nvar, f, x, record_type;
if (ctx->record_type == NULL)
@@ -5618,6 +6031,8 @@ expand_omp_for_generic (struct omp_region *region,
enum built_in_function next_fn,
gimple inner_stmt)
{
+ gcc_assert (gimple_omp_for_kind (fd->for_stmt) != GF_OMP_FOR_KIND_OACC_LOOP);
+
tree type, istart0, iend0, iend;
tree t, vmain, vback, bias = NULL_TREE;
basic_block entry_bb, cont_bb, exit_bb, l0_bb, l1_bb, collapse_bb;
@@ -5687,6 +6102,9 @@ expand_omp_for_generic (struct omp_region *region,
gcc_assert (gimple_code (gsi_stmt (gsi)) == GIMPLE_OMP_FOR);
if (fd->collapse > 1)
{
+ gcc_assert (gimple_omp_for_kind (gsi_stmt (gsi))
+ != GF_OMP_FOR_KIND_OACC_LOOP);
+
int first_zero_iter = -1;
basic_block zero_iter_bb = NULL, l2_dom_bb = NULL;
@@ -5715,6 +6133,9 @@ expand_omp_for_generic (struct omp_region *region,
}
if (in_combined_parallel)
{
+ gcc_assert (gimple_omp_for_kind (gsi_stmt (gsi))
+ != GF_OMP_FOR_KIND_OACC_LOOP);
+
/* In a combined parallel loop, emit a call to
GOMP_loop_foo_next. */
t = build_call_expr (builtin_decl_explicit (next_fn), 2,
@@ -5733,6 +6154,9 @@ expand_omp_for_generic (struct omp_region *region,
t0 = fd->loop.n1;
if (gimple_omp_for_combined_into_p (fd->for_stmt))
{
+ gcc_assert (gimple_omp_for_kind (gsi_stmt (gsi))
+ != GF_OMP_FOR_KIND_OACC_LOOP);
+
tree innerc = find_omp_clause (gimple_omp_for_clauses (fd->for_stmt),
OMP_CLAUSE__LOOPTEMP_);
gcc_assert (innerc);
@@ -6048,12 +6472,14 @@ expand_omp_for_static_nochunk (struct omp_region *region,
basic_block fin_bb;
gimple_stmt_iterator gsi;
edge ep;
- enum built_in_function get_num_threads = BUILT_IN_OMP_GET_NUM_THREADS;
- enum built_in_function get_thread_num = BUILT_IN_OMP_GET_THREAD_NUM;
bool broken_loop = region->cont == NULL;
tree *counts = NULL;
tree n1, n2, step;
+ gcc_assert ((gimple_omp_for_kind (fd->for_stmt)
+ != GF_OMP_FOR_KIND_OACC_LOOP)
+ || !inner_stmt);
+
itype = type = TREE_TYPE (fd->loop.v);
if (POINTER_TYPE_P (type))
itype = signed_type_for (type);
@@ -6077,12 +6503,6 @@ expand_omp_for_static_nochunk (struct omp_region *region,
gsi = gsi_last_bb (entry_bb);
gcc_assert (gimple_code (gsi_stmt (gsi)) == GIMPLE_OMP_FOR);
- if (gimple_omp_for_kind (fd->for_stmt) == GF_OMP_FOR_KIND_DISTRIBUTE)
- {
- get_num_threads = BUILT_IN_OMP_GET_NUM_TEAMS;
- get_thread_num = BUILT_IN_OMP_GET_TEAM_NUM;
- }
-
if (fd->collapse > 1)
{
int first_zero_iter = -1;
@@ -6141,14 +6561,30 @@ expand_omp_for_static_nochunk (struct omp_region *region,
gsi = gsi_last_bb (entry_bb);
}
- t = build_call_expr (builtin_decl_explicit (get_num_threads), 0);
- t = fold_convert (itype, t);
- nthreads = force_gimple_operand_gsi (&gsi, t, true, NULL_TREE,
+ switch (gimple_omp_for_kind (fd->for_stmt))
+ {
+ case GF_OMP_FOR_KIND_FOR:
+ nthreads = builtin_decl_explicit (BUILT_IN_OMP_GET_NUM_THREADS);
+ threadid = builtin_decl_explicit (BUILT_IN_OMP_GET_THREAD_NUM);
+ break;
+ case GF_OMP_FOR_KIND_DISTRIBUTE:
+ nthreads = builtin_decl_explicit (BUILT_IN_OMP_GET_NUM_TEAMS);
+ threadid = builtin_decl_explicit (BUILT_IN_OMP_GET_TEAM_NUM);
+ break;
+ case GF_OMP_FOR_KIND_OACC_LOOP:
+ nthreads = builtin_decl_explicit (BUILT_IN_GOACC_GET_NUM_THREADS);
+ threadid = builtin_decl_explicit (BUILT_IN_GOACC_GET_THREAD_NUM);
+ break;
+ default:
+ gcc_unreachable ();
+ }
+ nthreads = build_call_expr (nthreads, 0);
+ nthreads = fold_convert (itype, nthreads);
+ nthreads = force_gimple_operand_gsi (&gsi, nthreads, true, NULL_TREE,
true, GSI_SAME_STMT);
-
- t = build_call_expr (builtin_decl_explicit (get_thread_num), 0);
- t = fold_convert (itype, t);
- threadid = force_gimple_operand_gsi (&gsi, t, true, NULL_TREE,
+ threadid = build_call_expr (threadid, 0);
+ threadid = fold_convert (itype, threadid);
+ threadid = force_gimple_operand_gsi (&gsi, threadid, true, NULL_TREE,
true, GSI_SAME_STMT);
n1 = fd->loop.n1;
@@ -6156,6 +6592,9 @@ expand_omp_for_static_nochunk (struct omp_region *region,
step = fd->loop.step;
if (gimple_omp_for_combined_into_p (fd->for_stmt))
{
+ gcc_assert (gimple_omp_for_kind (fd->for_stmt)
+ != GF_OMP_FOR_KIND_OACC_LOOP);
+
tree innerc = find_omp_clause (gimple_omp_for_clauses (fd->for_stmt),
OMP_CLAUSE__LOOPTEMP_);
gcc_assert (innerc);
@@ -6323,7 +6762,10 @@ expand_omp_for_static_nochunk (struct omp_region *region,
if (!gimple_omp_return_nowait_p (gsi_stmt (gsi)))
{
t = gimple_omp_return_lhs (gsi_stmt (gsi));
- gsi_insert_after (&gsi, build_omp_barrier (t), GSI_SAME_STMT);
+ if (gimple_omp_for_kind (fd->for_stmt) == GF_OMP_FOR_KIND_OACC_LOOP)
+ gcc_assert (t == NULL_TREE);
+ else
+ gsi_insert_after (&gsi, build_omp_barrier (t), GSI_SAME_STMT);
}
gsi_remove (&gsi, true);
@@ -6341,6 +6783,9 @@ expand_omp_for_static_nochunk (struct omp_region *region,
ep = find_edge (cont_bb, body_bb);
if (gimple_omp_for_combined_p (fd->for_stmt))
{
+ gcc_assert (gimple_omp_for_kind (fd->for_stmt)
+ != GF_OMP_FOR_KIND_OACC_LOOP);
+
remove_edge (ep);
ep = NULL;
}
@@ -6424,12 +6869,14 @@ expand_omp_for_static_chunk (struct omp_region *region,
basic_block trip_update_bb = NULL, cont_bb, collapse_bb = NULL, fin_bb;
gimple_stmt_iterator gsi;
edge se;
- enum built_in_function get_num_threads = BUILT_IN_OMP_GET_NUM_THREADS;
- enum built_in_function get_thread_num = BUILT_IN_OMP_GET_THREAD_NUM;
bool broken_loop = region->cont == NULL;
tree *counts = NULL;
tree n1, n2, step;
+ gcc_assert ((gimple_omp_for_kind (fd->for_stmt)
+ != GF_OMP_FOR_KIND_OACC_LOOP)
+ || !inner_stmt);
+
itype = type = TREE_TYPE (fd->loop.v);
if (POINTER_TYPE_P (type))
itype = signed_type_for (type);
@@ -6457,12 +6904,6 @@ expand_omp_for_static_chunk (struct omp_region *region,
gsi = gsi_last_bb (entry_bb);
gcc_assert (gimple_code (gsi_stmt (gsi)) == GIMPLE_OMP_FOR);
- if (gimple_omp_for_kind (fd->for_stmt) == GF_OMP_FOR_KIND_DISTRIBUTE)
- {
- get_num_threads = BUILT_IN_OMP_GET_NUM_TEAMS;
- get_thread_num = BUILT_IN_OMP_GET_TEAM_NUM;
- }
-
if (fd->collapse > 1)
{
int first_zero_iter = -1;
@@ -6521,14 +6962,30 @@ expand_omp_for_static_chunk (struct omp_region *region,
gsi = gsi_last_bb (entry_bb);
}
- t = build_call_expr (builtin_decl_explicit (get_num_threads), 0);
- t = fold_convert (itype, t);
- nthreads = force_gimple_operand_gsi (&gsi, t, true, NULL_TREE,
+ switch (gimple_omp_for_kind (fd->for_stmt))
+ {
+ case GF_OMP_FOR_KIND_FOR:
+ nthreads = builtin_decl_explicit (BUILT_IN_OMP_GET_NUM_THREADS);
+ threadid = builtin_decl_explicit (BUILT_IN_OMP_GET_THREAD_NUM);
+ break;
+ case GF_OMP_FOR_KIND_DISTRIBUTE:
+ nthreads = builtin_decl_explicit (BUILT_IN_OMP_GET_NUM_TEAMS);
+ threadid = builtin_decl_explicit (BUILT_IN_OMP_GET_TEAM_NUM);
+ break;
+ case GF_OMP_FOR_KIND_OACC_LOOP:
+ nthreads = builtin_decl_explicit (BUILT_IN_GOACC_GET_NUM_THREADS);
+ threadid = builtin_decl_explicit (BUILT_IN_GOACC_GET_THREAD_NUM);
+ break;
+ default:
+ gcc_unreachable ();
+ }
+ nthreads = build_call_expr (nthreads, 0);
+ nthreads = fold_convert (itype, nthreads);
+ nthreads = force_gimple_operand_gsi (&gsi, nthreads, true, NULL_TREE,
true, GSI_SAME_STMT);
-
- t = build_call_expr (builtin_decl_explicit (get_thread_num), 0);
- t = fold_convert (itype, t);
- threadid = force_gimple_operand_gsi (&gsi, t, true, NULL_TREE,
+ threadid = build_call_expr (threadid, 0);
+ threadid = fold_convert (itype, threadid);
+ threadid = force_gimple_operand_gsi (&gsi, threadid, true, NULL_TREE,
true, GSI_SAME_STMT);
n1 = fd->loop.n1;
@@ -6536,6 +6993,9 @@ expand_omp_for_static_chunk (struct omp_region *region,
step = fd->loop.step;
if (gimple_omp_for_combined_into_p (fd->for_stmt))
{
+ gcc_assert (gimple_omp_for_kind (fd->for_stmt)
+ != GF_OMP_FOR_KIND_OACC_LOOP);
+
tree innerc = find_omp_clause (gimple_omp_for_clauses (fd->for_stmt),
OMP_CLAUSE__LOOPTEMP_);
gcc_assert (innerc);
@@ -6719,7 +7179,10 @@ expand_omp_for_static_chunk (struct omp_region *region,
if (!gimple_omp_return_nowait_p (gsi_stmt (gsi)))
{
t = gimple_omp_return_lhs (gsi_stmt (gsi));
- gsi_insert_after (&gsi, build_omp_barrier (t), GSI_SAME_STMT);
+ if (gimple_omp_for_kind (fd->for_stmt) == GF_OMP_FOR_KIND_OACC_LOOP)
+ gcc_assert (t == NULL_TREE);
+ else
+ gsi_insert_after (&gsi, build_omp_barrier (t), GSI_SAME_STMT);
}
gsi_remove (&gsi, true);
@@ -6732,6 +7195,9 @@ expand_omp_for_static_chunk (struct omp_region *region,
se = find_edge (cont_bb, body_bb);
if (gimple_omp_for_combined_p (fd->for_stmt))
{
+ gcc_assert (gimple_omp_for_kind (fd->for_stmt)
+ != GF_OMP_FOR_KIND_OACC_LOOP);
+
remove_edge (se);
se = NULL;
}
@@ -8274,7 +8740,7 @@ expand_omp_atomic (struct omp_region *region)
}
-/* Expand the OpenMP target{, data, update} directive starting at REGION. */
+/* Expand the GIMPLE_OMP_TARGET starting at REGION. */
static void
expand_omp_target (struct omp_region *region)
@@ -8283,47 +8749,91 @@ expand_omp_target (struct omp_region *region)
struct function *child_cfun = NULL;
tree child_fn = NULL_TREE, block, t;
gimple_stmt_iterator gsi;
- gomp_target *entry_stmt;
+ gimple entry_stmt;
gimple stmt;
edge e;
+ bool offloaded, data_region;
+ tree (*gimple_omp_child_fn) (const_gimple);
+ tree (*gimple_omp_data_arg) (const_gimple);
- entry_stmt = as_a <gomp_target *> (last_stmt (region->entry));
+ entry_stmt = last_stmt (region->entry);
new_bb = region->entry;
- int kind = gimple_omp_target_kind (entry_stmt);
- if (kind == GF_OMP_TARGET_KIND_REGION)
+
+ offloaded = is_gimple_omp_offloaded (entry_stmt);
+ data_region = false;
+ switch (region->type)
{
- child_fn = gimple_omp_target_child_fn (entry_stmt);
+ case GIMPLE_OACC_KERNELS:
+ gimple_omp_child_fn = gimple_oacc_kernels_child_fn;
+ gimple_omp_data_arg = gimple_oacc_kernels_data_arg;
+ break;
+ case GIMPLE_OACC_PARALLEL:
+ gimple_omp_child_fn = gimple_oacc_parallel_child_fn;
+ gimple_omp_data_arg = gimple_oacc_parallel_data_arg;
+ break;
+ case GIMPLE_OMP_TARGET:
+ switch (gimple_omp_target_kind (entry_stmt))
+ {
+ case GF_OMP_TARGET_KIND_DATA:
+ case GF_OMP_TARGET_KIND_OACC_DATA:
+ data_region = true;
+ break;
+ case GF_OMP_TARGET_KIND_REGION:
+ case GF_OMP_TARGET_KIND_UPDATE:
+ case GF_OMP_TARGET_KIND_OACC_UPDATE:
+ case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ break;
+ default:
+ gcc_unreachable ();
+ }
+
+ gimple_omp_child_fn = gimple_omp_target_child_fn;
+ gimple_omp_data_arg = gimple_omp_target_data_arg;
+ break;
+ default:
+ gcc_unreachable ();
+ }
+
+ child_fn = NULL_TREE;
+ child_cfun = NULL;
+ if (offloaded)
+ {
+ child_fn = gimple_omp_child_fn (entry_stmt);
child_cfun = DECL_STRUCT_FUNCTION (child_fn);
}
+ /* Supported by expand_omp_taskreg, but not here. */
+ if (child_cfun != NULL)
+ gcc_assert (!child_cfun->cfg);
+ gcc_assert (!gimple_in_ssa_p (cfun));
+
entry_bb = region->entry;
exit_bb = region->exit;
- if (kind == GF_OMP_TARGET_KIND_REGION)
+ if (offloaded)
{
unsigned srcidx, dstidx, num;
- /* If the target region needs data sent from the parent
+ /* If the offloading region needs data sent from the parent
function, then the very first statement (except possible
- tree profile counter updates) of the parallel body
+ tree profile counter updates) of the offloading body
is a copy assignment .OMP_DATA_I = &.OMP_DATA_O. Since
&.OMP_DATA_O is passed as an argument to the child function,
we need to replace it with the argument as seen by the child
function.
In most cases, this will end up being the identity assignment
- .OMP_DATA_I = .OMP_DATA_I. However, if the parallel body had
+ .OMP_DATA_I = .OMP_DATA_I. However, if the offloading body had
a function call that has been inlined, the original PARM_DECL
.OMP_DATA_I may have been converted into a different local
variable. In which case, we need to keep the assignment. */
- if (gimple_omp_target_data_arg (entry_stmt))
+ if (gimple_omp_data_arg (entry_stmt))
{
basic_block entry_succ_bb = single_succ (entry_bb);
gimple_stmt_iterator gsi;
tree arg;
gimple tgtcopy_stmt = NULL;
- tree sender
- = TREE_VEC_ELT (gimple_omp_target_data_arg (entry_stmt), 0);
+ tree sender = TREE_VEC_ELT (gimple_omp_data_arg (entry_stmt), 0);
for (gsi = gsi_start_bb (entry_succ_bb); ; gsi_next (&gsi))
{
@@ -8358,7 +8868,7 @@ expand_omp_target (struct omp_region *region)
/* Declare local variables needed in CHILD_CFUN. */
block = DECL_INITIAL (child_fn);
BLOCK_VARS (block) = vec2chain (child_cfun->local_decls);
- /* The gimplifier could record temporaries in target block
+ /* The gimplifier could record temporaries in the offloading block
rather than in containing function's local_decls chain,
which would mean cgraph missed finalizing them. Do it now. */
for (t = BLOCK_VARS (block); t; t = DECL_CHAIN (t))
@@ -8375,13 +8885,12 @@ expand_omp_target (struct omp_region *region)
for (t = DECL_ARGUMENTS (child_fn); t; t = DECL_CHAIN (t))
DECL_CONTEXT (t) = child_fn;
- /* Split ENTRY_BB at GIMPLE_OMP_TARGET,
+ /* Split ENTRY_BB at GIMPLE_*,
so that it can be moved to the child function. */
gsi = gsi_last_bb (entry_bb);
stmt = gsi_stmt (gsi);
- gcc_assert (stmt && gimple_code (stmt) == GIMPLE_OMP_TARGET
- && gimple_omp_target_kind (stmt)
- == GF_OMP_TARGET_KIND_REGION);
+ gcc_assert (stmt &&
+ gimple_code (stmt) == gimple_code (entry_stmt));
gsi_remove (&gsi, true);
e = split_block (entry_bb, stmt);
entry_bb = e->dest;
@@ -8398,7 +8907,7 @@ expand_omp_target (struct omp_region *region)
gsi_remove (&gsi, true);
}
- /* Move the target region into CHILD_CFUN. */
+ /* Move the offloading region into CHILD_CFUN. */
block = gimple_block (entry_stmt);
@@ -8433,6 +8942,9 @@ expand_omp_target (struct omp_region *region)
vec_safe_push (offload_funcs, child_fn);
#endif
+ /* Add the new function to the offload table. */
+ vec_safe_push (offload_funcs, child_fn);
+
/* Fix the callgraph edges for child_cfun. Those for cfun will be
fixed in a following pass. */
push_cfun (child_cfun);
@@ -8462,20 +8974,54 @@ expand_omp_target (struct omp_region *region)
pop_cfun ();
}
- /* Emit a library call to launch the target region, or do data
+ /* Emit a library call to launch the offloading region, or do data
transfers. */
tree t1, t2, t3, t4, device, cond, c, clauses;
enum built_in_function start_ix;
location_t clause_loc;
+ tree (*gimple_omp_clauses) (const_gimple);
- clauses = gimple_omp_target_clauses (entry_stmt);
+ switch (region->type)
+ {
+ case GIMPLE_OACC_KERNELS:
+ gimple_omp_clauses = gimple_oacc_kernels_clauses;
+ start_ix = BUILT_IN_GOACC_KERNELS;
+ break;
+ case GIMPLE_OACC_PARALLEL:
+ gimple_omp_clauses = gimple_oacc_parallel_clauses;
+ start_ix = BUILT_IN_GOACC_PARALLEL;
+ break;
+ case GIMPLE_OMP_TARGET:
+ gimple_omp_clauses = gimple_omp_target_clauses;
+ switch (gimple_omp_target_kind (entry_stmt))
+ {
+ case GF_OMP_TARGET_KIND_REGION:
+ start_ix = BUILT_IN_GOMP_TARGET;
+ break;
+ case GF_OMP_TARGET_KIND_DATA:
+ start_ix = BUILT_IN_GOMP_TARGET_DATA;
+ break;
+ case GF_OMP_TARGET_KIND_UPDATE:
+ start_ix = BUILT_IN_GOMP_TARGET_UPDATE;
+ break;
+ case GF_OMP_TARGET_KIND_OACC_DATA:
+ start_ix = BUILT_IN_GOACC_DATA_START;
+ break;
+ case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ start_ix = BUILT_IN_GOACC_ENTER_EXIT_DATA;
+ break;
+ case GF_OMP_TARGET_KIND_OACC_UPDATE:
+ start_ix = BUILT_IN_GOACC_UPDATE;
+ break;
+ default:
+ gcc_unreachable ();
+ }
+ break;
+ default:
+ gcc_unreachable ();
+ }
- if (kind == GF_OMP_TARGET_KIND_REGION)
- start_ix = BUILT_IN_GOMP_TARGET;
- else if (kind == GF_OMP_TARGET_KIND_DATA)
- start_ix = BUILT_IN_GOMP_TARGET_DATA;
- else
- start_ix = BUILT_IN_GOMP_TARGET_UPDATE;
+ clauses = gimple_omp_clauses (entry_stmt);
/* By default, the value of DEVICE is -1 (let runtime library choose)
and there is no conditional. */
@@ -8489,6 +9035,12 @@ expand_omp_target (struct omp_region *region)
c = find_omp_clause (clauses, OMP_CLAUSE_DEVICE);
if (c)
{
+ /* Even if we pass it to all library function calls, it is currently only
+ defined/used for the OpenMP target ones. */
+ gcc_assert (start_ix == BUILT_IN_GOMP_TARGET
+ || start_ix == BUILT_IN_GOMP_TARGET_DATA
+ || start_ix == BUILT_IN_GOMP_TARGET_UPDATE);
+
device = OMP_CLAUSE_DEVICE_ID (c);
clause_loc = OMP_CLAUSE_LOCATION (c);
}
@@ -8509,14 +9061,16 @@ expand_omp_target (struct omp_region *region)
tree tmp_var;
tmp_var = create_tmp_var (TREE_TYPE (device), NULL);
- if (kind != GF_OMP_TARGET_KIND_REGION)
+ if (offloaded)
+ {
+ e = split_block (new_bb, NULL);
+ }
+ else
{
gsi = gsi_last_bb (new_bb);
gsi_prev (&gsi);
e = split_block (new_bb, gsi_stmt (gsi));
}
- else
- e = split_block (new_bb, NULL);
cond_bb = e->src;
new_bb = e->dest;
remove_edge (e);
@@ -8550,7 +9104,7 @@ expand_omp_target (struct omp_region *region)
}
gsi = gsi_last_bb (new_bb);
- t = gimple_omp_target_data_arg (entry_stmt);
+ t = gimple_omp_data_arg (entry_stmt);
if (t == NULL)
{
t1 = size_zero_node;
@@ -8568,28 +9122,124 @@ expand_omp_target (struct omp_region *region)
}
gimple g;
- /* FIXME: This will be address of
- extern char __OPENMP_TARGET__[] __attribute__((visibility ("hidden")))
- symbol, as soon as the linker plugin is able to create it for us. */
- tree openmp_target = build_zero_cst (ptr_type_node);
- if (kind == GF_OMP_TARGET_KIND_REGION)
- {
- tree fnaddr = build_fold_addr_expr (child_fn);
- g = gimple_build_call (builtin_decl_explicit (start_ix), 7,
- device, fnaddr, openmp_target, t1, t2, t3, t4);
+ tree openmp_target = get_offload_symbol_decl ();
+ vec<tree> *args;
+ /* The maximum number used by any start_ix, without varargs. */
+ unsigned int argcnt = 12;
+
+ vec_alloc (args, argcnt);
+ args->quick_push (device);
+ if (offloaded)
+ args->quick_push (build_fold_addr_expr (child_fn));
+ args->quick_push (build_fold_addr_expr (openmp_target));
+ args->quick_push (t1);
+ args->quick_push (t2);
+ args->quick_push (t3);
+ args->quick_push (t4);
+ switch (start_ix)
+ {
+ case BUILT_IN_GOACC_DATA_START:
+ case BUILT_IN_GOMP_TARGET:
+ case BUILT_IN_GOMP_TARGET_DATA:
+ case BUILT_IN_GOMP_TARGET_UPDATE:
+ break;
+ case BUILT_IN_GOACC_KERNELS:
+ case BUILT_IN_GOACC_PARALLEL:
+ {
+ tree t_num_gangs, t_num_workers, t_vector_length;
+
+ /* Default values for num_gangs, num_workers, and vector_length. */
+ t_num_gangs = t_num_workers = t_vector_length
+ = fold_convert_loc (gimple_location (entry_stmt),
+ integer_type_node, integer_one_node);
+ /* ..., but if present, use the value specified by the respective
+ clause, making sure that are of the correct type. */
+ c = find_omp_clause (clauses, OMP_CLAUSE_NUM_GANGS);
+ if (c)
+ t_num_gangs = fold_convert_loc (OMP_CLAUSE_LOCATION (c),
+ integer_type_node,
+ OMP_CLAUSE_NUM_GANGS_EXPR (c));
+ c = find_omp_clause (clauses, OMP_CLAUSE_NUM_WORKERS);
+ if (c)
+ t_num_workers = fold_convert_loc (OMP_CLAUSE_LOCATION (c),
+ integer_type_node,
+ OMP_CLAUSE_NUM_WORKERS_EXPR (c));
+ c = find_omp_clause (clauses, OMP_CLAUSE_VECTOR_LENGTH);
+ if (c)
+ t_vector_length = fold_convert_loc (OMP_CLAUSE_LOCATION (c),
+ integer_type_node,
+ OMP_CLAUSE_VECTOR_LENGTH_EXPR (c));
+ args->quick_push (t_num_gangs);
+ args->quick_push (t_num_workers);
+ args->quick_push (t_vector_length);
+ }
+ /* FALLTHRU */
+ case BUILT_IN_GOACC_ENTER_EXIT_DATA:
+ case BUILT_IN_GOACC_UPDATE:
+ {
+ tree t_async;
+ int t_wait_idx;
+
+ /* Default values for t_async. */
+ t_async = fold_convert_loc (gimple_location (entry_stmt),
+ integer_type_node,
+ build_int_cst (integer_type_node, -2));
+ /* ..., but if present, use the value specified by the respective
+ clause, making sure that is of the correct type. */
+ c = find_omp_clause (clauses, OMP_CLAUSE_ASYNC);
+ if (c)
+ t_async = fold_convert_loc (OMP_CLAUSE_LOCATION (c),
+ integer_type_node,
+ OMP_CLAUSE_ASYNC_EXPR (c));
+
+ args->quick_push (t_async);
+ /* Save the index, and... */
+ t_wait_idx = args->length ();
+ /* ... push a default value. */
+ args->quick_push (fold_convert_loc (gimple_location (entry_stmt),
+ integer_type_node,
+ integer_zero_node));
+ c = find_omp_clause (clauses, OMP_CLAUSE_WAIT);
+ if (c)
+ {
+ int n = 0;
+
+ for (; c; c = OMP_CLAUSE_CHAIN (c))
+ {
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_WAIT)
+ {
+ args->safe_push (fold_convert_loc (OMP_CLAUSE_LOCATION (c),
+ integer_type_node,
+ OMP_CLAUSE_WAIT_EXPR (c)));
+ n++;
+ }
+ }
+
+ /* Now that we know the number, replace the default value. */
+ args->ordered_remove (t_wait_idx);
+ args->quick_insert (t_wait_idx,
+ fold_convert_loc (gimple_location (entry_stmt),
+ integer_type_node,
+ build_int_cst (integer_type_node, n)));
+ }
+ }
+ break;
+ default:
+ gcc_unreachable ();
}
- else
- g = gimple_build_call (builtin_decl_explicit (start_ix), 6,
- device, openmp_target, t1, t2, t3, t4);
+
+ g = gimple_build_call_vec (builtin_decl_explicit (start_ix), *args);
+ args->release ();
gimple_set_location (g, gimple_location (entry_stmt));
gsi_insert_before (&gsi, g, GSI_SAME_STMT);
- if (kind != GF_OMP_TARGET_KIND_REGION)
+ if (!offloaded)
{
g = gsi_stmt (gsi);
gcc_assert (g && gimple_code (g) == GIMPLE_OMP_TARGET);
gsi_remove (&gsi, true);
}
- if (kind == GF_OMP_TARGET_KIND_DATA && region->exit)
+ if (data_region
+ && region->exit)
{
gsi = gsi_last_bb (region->exit);
g = gsi_stmt (gsi);
@@ -8665,6 +9315,8 @@ expand_omp (struct omp_region *region)
expand_omp_atomic (region);
break;
+ case GIMPLE_OACC_KERNELS:
+ case GIMPLE_OACC_PARALLEL:
case GIMPLE_OMP_TARGET:
expand_omp_target (region);
break;
@@ -8734,7 +9386,9 @@ build_omp_regions_1 (basic_block bb, struct omp_region *parent,
;
}
else if (code == GIMPLE_OMP_TARGET
- && gimple_omp_target_kind (stmt) == GF_OMP_TARGET_KIND_UPDATE)
+ && (gimple_omp_target_kind (stmt) == GF_OMP_TARGET_KIND_UPDATE
+ || (gimple_omp_target_kind (stmt)
+ == GF_OMP_TARGET_KIND_OACC_UPDATE)))
new_omp_region (bb, code, parent);
else
{
@@ -8850,8 +9504,9 @@ public:
/* opt_pass methods: */
virtual unsigned int execute (function *)
{
- bool gate = ((flag_openmp != 0 || flag_openmp_simd != 0
- || flag_cilkplus != 0) && !seen_error ());
+ bool gate = ((flag_openmp != 0 || flag_openacc != 0
+ || flag_openmp_simd != 0 || flag_cilkplus != 0)
+ && !seen_error ());
/* This pass always runs, to provide PROP_gimple_eomp.
But there is nothing to do unless -fopenmp is given. */
@@ -8912,6 +9567,397 @@ make_pass_expand_omp_ssa (gcc::context *ctxt)
/* Routines to lower OpenMP directives into OMP-GIMPLE. */
+/* Helper function to preform, potentially COMPLEX_TYPE, operation and
+ convert it to gimple. */
+static void
+omp_gimple_assign_with_ops (tree_code op, tree dest, tree src, gimple_seq *seq)
+{
+ gimple stmt;
+
+ if (TREE_CODE (TREE_TYPE (dest)) != COMPLEX_TYPE)
+ {
+ stmt = gimple_build_assign_with_ops (op, dest, dest, src);
+ gimple_seq_add_stmt (seq, stmt);
+ return;
+ }
+
+ tree t = create_tmp_var (TREE_TYPE (TREE_TYPE (dest)), NULL);
+ tree rdest = fold_build1 (REALPART_EXPR, TREE_TYPE (TREE_TYPE (dest)), dest);
+ gimplify_assign (t, rdest, seq);
+ rdest = t;
+
+ t = create_tmp_var (TREE_TYPE (TREE_TYPE (dest)), NULL);
+ tree idest = fold_build1 (IMAGPART_EXPR, TREE_TYPE (TREE_TYPE (dest)), dest);
+ gimplify_assign (t, idest, seq);
+ idest = t;
+
+ t = create_tmp_var (TREE_TYPE (TREE_TYPE (src)), NULL);
+ tree rsrc = fold_build1 (REALPART_EXPR, TREE_TYPE (TREE_TYPE (src)), src);
+ gimplify_assign (t, rsrc, seq);
+ rsrc = t;
+
+ t = create_tmp_var (TREE_TYPE (TREE_TYPE (src)), NULL);
+ tree isrc = fold_build1 (IMAGPART_EXPR, TREE_TYPE (TREE_TYPE (src)), src);
+ gimplify_assign (t, isrc, seq);
+ isrc = t;
+
+ tree r = create_tmp_var (TREE_TYPE (TREE_TYPE (dest)), NULL);
+ tree i = create_tmp_var (TREE_TYPE (TREE_TYPE (dest)), NULL);
+ tree result;
+
+ gcc_assert (op == PLUS_EXPR || op == MULT_EXPR);
+
+ if (op == PLUS_EXPR)
+ {
+ stmt = gimple_build_assign_with_ops (op, r, rdest, rsrc);
+ gimple_seq_add_stmt (seq, stmt);
+
+ stmt = gimple_build_assign_with_ops (op, i, idest, isrc);
+ gimple_seq_add_stmt (seq, stmt);
+ }
+ else if (op == MULT_EXPR)
+ {
+ /* Let x = a + ib = dest, y = c + id = src.
+ x * y = (ac - bd) + i(ad + bc) */
+ tree ac = create_tmp_var (TREE_TYPE (TREE_TYPE (dest)), NULL);
+ tree bd = create_tmp_var (TREE_TYPE (TREE_TYPE (dest)), NULL);
+ tree ad = create_tmp_var (TREE_TYPE (TREE_TYPE (dest)), NULL);
+ tree bc = create_tmp_var (TREE_TYPE (TREE_TYPE (dest)), NULL);
+
+ stmt = gimple_build_assign_with_ops (MULT_EXPR, ac, rdest, rsrc);
+ gimple_seq_add_stmt (seq, stmt);
+
+ stmt = gimple_build_assign_with_ops (MULT_EXPR, bd, idest, isrc);
+ gimple_seq_add_stmt (seq, stmt);
+
+ stmt = gimple_build_assign_with_ops (MINUS_EXPR, r, ac, bd);
+ gimple_seq_add_stmt (seq, stmt);
+
+ stmt = gimple_build_assign_with_ops (MULT_EXPR, ad, rdest, isrc);
+ gimple_seq_add_stmt (seq, stmt);
+
+ stmt = gimple_build_assign_with_ops (MULT_EXPR, bd, idest, rsrc);
+ gimple_seq_add_stmt (seq, stmt);
+
+ stmt = gimple_build_assign_with_ops (PLUS_EXPR, i, ad, bc);
+ gimple_seq_add_stmt (seq, stmt);
+ }
+
+ result = build2 (COMPLEX_EXPR, TREE_TYPE (dest), r, i);
+ gimplify_assign (dest, result, seq);
+}
+
+/* Helper function to initialize local data for the reduction arrays.
+ The reduction arrays need to be placed inside the calling function
+ for accelerators, or else the host won't be able to preform the final
+ reduction. */
+
+static void
+initialize_reduction_data (tree clauses, tree nthreads, gimple_seq *stmt_seqp,
+ omp_context *ctx)
+{
+ gcc_assert (is_gimple_omp_oacc_specifically (ctx->stmt));
+
+ tree c, t, oc;
+ gimple stmt;
+ omp_context *octx;
+ tree (*gimple_omp_clauses) (const_gimple);
+ void (*gimple_omp_set_clauses) (gimple, tree);
+
+ /* Find the innermost GIMPLE_OACC_PARALLEL ctx. */
+ if (gimple_code (ctx->stmt) == GIMPLE_OACC_PARALLEL)
+ octx = ctx;
+ else
+ octx = ctx->outer;
+ gcc_assert (gimple_code (octx->stmt) == GIMPLE_OACC_PARALLEL);
+
+ gimple_omp_clauses = gimple_oacc_parallel_clauses;
+ gimple_omp_set_clauses = gimple_oacc_parallel_set_clauses;
+
+ /* Extract the clauses. */
+ oc = gimple_omp_clauses (octx->stmt);
+
+ /* Find the last outer clause. */
+ for (; oc && OMP_CLAUSE_CHAIN (oc); oc = OMP_CLAUSE_CHAIN (oc))
+ ;
+
+ /* Allocate arrays for each reduction variable. */
+ for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+ {
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_REDUCTION)
+ continue;
+
+ tree var = OMP_CLAUSE_DECL (c);
+ tree type = get_base_type (var);
+ tree array = lookup_reduction (omp_get_id (var), ctx);
+ tree size, call;
+
+ /* Calculate size of the reduction array. */
+ t = create_tmp_var (TREE_TYPE (nthreads), NULL);
+ stmt = gimple_build_assign_with_ops (MULT_EXPR, t, nthreads,
+ fold_convert (TREE_TYPE (nthreads),
+ TYPE_SIZE_UNIT (type)));
+ gimple_seq_add_stmt (stmt_seqp, stmt);
+
+ size = create_tmp_var (sizetype, NULL);
+ gimplify_assign (size, fold_build1 (NOP_EXPR, sizetype, t), stmt_seqp);
+
+ /* Now allocate memory for it. */
+ call = unshare_expr (builtin_decl_explicit (BUILT_IN_ALLOCA));
+ stmt = gimple_build_call (call, 1, size);
+ gimple_call_set_lhs (stmt, array);
+ gimple_seq_add_stmt (stmt_seqp, stmt);
+
+ /* Map this array into the accelerator. */
+
+ /* Add the reduction array to the list of clauses. */
+ tree x = array;
+ t = build_omp_clause (gimple_location (ctx->stmt), OMP_CLAUSE_MAP);
+ OMP_CLAUSE_MAP_KIND (t) = OMP_CLAUSE_MAP_FORCE_FROM;
+ OMP_CLAUSE_DECL (t) = x;
+ OMP_CLAUSE_CHAIN (t) = NULL;
+ if (oc)
+ OMP_CLAUSE_CHAIN (oc) = t;
+ else
+ gimple_omp_set_clauses (octx->stmt, t);
+ OMP_CLAUSE_SIZE (t) = size;
+ oc = t;
+ }
+}
+
+/* Helper function to process the array of partial reductions. Nthreads
+ indicates the number of threads. Unfortunately, GOACC_GET_NUM_THREADS
+ cannot be used here, because nthreads on the host may be different than
+ on the accelerator. */
+
+static void
+finalize_reduction_data (tree clauses, tree nthreads, gimple_seq *stmt_seqp,
+ omp_context *ctx)
+{
+ gcc_assert (is_gimple_omp_oacc_specifically (ctx->stmt));
+
+ tree c, x, var, array, loop_header, loop_body, loop_exit, type;
+ gimple stmt;
+
+ /* Create for loop.
+
+ let var = the original reduction variable
+ let array = reduction variable array
+
+ for (i = 0; i < nthreads; i++)
+ var op= array[i]
+ */
+
+ loop_header = create_artificial_label (UNKNOWN_LOCATION);
+ loop_body = create_artificial_label (UNKNOWN_LOCATION);
+ loop_exit = create_artificial_label (UNKNOWN_LOCATION);
+
+ /* Create and initialize an index variable. */
+ tree ix = create_tmp_var (sizetype, NULL);
+ gimplify_assign (ix, fold_build1 (NOP_EXPR, sizetype, integer_zero_node),
+ stmt_seqp);
+
+ /* Insert the loop header label here. */
+ gimple_seq_add_stmt (stmt_seqp, gimple_build_label (loop_header));
+
+ /* Exit loop if ix >= nthreads. */
+ x = create_tmp_var (sizetype, NULL);
+ gimplify_assign (x, fold_build1 (NOP_EXPR, sizetype, nthreads), stmt_seqp);
+ stmt = gimple_build_cond (GE_EXPR, ix, x, loop_exit, loop_body);
+ gimple_seq_add_stmt (stmt_seqp, stmt);
+
+ /* Insert the loop body label here. */
+ gimple_seq_add_stmt (stmt_seqp, gimple_build_label (loop_body));
+
+ /* Collapse each reduction array, one element at a time. */
+ for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+ {
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_REDUCTION)
+ continue;
+
+ tree_code reduction_code = OMP_CLAUSE_REDUCTION_CODE (c);
+
+ /* reduction(-:var) sums up the partial results, so it acts
+ identically to reduction(+:var). */
+ if (reduction_code == MINUS_EXPR)
+ reduction_code = PLUS_EXPR;
+
+ /* Set up reduction variable var. */
+ var = OMP_CLAUSE_DECL (c);
+ type = get_base_type (var);
+ array = lookup_reduction (omp_get_id (OMP_CLAUSE_DECL (c)), ctx);
+
+ /* Calculate the array offset. */
+ tree offset = create_tmp_var (sizetype, NULL);
+ gimplify_assign (offset, TYPE_SIZE_UNIT (type), stmt_seqp);
+ stmt = gimple_build_assign_with_ops (MULT_EXPR, offset, offset, ix);
+ gimple_seq_add_stmt (stmt_seqp, stmt);
+
+ tree ptr = create_tmp_var (TREE_TYPE (array), NULL);
+ stmt = gimple_build_assign_with_ops (POINTER_PLUS_EXPR, ptr, array,
+ offset);
+ gimple_seq_add_stmt (stmt_seqp, stmt);
+
+ /* Extract array[ix] into mem. */
+ tree mem = create_tmp_var (type, NULL);
+ gimplify_assign (mem, build_simple_mem_ref (ptr), stmt_seqp);
+
+ /* Find the original reduction variable. */
+ if (is_reference (var))
+ var = build_simple_mem_ref (var);
+
+ tree t = create_tmp_var (type, NULL);
+
+ x = lang_hooks.decls.omp_clause_assign_op (c, t, var);
+ gimplify_and_add (unshare_expr(x), stmt_seqp);
+
+ /* var = var op mem */
+ switch (OMP_CLAUSE_REDUCTION_CODE (c))
+ {
+ case TRUTH_ANDIF_EXPR:
+ case TRUTH_ORIF_EXPR:
+ t = fold_build2 (OMP_CLAUSE_REDUCTION_CODE (c), integer_type_node,
+ t, mem);
+ gimplify_and_add (t, stmt_seqp);
+ break;
+ default:
+ /* The lhs isn't a gimple_reg when var is COMPLEX_TYPE. */
+ omp_gimple_assign_with_ops (OMP_CLAUSE_REDUCTION_CODE (c),
+ t, mem, stmt_seqp);
+ }
+
+ t = fold_build1 (NOP_EXPR, TREE_TYPE (var), t);
+ x = lang_hooks.decls.omp_clause_assign_op (c, var, t);
+ gimplify_and_add (unshare_expr(x), stmt_seqp);
+ }
+
+ /* Increment the induction variable. */
+ tree one = fold_build1 (NOP_EXPR, sizetype, integer_one_node);
+ stmt = gimple_build_assign_with_ops (PLUS_EXPR, ix, ix, one);
+ gimple_seq_add_stmt (stmt_seqp, stmt);
+
+ /* Go back to the top of the loop. */
+ gimple_seq_add_stmt (stmt_seqp, gimple_build_goto (loop_header));
+
+ /* Place the loop exit label here. */
+ gimple_seq_add_stmt (stmt_seqp, gimple_build_label (loop_exit));
+}
+
+/* Scan through all of the gimple stmts searching for an OMP_FOR_EXPR, and
+ scan that for reductions. */
+
+static void
+process_reduction_data (gimple_seq *body, gimple_seq *in_stmt_seqp,
+ gimple_seq *out_stmt_seqp, omp_context *ctx)
+{
+ gcc_assert (is_gimple_omp_oacc_specifically (ctx->stmt));
+
+ gimple_stmt_iterator gsi;
+ gimple_seq inner = NULL;
+ gimple stmt;
+
+ /* A collapse clause may have inserted a new bind block. */
+ gsi = gsi_start (*body);
+ while (!gsi_end_p (gsi))
+ {
+ stmt = gsi_stmt (gsi);
+ if (gimple_code (stmt) == GIMPLE_BIND)
+ {
+ inner = gimple_bind_body (as_a <gbind *> (stmt));
+ body = &inner;
+ gsi = gsi_start (*body);
+ }
+ else if (gimple_code (stmt) == GIMPLE_OMP_FOR)
+ break;
+ else
+ gsi_next (&gsi);
+ }
+
+ for (gsi = gsi_start (*body); !gsi_end_p (gsi); gsi_next (&gsi))
+ {
+ tree clauses, nthreads, t, c, acc_device, acc_device_host, call,
+ enter, exit;
+ bool reduction_found = false;
+
+ stmt = gsi_stmt (gsi);
+
+ switch (gimple_code (stmt))
+ {
+ case GIMPLE_OMP_FOR:
+ clauses = gimple_omp_for_clauses (stmt);
+
+ /* Search for a reduction clause. */
+ for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
+ {
+ reduction_found = true;
+ break;
+ }
+
+ if (!reduction_found)
+ break;
+
+ ctx = maybe_lookup_ctx (stmt);
+ t = NULL_TREE;
+
+ /* Extract the number of threads. */
+ nthreads = create_tmp_var (sizetype, NULL);
+ t = oacc_max_threads (ctx);
+ gimplify_assign (nthreads, t, in_stmt_seqp);
+
+ /* Determine if this is kernel will be executed on the host. */
+ call = builtin_decl_explicit (BUILT_IN_ACC_GET_DEVICE_TYPE);
+ acc_device = create_tmp_var (integer_type_node, ".acc_device_type");
+ stmt = gimple_build_call (call, 0);
+ gimple_call_set_lhs (stmt, acc_device);
+ gimple_seq_add_stmt (in_stmt_seqp, stmt);
+
+ /* Set nthreads = 1 for ACC_DEVICE_TYPE=host. */
+ acc_device_host = create_tmp_var (integer_type_node,
+ ".acc_device_host");
+ gimplify_assign (acc_device_host, build_int_cst (integer_type_node,
+ 2),
+ in_stmt_seqp);
+
+ enter = create_artificial_label (UNKNOWN_LOCATION);
+ exit = create_artificial_label (UNKNOWN_LOCATION);
+
+ stmt = gimple_build_cond (EQ_EXPR, acc_device, acc_device_host,
+ enter, exit);
+ gimple_seq_add_stmt (in_stmt_seqp, stmt);
+ gimple_seq_add_stmt (in_stmt_seqp, gimple_build_label (enter));
+ gimplify_assign (nthreads, fold_build1 (NOP_EXPR, sizetype,
+ integer_one_node),
+ in_stmt_seqp);
+ gimple_seq_add_stmt (in_stmt_seqp, gimple_build_label (exit));
+
+ /* Also, set nthreads = 1 for ACC_DEVICE_TYPE=host_nonshm. */
+ gimplify_assign (acc_device_host, build_int_cst (integer_type_node,
+ 3),
+ in_stmt_seqp);
+
+ enter = create_artificial_label (UNKNOWN_LOCATION);
+ exit = create_artificial_label (UNKNOWN_LOCATION);
+
+ stmt = gimple_build_cond (EQ_EXPR, acc_device, acc_device_host,
+ enter, exit);
+ gimple_seq_add_stmt (in_stmt_seqp, stmt);
+ gimple_seq_add_stmt (in_stmt_seqp, gimple_build_label (enter));
+ gimplify_assign (nthreads, fold_build1 (NOP_EXPR, sizetype,
+ integer_one_node),
+ in_stmt_seqp);
+ gimple_seq_add_stmt (in_stmt_seqp, gimple_build_label (exit));
+
+ initialize_reduction_data (clauses, nthreads, in_stmt_seqp, ctx);
+ finalize_reduction_data (clauses, nthreads, out_stmt_seqp, ctx);
+ break;
+ default:
+ // Scan for other directives which support reduction here.
+ break;
+ }
+ }
+}
+
/* If ctx is a worksharing context inside of a cancellable parallel
region and it isn't nowait, add lhs to its GIMPLE_OMP_RETURN
and conditional branch to parallel's cancel_label to handle
@@ -9527,6 +10573,8 @@ lower_omp_for (gimple_stmt_iterator *gsi_p, omp_context *ctx)
if (gimple_omp_for_combined_into_p (stmt))
{
+ gcc_assert (gimple_omp_for_kind (stmt) != GF_OMP_FOR_KIND_OACC_LOOP);
+
extract_omp_for_data (stmt, &fd, NULL);
fdp = &fd;
@@ -9716,6 +10764,8 @@ task_copyfn_remap_type (struct omp_taskcopy_context *tcctx, tree orig_type)
static void
create_task_copyfn (gomp_task *task_stmt, omp_context *ctx)
{
+ gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt));
+
struct function *child_cfun;
tree child_fn, t, c, src, dst, f, sf, arg, sarg, decl;
tree record_type, srecord_type, bind, list;
@@ -10126,7 +11176,7 @@ lower_omp_taskreg (gimple_stmt_iterator *gsi_p, omp_context *ctx)
}
}
-/* Lower the OpenMP target directive in the current statement
+/* Lower the GIMPLE_OMP_TARGET in the current statement
in GSI_P. CTX holds context information for the directive. */
static void
@@ -10134,25 +11184,77 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
{
tree clauses;
tree child_fn, t, c;
- gomp_target *stmt = as_a <gomp_target *> (gsi_stmt (*gsi_p));
+ gimple stmt = gsi_stmt (*gsi_p);
gbind *tgt_bind = NULL, *bind;
- gimple_seq tgt_body = NULL, olist, ilist, new_body;
+ gimple_seq tgt_body = NULL, olist, ilist, orlist, irlist, new_body;
location_t loc = gimple_location (stmt);
- int kind = gimple_omp_target_kind (stmt);
+ bool offloaded, data_region;
unsigned int map_cnt = 0;
+ tree (*gimple_omp_clauses) (const_gimple);
+ void (*gimple_omp_set_data_arg) (gimple, tree);
+
+ offloaded = is_gimple_omp_offloaded (stmt);
+ data_region = false;
+ switch (gimple_code (stmt))
+ {
+ case GIMPLE_OACC_KERNELS:
+ gimple_omp_clauses = gimple_oacc_kernels_clauses;
+ gimple_omp_set_data_arg = gimple_oacc_kernels_set_data_arg;
+ break;
+ case GIMPLE_OACC_PARALLEL:
+ gimple_omp_clauses = gimple_oacc_parallel_clauses;
+ gimple_omp_set_data_arg = gimple_oacc_parallel_set_data_arg;
+ break;
+ case GIMPLE_OMP_TARGET:
+ switch (gimple_omp_target_kind (stmt))
+ {
+ case GF_OMP_TARGET_KIND_DATA:
+ case GF_OMP_TARGET_KIND_OACC_DATA:
+ data_region = true;
+ break;
+ case GF_OMP_TARGET_KIND_REGION:
+ case GF_OMP_TARGET_KIND_UPDATE:
+ case GF_OMP_TARGET_KIND_OACC_UPDATE:
+ case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ break;
+ default:
+ gcc_unreachable ();
+ }
+
+ gimple_omp_clauses = gimple_omp_target_clauses;
+ gimple_omp_set_data_arg = gimple_omp_target_set_data_arg;
+ break;
+ default:
+ gcc_unreachable ();
+ }
- clauses = gimple_omp_target_clauses (stmt);
- if (kind == GF_OMP_TARGET_KIND_REGION)
+ clauses = gimple_omp_clauses (stmt);
+
+ tgt_bind = NULL;
+ tgt_body = NULL;
+ if (offloaded)
{
tgt_bind = gimple_seq_first_stmt_as_a_bind (gimple_omp_body (stmt));
tgt_body = gimple_bind_body (tgt_bind);
}
- else if (kind == GF_OMP_TARGET_KIND_DATA)
+ else if (data_region)
tgt_body = gimple_omp_body (stmt);
child_fn = ctx->cb.dst_fn;
push_gimplify_context ();
+ irlist = NULL;
+ orlist = NULL;
+ switch (gimple_code (stmt))
+ {
+ case GIMPLE_OACC_KERNELS:
+ case GIMPLE_OACC_PARALLEL:
+ process_reduction_data (&tgt_body, &irlist, &orlist, ctx);
+ break;
+ default:
+ break;
+ }
+
for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
switch (OMP_CLAUSE_CODE (c))
{
@@ -10161,8 +11263,37 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
default:
break;
case OMP_CLAUSE_MAP:
+#ifdef ENABLE_CHECKING
+ /* First check what we're prepared to handle in the following. */
+ switch (OMP_CLAUSE_MAP_KIND (c))
+ {
+ case OMP_CLAUSE_MAP_ALLOC:
+ case OMP_CLAUSE_MAP_TO:
+ case OMP_CLAUSE_MAP_FROM:
+ case OMP_CLAUSE_MAP_TOFROM:
+ case OMP_CLAUSE_MAP_POINTER:
+ case OMP_CLAUSE_MAP_TO_PSET:
+ break;
+ case OMP_CLAUSE_MAP_FORCE_ALLOC:
+ case OMP_CLAUSE_MAP_FORCE_TO:
+ case OMP_CLAUSE_MAP_FORCE_FROM:
+ case OMP_CLAUSE_MAP_FORCE_TOFROM:
+ case OMP_CLAUSE_MAP_FORCE_PRESENT:
+ case OMP_CLAUSE_MAP_FORCE_DEALLOC:
+ case OMP_CLAUSE_MAP_FORCE_DEVICEPTR:
+ gcc_assert (is_gimple_omp_oacc_specifically (stmt));
+ break;
+ default:
+ gcc_unreachable ();
+ }
+#endif
+ /* FALLTHRU */
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+ gcc_assert (gimple_code (stmt) == GIMPLE_OMP_TARGET
+ && (gimple_omp_target_kind (stmt)
+ == GF_OMP_TARGET_KIND_UPDATE));
var = OMP_CLAUSE_DECL (c);
if (!DECL_P (var))
{
@@ -10185,12 +11316,15 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
if (!maybe_lookup_field (var, ctx))
continue;
- if (kind == GF_OMP_TARGET_KIND_REGION)
+ if (offloaded)
{
x = build_receiver_ref (var, true, ctx);
tree new_var = lookup_decl (var, ctx);
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
+ gcc_assert (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP);
+ gcc_assert ((OMP_CLAUSE_MAP_KIND (c)
+ != OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
+ || TREE_CODE (TREE_TYPE (var)) != ARRAY_TYPE);
+ if (OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
&& TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
x = build_simple_mem_ref (x);
@@ -10200,16 +11334,16 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
map_cnt++;
}
- if (kind == GF_OMP_TARGET_KIND_REGION)
+ if (offloaded)
{
target_nesting_level++;
lower_omp (&tgt_body, ctx);
target_nesting_level--;
}
- else if (kind == GF_OMP_TARGET_KIND_DATA)
+ else if (data_region)
lower_omp (&tgt_body, ctx);
- if (kind == GF_OMP_TARGET_KIND_REGION)
+ if (offloaded)
{
/* Declare all the variables created by mapping and the variables
declared in the scope of the target body. */
@@ -10233,14 +11367,25 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
DECL_NAMELESS (TREE_VEC_ELT (t, 1)) = 1;
TREE_ADDRESSABLE (TREE_VEC_ELT (t, 1)) = 1;
TREE_STATIC (TREE_VEC_ELT (t, 1)) = 1;
+ tree tkind_type;
+ int talign_shift;
+ if (is_gimple_omp_oacc_specifically (stmt))
+ {
+ tkind_type = short_unsigned_type_node;
+ talign_shift = 8;
+ }
+ else
+ {
+ tkind_type = unsigned_char_type_node;
+ talign_shift = 3;
+ }
TREE_VEC_ELT (t, 2)
- = create_tmp_var (build_array_type_nelts (unsigned_char_type_node,
- map_cnt),
+ = create_tmp_var (build_array_type_nelts (tkind_type, map_cnt),
".omp_data_kinds");
DECL_NAMELESS (TREE_VEC_ELT (t, 2)) = 1;
TREE_ADDRESSABLE (TREE_VEC_ELT (t, 2)) = 1;
TREE_STATIC (TREE_VEC_ELT (t, 2)) = 1;
- gimple_omp_target_set_data_arg (stmt, t);
+ gimple_omp_set_data_arg (stmt, t);
vec<constructor_elt, va_gc> *vsize;
vec<constructor_elt, va_gc> *vkind;
@@ -10301,12 +11446,22 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
{
tree var = lookup_decl_in_outer_ctx (ovar, ctx);
tree x = build_sender_ref (ovar, ctx);
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
- && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
- && TREE_CODE (TREE_TYPE (ovar)) == ARRAY_TYPE)
+ gcc_assert (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+ || (OMP_CLAUSE_MAP_KIND (c)
+ != OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
+ || TREE_CODE (TREE_TYPE (ovar)) != ARRAY_TYPE);
+ if (maybe_lookup_reduction (var, ctx))
+ {
+ gcc_assert (gimple_code (stmt) == GIMPLE_OACC_KERNELS
+ || gimple_code (stmt) == GIMPLE_OACC_PARALLEL);
+ gimplify_assign (x, var, &ilist);
+ }
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
+ && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
+ && TREE_CODE (TREE_TYPE (ovar)) == ARRAY_TYPE)
{
- gcc_assert (kind == GF_OMP_TARGET_KIND_REGION);
+ gcc_assert (offloaded);
tree avar
= create_tmp_var (TREE_TYPE (TREE_TYPE (x)), NULL);
mark_addressable (avar);
@@ -10317,16 +11472,22 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
}
else if (is_gimple_reg (var))
{
- gcc_assert (kind == GF_OMP_TARGET_KIND_REGION);
+ gcc_assert (offloaded);
tree avar = create_tmp_var (TREE_TYPE (var), NULL);
mark_addressable (avar);
- if (OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_ALLOC
- && OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_FROM)
+ enum omp_clause_map_kind map_kind
+ = OMP_CLAUSE_MAP_KIND (c);
+ if ((!(map_kind & OMP_CLAUSE_MAP_SPECIAL)
+ && (map_kind & OMP_CLAUSE_MAP_TO))
+ || map_kind == OMP_CLAUSE_MAP_POINTER
+ || map_kind == OMP_CLAUSE_MAP_TO_PSET
+ || map_kind == OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
gimplify_assign (avar, var, &ilist);
avar = build_fold_addr_expr (avar);
gimplify_assign (x, avar, &ilist);
- if ((OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_FROM
- || OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_TOFROM)
+ if (((!(map_kind & OMP_CLAUSE_MAP_SPECIAL)
+ && (map_kind & OMP_CLAUSE_MAP_FROM))
+ || map_kind == OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
&& !TYPE_READONLY (TREE_TYPE (var)))
{
x = build_sender_ref (ovar, ctx);
@@ -10349,7 +11510,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
if (TREE_CODE (s) != INTEGER_CST)
TREE_STATIC (TREE_VEC_ELT (t, 1)) = 0;
- unsigned char tkind = 0;
+ unsigned HOST_WIDE_INT tkind;
switch (OMP_CLAUSE_CODE (c))
{
case OMP_CLAUSE_MAP:
@@ -10364,11 +11525,12 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
default:
gcc_unreachable ();
}
+ gcc_assert (tkind < (HOST_WIDE_INT_C (1U) << talign_shift));
talign = ceil_log2 (talign);
- tkind |= talign << 3;
+ tkind |= talign << talign_shift;
+ gcc_assert (tkind <= tree_to_uhwi (TYPE_MAX_VALUE (tkind_type)));
CONSTRUCTOR_APPEND_ELT (vkind, purpose,
- build_int_cst (unsigned_char_type_node,
- tkind));
+ build_int_cstu (tkind_type, tkind));
if (nc && nc != c)
c = nc;
}
@@ -10406,7 +11568,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
new_body = NULL;
- if (ctx->record_type && kind == GF_OMP_TARGET_KIND_REGION)
+ if (offloaded
+ && ctx->record_type)
{
t = build_fold_addr_expr_loc (loc, ctx->sender_decl);
/* fixup_child_record_type might have changed receiver_decl's type. */
@@ -10415,14 +11578,14 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
gimple_build_assign (ctx->receiver_decl, t));
}
- if (kind == GF_OMP_TARGET_KIND_REGION)
+ if (offloaded)
{
gimple_seq_add_seq (&new_body, tgt_body);
new_body = maybe_catch_exception (new_body);
}
- else if (kind == GF_OMP_TARGET_KIND_DATA)
+ else if (data_region)
new_body = tgt_body;
- if (kind != GF_OMP_TARGET_KIND_UPDATE)
+ if (offloaded || data_region)
{
gimple_seq_add_stmt (&new_body, gimple_build_omp_return (false));
gimple_omp_set_body (stmt, new_body);
@@ -10432,9 +11595,11 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
tgt_bind ? gimple_bind_block (tgt_bind)
: NULL_TREE);
gsi_replace (gsi_p, bind, true);
+ gimple_bind_add_seq (bind, irlist);
gimple_bind_add_seq (bind, ilist);
gimple_bind_add_stmt (bind, stmt);
gimple_bind_add_seq (bind, olist);
+ gimple_bind_add_seq (bind, orlist);
pop_gimplify_context (NULL);
}
@@ -10641,9 +11806,13 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context *ctx)
lower_omp_regimplify_p, ctx ? NULL : &wi, NULL))
gimple_regimplify_operands (stmt, gsi_p);
break;
+ case GIMPLE_OACC_KERNELS:
+ case GIMPLE_OACC_PARALLEL:
case GIMPLE_OMP_TARGET:
ctx = maybe_lookup_ctx (stmt);
gcc_assert (ctx);
+ if (is_gimple_omp_oacc_specifically (stmt))
+ gcc_assert (!ctx->cancellable);
lower_omp_target (gsi_p, ctx);
break;
case GIMPLE_OMP_TEAMS:
@@ -10733,8 +11902,8 @@ lower_omp (gimple_seq *body, omp_context *ctx)
gimple_stmt_iterator gsi;
for (gsi = gsi_start (*body); !gsi_end_p (gsi); gsi_next (&gsi))
lower_omp_1 (&gsi, ctx);
- /* During gimplification, we have not always invoked fold_stmt
- (gimplify.c:maybe_fold_stmt); call it now. */
+ /* During gimplification, we haven't folded statments inside offloading
+ regions (gimplify.c:maybe_fold_stmt); do that now. */
if (target_nesting_level)
for (gsi = gsi_start (*body); !gsi_end_p (gsi); gsi_next (&gsi))
fold_stmt (&gsi);
@@ -10751,8 +11920,9 @@ execute_lower_omp (void)
omp_context *ctx;
/* This pass always runs, to provide PROP_gimple_lomp.
- But there is nothing to do unless -fopenmp is given. */
- if (flag_openmp == 0 && flag_openmp_simd == 0 && flag_cilkplus == 0)
+ But often, there is nothing to do. */
+ if (flag_cilkplus == 0 && flag_openacc == 0 && flag_openmp == 0
+ && flag_openmp_simd == 0)
return 0;
all_contexts = splay_tree_new (splay_tree_compare_pointers, 0,
@@ -10818,7 +11988,7 @@ make_pass_lower_omp (gcc::context *ctxt)
return new pass_lower_omp (ctxt);
}
-/* The following is a utility to diagnose OpenMP structured block violations.
+/* The following is a utility to diagnose structured block violations.
It is not part of the "omplower" pass, as that's invoked too late. It
should be invoked by the respective front ends after gimplification. */
@@ -10831,9 +12001,38 @@ static bool
diagnose_sb_0 (gimple_stmt_iterator *gsi_p,
gimple branch_ctx, gimple label_ctx)
{
+ gcc_assert (!branch_ctx || is_gimple_omp (branch_ctx));
+ gcc_assert (!label_ctx || is_gimple_omp (label_ctx));
+
if (label_ctx == branch_ctx)
return false;
+ const char* kind = NULL;
+
+ if (flag_cilkplus)
+ {
+ if ((branch_ctx
+ && gimple_code (branch_ctx) == GIMPLE_OMP_FOR
+ && gimple_omp_for_kind (branch_ctx) == GF_OMP_FOR_KIND_CILKSIMD)
+ || (label_ctx
+ && gimple_code (label_ctx) == GIMPLE_OMP_FOR
+ && gimple_omp_for_kind (label_ctx) == GF_OMP_FOR_KIND_CILKSIMD))
+ kind = "Cilk Plus";
+ }
+ if (flag_openacc)
+ {
+ if ((branch_ctx && is_gimple_omp_oacc_specifically (branch_ctx))
+ || (label_ctx && is_gimple_omp_oacc_specifically (label_ctx)))
+ {
+ gcc_assert (kind == NULL);
+ kind = "OpenACC";
+ }
+ }
+ if (kind == NULL)
+ {
+ gcc_assert (flag_openmp);
+ kind = "OpenMP";
+ }
/*
Previously we kept track of the label's entire context in diagnose_sb_[12]
@@ -10866,45 +12065,25 @@ diagnose_sb_0 (gimple_stmt_iterator *gsi_p,
}
if (exit_p)
- error ("invalid exit from OpenMP structured block");
+ error ("invalid exit from %s structured block", kind);
else
- error ("invalid entry to OpenMP structured block");
+ error ("invalid entry to %s structured block", kind);
#endif
- bool cilkplus_block = false;
- if (flag_cilkplus)
- {
- if ((branch_ctx
- && gimple_code (branch_ctx) == GIMPLE_OMP_FOR
- && gimple_omp_for_kind (branch_ctx) == GF_OMP_FOR_KIND_CILKSIMD)
- || (label_ctx
- && gimple_code (label_ctx) == GIMPLE_OMP_FOR
- && gimple_omp_for_kind (label_ctx) == GF_OMP_FOR_KIND_CILKSIMD))
- cilkplus_block = true;
- }
-
/* If it's obvious we have an invalid entry, be specific about the error. */
if (branch_ctx == NULL)
- {
- if (cilkplus_block)
- error ("invalid entry to Cilk Plus structured block");
- else
- error ("invalid entry to OpenMP structured block");
- }
+ error ("invalid entry to %s structured block", kind);
else
{
/* Otherwise, be vague and lazy, but efficient. */
- if (cilkplus_block)
- error ("invalid branch to/from a Cilk Plus structured block");
- else
- error ("invalid branch to/from an OpenMP structured block");
+ error ("invalid branch to/from %s structured block", kind);
}
gsi_replace (gsi_p, gimple_build_nop (), false);
return true;
}
-/* Pass 1: Create a minimal tree of OpenMP structured blocks, and record
+/* Pass 1: Create a minimal tree of structured blocks, and record
where each label is found. */
static tree
@@ -10917,10 +12096,12 @@ diagnose_sb_1 (gimple_stmt_iterator *gsi_p, bool *handled_ops_p,
*handled_ops_p = true;
- switch (gimple_code (stmt))
+ switch (gimple_code (stmt))
{
WALK_SUBSTMTS;
+ case GIMPLE_OACC_KERNELS:
+ case GIMPLE_OACC_PARALLEL:
case GIMPLE_OMP_PARALLEL:
case GIMPLE_OMP_TASK:
case GIMPLE_OMP_SECTIONS:
@@ -10981,6 +12162,8 @@ diagnose_sb_2 (gimple_stmt_iterator *gsi_p, bool *handled_ops_p,
{
WALK_SUBSTMTS;
+ case GIMPLE_OACC_KERNELS:
+ case GIMPLE_OACC_PARALLEL:
case GIMPLE_OMP_PARALLEL:
case GIMPLE_OMP_TASK:
case GIMPLE_OMP_SECTIONS:
@@ -11065,8 +12248,8 @@ diagnose_sb_2 (gimple_stmt_iterator *gsi_p, bool *handled_ops_p,
return NULL_TREE;
}
-/* Called from tree-cfg.c::make_edges to create cfg edges for all GIMPLE_OMP
- codes. */
+/* Called from tree-cfg.c::make_edges to create cfg edges for all relevant
+ GIMPLE_* codes. */
bool
make_gimple_omp_edges (basic_block bb, struct omp_region **region,
int *region_idx)
@@ -11078,6 +12261,8 @@ make_gimple_omp_edges (basic_block bb, struct omp_region **region,
switch (code)
{
+ case GIMPLE_OACC_KERNELS:
+ case GIMPLE_OACC_PARALLEL:
case GIMPLE_OMP_PARALLEL:
case GIMPLE_OMP_TASK:
case GIMPLE_OMP_FOR:
@@ -11095,7 +12280,10 @@ make_gimple_omp_edges (basic_block bb, struct omp_region **region,
case GIMPLE_OMP_TARGET:
cur_region = new_omp_region (bb, code, cur_region);
fallthru = true;
- if (gimple_omp_target_kind (last) == GF_OMP_TARGET_KIND_UPDATE)
+ if (gimple_omp_target_kind (last) == GF_OMP_TARGET_KIND_UPDATE
+ || (gimple_omp_target_kind (last)
+ == GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA)
+ || gimple_omp_target_kind (last) == GF_OMP_TARGET_KIND_OACC_UPDATE)
cur_region = cur_region->outer;
break;
@@ -11233,7 +12421,10 @@ public:
{}
/* opt_pass methods: */
- virtual bool gate (function *) { return flag_openmp || flag_cilkplus; }
+ virtual bool gate (function *)
+ {
+ return flag_cilkplus || flag_openacc || flag_openmp;
+ }
virtual unsigned int execute (function *)
{
return diagnose_omp_structured_block_errors ();
@@ -21,5 +21,5 @@ void foo()
}
}
-// { dg-message "error: invalid branch to/from an OpenMP structured block" "" { target *-*-* } 7 }
+// { dg-message "error: invalid branch to/from OpenMP structured block" "" { target *-*-* } 7 }
// { dg-message "error: invalid entry to OpenMP structured block" "" { target *-*-* } 9 }
@@ -31,5 +31,5 @@ void foo()
continue;
}
-// { dg-message "error: invalid branch to/from an OpenMP structured block" "" { target *-*-* } 14 }
+// { dg-message "error: invalid branch to/from OpenMP structured block" "" { target *-*-* } 14 }
// { dg-message "error: invalid entry to OpenMP structured block" "" { target *-*-* } 16 }
@@ -58,6 +58,6 @@ void foo()
}
}
-// { dg-message "error: invalid branch to/from an OpenMP structured block" "" { target *-*-* } 21 }
-// { dg-message "error: invalid branch to/from an OpenMP structured block" "" { target *-*-* } 26 }
+// { dg-message "error: invalid branch to/from OpenMP structured block" "" { target *-*-* } 21 }
+// { dg-message "error: invalid branch to/from OpenMP structured block" "" { target *-*-* } 26 }
// { dg-message "error: invalid entry to OpenMP structured block" "" { target *-*-* } 30 }
@@ -14,4 +14,4 @@ void foo()
}
}
-// { dg-message "error: invalid branch to/from an OpenMP structured block" "" { target *-*-* } 7 }
+// { dg-message "error: invalid branch to/from OpenMP structured block" "" { target *-*-* } 7 }
@@ -28,5 +28,5 @@ foo (int x)
}
}
-// { dg-error "invalid branch to/from an OpenMP structured block" "" { target *-*-* } 8 }
+// { dg-error "invalid branch to/from OpenMP structured block" "" { target *-*-* } 8 }
// { dg-error "invalid entry to OpenMP structured block" "" { target *-*-* } 10 }
@@ -28,5 +28,5 @@ foo (int x, int y)
}
}
-// { dg-error "invalid branch to/from an OpenMP structured block" "" { target *-*-* } 8 }
+// { dg-error "invalid branch to/from OpenMP structured block" "" { target *-*-* } 8 }
// { dg-error "invalid entry to OpenMP structured block" "" { target *-*-* } 10 }
@@ -28,5 +28,5 @@ foo (int x)
}
}
-// { dg-error "invalid branch to/from an OpenMP structured block" "" { target *-*-* } 8 }
+// { dg-error "invalid branch to/from OpenMP structured block" "" { target *-*-* } 8 }
// { dg-error "invalid entry to OpenMP structured block" "" { target *-*-* } 10 }
@@ -60,7 +60,7 @@ bar (int x)
}
}
-// { dg-error "invalid branch to/from an OpenMP structured block" "" { target *-*-* } 8 }
+// { dg-error "invalid branch to/from OpenMP structured block" "" { target *-*-* } 8 }
// { dg-error "invalid entry to OpenMP structured block" "" { target *-*-* } 10 }
-// { dg-error "invalid branch to/from an OpenMP structured block" "" { target *-*-* } 37 }
+// { dg-error "invalid branch to/from OpenMP structured block" "" { target *-*-* } 37 }
// { dg-error "invalid entry to OpenMP structured block" "" { target *-*-* } 39 }
@@ -11,7 +11,7 @@ void foo()
{
a[i] = b[i];
if (c == 5)
- return; /* { dg-error "invalid branch to/from a Cilk Plus structured block" } */
+ return; /* { dg-error "invalid branch to/from Cilk Plus structured block" } */
}
}
@@ -31,7 +31,7 @@ void baz()
{
bad1:
#pragma omp parallel
- goto bad1; /* { dg-error "invalid branch to/from an OpenMP structured block" } */
+ goto bad1; /* { dg-error "invalid branch to/from OpenMP structured block" } */
goto bad2; /* { dg-error "invalid entry to OpenMP structured block" } */
#pragma omp parallel
@@ -10,7 +10,7 @@ void foo()
{
a[i] = b[i];
if (c == 5)
- return; /* { dg-error "invalid branch to.from a Cilk" } */
+ return; /* { dg-error "invalid branch to/from Cilk Plus structured block" } */
}
}
@@ -23,5 +23,5 @@ void bar()
a[i] = b[i];
}
if (c == 6)
- goto lab; /* { dg-error "invalid entry to Cilk Plus" } */
+ goto lab; /* { dg-error "invalid entry to Cilk Plus structured block" } */
}
@@ -4,9 +4,9 @@ void foo()
{
bad1:
#pragma omp parallel
- goto bad1; // { dg-error "invalid branch" }
+ goto bad1; // { dg-error "invalid branch to/from OpenMP structured block" }
- goto bad2; // { dg-error "invalid entry" }
+ goto bad2; // { dg-error "invalid entry to OpenMP structured block" }
#pragma omp parallel
{
bad2: ;
@@ -3,28 +3,28 @@
void foo(int i)
{
int j;
- switch (i) // { dg-error "invalid entry" }
+ switch (i) // { dg-error "invalid entry to OpenMP structured block" }
{
#pragma omp parallel
{ case 0:; }
}
- switch (i) // { dg-error "invalid entry" }
+ switch (i) // { dg-error "invalid entry to OpenMP structured block" }
{
#pragma omp for
for (j = 0; j < 10; ++ j)
{ case 1:; }
}
- switch (i) // { dg-error "invalid entry" }
+ switch (i) // { dg-error "invalid entry to OpenMP structured block" }
{
#pragma omp critical
{ case 2:; }
}
- switch (i) // { dg-error "invalid entry" }
+ switch (i) // { dg-error "invalid entry to OpenMP structured block" }
{
#pragma omp master
{ case 3:; }
}
- switch (i) // { dg-error "invalid entry" }
+ switch (i) // { dg-error "invalid entry to OpenMP structured block" }
{
#pragma omp sections
{ case 4:;
@@ -32,7 +32,7 @@ void foo(int i)
{ case 5:; }
}
}
- switch (i) // { dg-error "invalid entry" }
+ switch (i) // { dg-error "invalid entry to OpenMP structured block" }
{
#pragma omp ordered
{ default:; }
@@ -11,9 +11,9 @@ void foo()
bad1:
#pragma omp for
for (i = 0; i < 10; ++i)
- goto bad1; // { dg-error "invalid branch" }
+ goto bad1; // { dg-error "invalid branch to/from OpenMP structured block" }
- goto bad2; // { dg-error "invalid entry" }
+ goto bad2; // { dg-error "invalid entry to OpenMP structured block" }
#pragma omp for
for (i = 0; i < 10; ++i)
{
@@ -9,7 +9,7 @@ void foo()
{
#pragma omp sections
{
- continue; // { dg-error "invalid branch" }
+ continue; // { dg-error "invalid branch to/from OpenMP structured block" }
}
}
@@ -18,16 +18,16 @@ void foo()
#pragma omp section
{ bad1: ; }
#pragma omp section
- goto bad1; // { dg-error "invalid branch" }
+ goto bad1; // { dg-error "invalid branch to/from OpenMP structured block" }
}
#pragma omp sections
{
- goto bad2; // { dg-error "invalid branch" }
+ goto bad2; // { dg-error "invalid branch to/from OpenMP structured block" }
}
bad2:;
- goto bad3; // { dg-error "invalid entry" }
+ goto bad3; // { dg-error "invalid entry to OpenMP structured block" }
#pragma omp sections
{
bad3: ;
@@ -4,6 +4,6 @@ void foo()
{
#pragma omp critical
{
- return; // { dg-error "invalid branch" }
+ return; // { dg-error "invalid branch to/from OpenMP structured block" }
}
}
@@ -4,12 +4,12 @@ void foo()
{
#pragma omp master
{
- goto bad1; // { dg-error "invalid branch" }
+ goto bad1; // { dg-error "invalid branch to/from OpenMP structured block" }
}
#pragma omp master
{
bad1:
- return; // { dg-error "invalid branch" }
+ return; // { dg-error "invalid branch to/from OpenMP structured block" }
}
}
@@ -4,6 +4,6 @@ void foo()
{
#pragma omp ordered
{
- return; // { dg-error "invalid branch" }
+ return; // { dg-error "invalid branch to/from OpenMP structured block" }
}
}
@@ -6,15 +6,15 @@ void foo()
for (i = 0; i < 10; ++i)
{
#pragma omp for
- for (j = ({ continue; 0; }); // { dg-error "invalid branch" }
- j < ({ continue; 10; }); // { dg-error "invalid branch" }
- j += ({ continue; 1; })) // { dg-error "invalid branch" }
+ for (j = ({ continue; 0; }); // { dg-error "invalid branch to/from OpenMP structured block" }
+ j < ({ continue; 10; }); // { dg-error "invalid branch to/from OpenMP structured block" }
+ j += ({ continue; 1; })) // { dg-error "invalid branch to/from OpenMP structured block" }
continue;
#pragma omp for
- for (j = ({ break; 0; }); // { dg-error "invalid branch" }
- j < ({ break; 10; }); // { dg-error "invalid branch" }
- j += ({ break; 1; })) // { dg-error "invalid branch" }
+ for (j = ({ break; 0; }); // { dg-error "invalid branch to/from OpenMP structured block" }
+ j < ({ break; 10; }); // { dg-error "invalid branch to/from OpenMP structured block" }
+ j += ({ break; 1; })) // { dg-error "invalid branch to/from OpenMP structured block" }
break; // { dg-error "break" }
}
}
@@ -7,5 +7,5 @@ int foo()
#pragma omp parallel for
for (i = 0; i < 10; ++i)
- return 0; // { dg-error "invalid branch" }
+ return 0; // { dg-error "invalid branch to/from OpenMP structured block" }
}
@@ -3,7 +3,7 @@
void foo(int i)
{
int j;
- switch (i) // { dg-error "invalid entry" }
+ switch (i) // { dg-error "invalid entry to OpenMP structured block" }
{
#pragma omp parallel
{ case 0:; }
@@ -5,9 +5,9 @@ foo (int x)
{
bad1:
#pragma omp target
- goto bad1; /* { dg-error "invalid branch" } */
+ goto bad1; // { dg-error "invalid branch to/from OpenMP structured block" }
- goto bad2; /* { dg-error "invalid entry" } */
+ goto bad2; // { dg-error "invalid entry to OpenMP structured block" }
#pragma omp target
{
bad2: ;
@@ -21,7 +21,7 @@ foo (int x)
{ ok1: break; }
}
- switch (x) /* { dg-error "invalid entry" } */
+ switch (x) // { dg-error "invalid entry to OpenMP structured block" }
{
#pragma omp target
{ case 0:; }
@@ -5,9 +5,9 @@ foo (int x, int y)
{
bad1:
#pragma omp target data map(tofrom: y)
- goto bad1; /* { dg-error "invalid branch" } */
+ goto bad1; // { dg-error "invalid branch to/from OpenMP structured block" }
- goto bad2; /* { dg-error "invalid entry" } */
+ goto bad2; // { dg-error "invalid entry to OpenMP structured block" }
#pragma omp target data map(tofrom: y)
{
bad2: ;
@@ -21,7 +21,7 @@ foo (int x, int y)
{ ok1: break; }
}
- switch (x) /* { dg-error "invalid entry" } */
+ switch (x) // { dg-error "invalid entry to OpenMP structured block" }
{
#pragma omp target data map(tofrom: y)
{ case 0:; }
@@ -5,9 +5,9 @@ foo (int x)
{
bad1:
#pragma omp taskgroup
- goto bad1; /* { dg-error "invalid branch" } */
+ goto bad1; // { dg-error "invalid branch to/from OpenMP structured block" }
- goto bad2; /* { dg-error "invalid entry" } */
+ goto bad2; // { dg-error "invalid entry to OpenMP structured block" }
#pragma omp taskgroup
{
bad2: ;
@@ -21,7 +21,7 @@ foo (int x)
{ ok1: break; }
}
- switch (x) /* { dg-error "invalid entry" } */
+ switch (x) // { dg-error "invalid entry to OpenMP structured block" }
{
#pragma omp taskgroup
{ case 0:; }
@@ -5,9 +5,9 @@ foo (int x)
{
bad1:
#pragma omp target teams
- goto bad1; /* { dg-error "invalid branch" } */
+ goto bad1; // { dg-error "invalid branch to/from OpenMP structured block" }
- goto bad2; /* { dg-error "invalid entry" } */
+ goto bad2; // { dg-error "invalid entry to OpenMP structured block" }
#pragma omp target teams
{
bad2: ;
@@ -21,7 +21,7 @@ foo (int x)
{ ok1: break; }
}
- switch (x) /* { dg-error "invalid entry" } */
+ switch (x) // { dg-error "invalid entry to OpenMP structured block" }
{
#pragma omp target teams
{ case 0:; }
@@ -34,9 +34,9 @@ bar (int x)
bad1:
#pragma omp target
#pragma omp teams
- goto bad1; /* { dg-error "invalid branch" } */
+ goto bad1; // { dg-error "invalid branch to/from OpenMP structured block" }
- goto bad2; /* { dg-error "invalid entry" } */
+ goto bad2; // { dg-error "invalid entry to OpenMP structured block" }
#pragma omp target
#pragma omp teams
{
@@ -52,7 +52,7 @@ bar (int x)
{ ok1: break; }
}
- switch (x) /* { dg-error "invalid entry" } */
+ switch (x) // { dg-error "invalid entry to OpenMP structured block" }
{
#pragma omp target
#pragma omp teams
@@ -221,19 +221,19 @@ enum omp_clause_code {
(c_parser_omp_variable_list). */
OMP_CLAUSE_ERROR = 0,
- /* OpenMP clause: private (variable_list). */
+ /* OpenACC/OpenMP clause: private (variable_list). */
OMP_CLAUSE_PRIVATE,
/* OpenMP clause: shared (variable_list). */
OMP_CLAUSE_SHARED,
- /* OpenMP clause: firstprivate (variable_list). */
+ /* OpenACC/OpenMP clause: firstprivate (variable_list). */
OMP_CLAUSE_FIRSTPRIVATE,
/* OpenMP clause: lastprivate (variable_list). */
OMP_CLAUSE_LASTPRIVATE,
- /* OpenMP clause: reduction (operator:variable_list).
+ /* OpenACC/OpenMP clause: reduction (operator:variable_list).
OMP_CLAUSE_REDUCTION_CODE: The tree_code of the operator.
Operand 1: OMP_CLAUSE_REDUCTION_INIT: Stmt-list to initialize the var.
Operand 2: OMP_CLAUSE_REDUCTION_MERGE: Stmt-list to merge private var
@@ -266,13 +266,42 @@ enum omp_clause_code {
/* OpenMP clause: to (variable-list). */
OMP_CLAUSE_TO,
- /* OpenMP clause: map ({alloc:,to:,from:,tofrom:,}variable-list). */
+ /* OpenACC clauses: {copy, copyin, copyout, create, delete, deviceptr,
+ device, host (self), present, present_or_copy (pcopy), present_or_copyin
+ (pcopyin), present_or_copyout (pcopyout), present_or_create (pcreate)}
+ (variable-list).
+
+ OpenMP clause: map ({alloc:,to:,from:,tofrom:,}variable-list). */
OMP_CLAUSE_MAP,
+ /* Internal structure to hold OpenACC cache directive's variable-list.
+ #pragma acc cache (variable-list). */
+ OMP_CLAUSE__CACHE_,
+
+ /* OpenACC clause: device_resident (variable_list). */
+ OMP_CLAUSE_DEVICE_RESIDENT,
+
+ /* OpenACC clause: use_device (variable_list). */
+ OMP_CLAUSE_USE_DEVICE,
+
+ /* OpenACC clause: gang [(gang-argument-list)].
+ Where
+ gang-argument-list: [gang-argument-list, ] gang-argument
+ gang-argument: [num:] integer-expression
+ | static: size-expression
+ size-expression: * | integer-expression. */
+ OMP_CLAUSE_GANG,
+
+ /* OpenACC clause: async [(integer-expression)]. */
+ OMP_CLAUSE_ASYNC,
+
+ /* OpenACC clause: wait [(integer-expression-list)]. */
+ OMP_CLAUSE_WAIT,
+
/* Internal clause: temporary for combined loops expansion. */
OMP_CLAUSE__LOOPTEMP_,
- /* OpenMP clause: if (scalar-expression). */
+ /* OpenACC/OpenMP clause: if (scalar-expression). */
OMP_CLAUSE_IF,
/* OpenMP clause: num_threads (integer-expression). */
@@ -284,13 +313,15 @@ enum omp_clause_code {
/* OpenMP clause: nowait. */
OMP_CLAUSE_NOWAIT,
- /* OpenMP clause: ordered. */
+ /* OpenACC clause: seq.
+
+ OpenMP clause: ordered. */
OMP_CLAUSE_ORDERED,
/* OpenMP clause: default. */
OMP_CLAUSE_DEFAULT,
- /* OpenMP clause: collapse (constant-integer-expression). */
+ /* OpenACC/OpenMP clause: collapse (constant-integer-expression). */
OMP_CLAUSE_COLLAPSE,
/* OpenMP clause: untied. */
@@ -346,7 +377,25 @@ enum omp_clause_code {
/* Internally used only clause, holding _Cilk_for # of iterations
on OMP_PARALLEL. */
- OMP_CLAUSE__CILK_FOR_COUNT_
+ OMP_CLAUSE__CILK_FOR_COUNT_,
+
+ /* OpenACC clause: independent. */
+ OMP_CLAUSE_INDEPENDENT,
+
+ /* OpenACC clause: worker [( [num:] integer-expression)]. */
+ OMP_CLAUSE_WORKER,
+
+ /* OpenACC clause: vector [( [length:] integer-expression)]. */
+ OMP_CLAUSE_VECTOR,
+
+ /* OpenACC clause: num_gangs (integer-expression). */
+ OMP_CLAUSE_NUM_GANGS,
+
+ /* OpenACC clause: num_workers (integer-expression). */
+ OMP_CLAUSE_NUM_WORKERS,
+
+ /* OpenACC clause: vector_length (integer-expression). */
+ OMP_CLAUSE_VECTOR_LENGTH
};
#undef DEFTREESTRUCT
@@ -1177,19 +1226,45 @@ enum omp_clause_depend_kind
enum omp_clause_map_kind
{
- OMP_CLAUSE_MAP_ALLOC,
- OMP_CLAUSE_MAP_TO,
- OMP_CLAUSE_MAP_FROM,
- OMP_CLAUSE_MAP_TOFROM,
+ /* If not already present, allocate. */
+ OMP_CLAUSE_MAP_ALLOC = 0,
+ /* ..., and copy to device. */
+ OMP_CLAUSE_MAP_TO = 1 << 0,
+ /* ..., and copy from device. */
+ OMP_CLAUSE_MAP_FROM = 1 << 1,
+ /* ..., and copy to and from device. */
+ OMP_CLAUSE_MAP_TOFROM = OMP_CLAUSE_MAP_TO | OMP_CLAUSE_MAP_FROM,
+ /* Special map kinds. */
+ OMP_CLAUSE_MAP_SPECIAL = 1 << 2,
/* The following kind is an internal only map kind, used for pointer based
array sections. OMP_CLAUSE_SIZE for these is not the pointer size,
which is implicitly POINTER_SIZE_UNITS, but the bias. */
- OMP_CLAUSE_MAP_POINTER,
+ OMP_CLAUSE_MAP_POINTER = OMP_CLAUSE_MAP_SPECIAL,
/* Also internal, behaves like OMP_CLAUS_MAP_TO, but additionally any
OMP_CLAUSE_MAP_POINTER records consecutive after it which have addresses
falling into that range will not be ignored if OMP_CLAUSE_MAP_TO_PSET
wasn't mapped already. */
OMP_CLAUSE_MAP_TO_PSET,
+ /* The following are only valid for OpenACC. */
+ /* Flag to force a specific behavior (or else, a run-time error). */
+ OMP_CLAUSE_MAP_FORCE = 1 << 3,
+ /* Allocate. */
+ OMP_CLAUSE_MAP_FORCE_ALLOC = OMP_CLAUSE_MAP_FORCE | OMP_CLAUSE_MAP_ALLOC,
+ /* ..., and copy to device. */
+ OMP_CLAUSE_MAP_FORCE_TO = OMP_CLAUSE_MAP_FORCE | OMP_CLAUSE_MAP_TO,
+ /* ..., and copy from device. */
+ OMP_CLAUSE_MAP_FORCE_FROM = OMP_CLAUSE_MAP_FORCE | OMP_CLAUSE_MAP_FROM,
+ /* ..., and copy to and from device. */
+ OMP_CLAUSE_MAP_FORCE_TOFROM = OMP_CLAUSE_MAP_FORCE | OMP_CLAUSE_MAP_TOFROM,
+ /* Must already be present. */
+ OMP_CLAUSE_MAP_FORCE_PRESENT = OMP_CLAUSE_MAP_FORCE | OMP_CLAUSE_MAP_SPECIAL,
+ /* Deallocate a mapping, without copying from device. */
+ OMP_CLAUSE_MAP_FORCE_DEALLOC,
+ /* Is a device pointer. OMP_CLAUSE_SIZE for these is unused; is implicitly
+ POINTER_SIZE_UNITS. */
+ OMP_CLAUSE_MAP_FORCE_DEVICEPTR,
+
+ /* End marker. */
OMP_CLAUSE_MAP_LAST
};
@@ -1398,6 +1398,10 @@ remap_gimple_stmt (gimple stmt, copy_body_data *id)
copy = gimple_build_wce (s1);
break;
+ case GIMPLE_OACC_KERNELS:
+ case GIMPLE_OACC_PARALLEL:
+ gcc_unreachable ();
+
case GIMPLE_OMP_PARALLEL:
{
gomp_parallel *omp_par_stmt = as_a <gomp_parallel *> (stmt);
@@ -1423,6 +1427,7 @@ remap_gimple_stmt (gimple stmt, copy_body_data *id)
break;
case GIMPLE_OMP_FOR:
+ gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
s1 = remap_gimple_seq (gimple_omp_body (stmt), id);
s2 = remap_gimple_seq (gimple_omp_for_pre_body (stmt), id);
copy = gimple_build_omp_for (s1, gimple_omp_for_kind (stmt),
@@ -1479,6 +1484,7 @@ remap_gimple_stmt (gimple stmt, copy_body_data *id)
break;
case GIMPLE_OMP_TARGET:
+ gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
s1 = remap_gimple_seq (gimple_omp_body (stmt), id);
copy = gimple_build_omp_target
(s1, gimple_omp_target_kind (stmt),
@@ -4136,6 +4142,8 @@ estimate_num_insns (gimple stmt, eni_weights *weights)
+ estimate_num_insns_seq (gimple_omp_body (stmt), weights)
+ estimate_num_insns_seq (gimple_omp_for_pre_body (stmt), weights));
+ case GIMPLE_OACC_KERNELS:
+ case GIMPLE_OACC_PARALLEL:
case GIMPLE_OMP_PARALLEL:
case GIMPLE_OMP_TASK:
case GIMPLE_OMP_CRITICAL:
@@ -627,6 +627,8 @@ walk_gimple_omp_for (gomp_for *for_stmt,
walk_stmt_fn callback_stmt, walk_tree_fn callback_op,
struct nesting_info *info)
{
+ gcc_assert (!is_gimple_omp_oacc_specifically (for_stmt));
+
struct walk_stmt_info wi;
gimple_seq seq;
tree t;
@@ -1323,6 +1325,10 @@ convert_nonlocal_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
}
break;
+ case GIMPLE_OACC_KERNELS:
+ case GIMPLE_OACC_PARALLEL:
+ gcc_unreachable ();
+
case GIMPLE_OMP_PARALLEL:
case GIMPLE_OMP_TASK:
save_suppress = info->suppress_expansion;
@@ -1353,6 +1359,7 @@ convert_nonlocal_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
break;
case GIMPLE_OMP_FOR:
+ gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
save_suppress = info->suppress_expansion;
convert_nonlocal_omp_clauses (gimple_omp_for_clauses_ptr (stmt), wi);
walk_gimple_omp_for (as_a <gomp_for *> (stmt),
@@ -1380,6 +1387,7 @@ convert_nonlocal_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
break;
case GIMPLE_OMP_TARGET:
+ gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
if (gimple_omp_target_kind (stmt) != GF_OMP_TARGET_KIND_REGION)
{
save_suppress = info->suppress_expansion;
@@ -1893,6 +1901,10 @@ convert_local_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
switch (gimple_code (stmt))
{
+ case GIMPLE_OACC_KERNELS:
+ case GIMPLE_OACC_PARALLEL:
+ gcc_unreachable ();
+
case GIMPLE_OMP_PARALLEL:
case GIMPLE_OMP_TASK:
save_suppress = info->suppress_expansion;
@@ -1922,6 +1934,7 @@ convert_local_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
break;
case GIMPLE_OMP_FOR:
+ gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
save_suppress = info->suppress_expansion;
convert_local_omp_clauses (gimple_omp_for_clauses_ptr (stmt), wi);
walk_gimple_omp_for (as_a <gomp_for *> (stmt),
@@ -1949,6 +1962,7 @@ convert_local_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
break;
case GIMPLE_OMP_TARGET:
+ gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
if (gimple_omp_target_kind (stmt) != GF_OMP_TARGET_KIND_REGION)
{
save_suppress = info->suppress_expansion;
@@ -2281,7 +2295,12 @@ convert_tramp_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
break;
}
+ case GIMPLE_OACC_KERNELS:
+ case GIMPLE_OACC_PARALLEL:
+ gcc_unreachable ();
+
case GIMPLE_OMP_TARGET:
+ gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
if (gimple_omp_target_kind (stmt) != GF_OMP_TARGET_KIND_REGION)
{
*handled_ops_p = false;
@@ -2348,6 +2367,10 @@ convert_gimple_call (gimple_stmt_iterator *gsi, bool *handled_ops_p,
}
break;
+ case GIMPLE_OACC_KERNELS:
+ case GIMPLE_OACC_PARALLEL:
+ gcc_unreachable ();
+
case GIMPLE_OMP_PARALLEL:
case GIMPLE_OMP_TASK:
save_static_chain_added = info->static_chain_added;
@@ -2381,6 +2404,7 @@ convert_gimple_call (gimple_stmt_iterator *gsi, bool *handled_ops_p,
break;
case GIMPLE_OMP_TARGET:
+ gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
if (gimple_omp_target_kind (stmt) != GF_OMP_TARGET_KIND_REGION)
{
walk_body (convert_gimple_call, NULL, info, gimple_omp_body_ptr (stmt));
@@ -2418,6 +2442,7 @@ convert_gimple_call (gimple_stmt_iterator *gsi, bool *handled_ops_p,
break;
case GIMPLE_OMP_FOR:
+ gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
walk_body (convert_gimple_call, NULL, info,
gimple_omp_for_pre_body_ptr (stmt));
/* FALLTHRU */
@@ -2429,6 +2454,7 @@ convert_gimple_call (gimple_stmt_iterator *gsi, bool *handled_ops_p,
case GIMPLE_OMP_TASKGROUP:
case GIMPLE_OMP_ORDERED:
case GIMPLE_OMP_CRITICAL:
+ gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
walk_body (convert_gimple_call, NULL, info, gimple_omp_body_ptr (stmt));
break;
@@ -335,6 +335,12 @@ dump_omp_clause (pretty_printer *buffer, tree clause, int spc, int flags)
case OMP_CLAUSE__LOOPTEMP_:
name = "_looptemp_";
goto print_remap;
+ case OMP_CLAUSE_DEVICE_RESIDENT:
+ name = "device_resident";
+ goto print_remap;
+ case OMP_CLAUSE_USE_DEVICE:
+ name = "use_device";
+ goto print_remap;
print_remap:
pp_string (buffer, name);
pp_left_paren (buffer);
@@ -527,6 +533,27 @@ dump_omp_clause (pretty_printer *buffer, tree clause, int spc, int flags)
case OMP_CLAUSE_MAP_TOFROM:
pp_string (buffer, "tofrom");
break;
+ case OMP_CLAUSE_MAP_FORCE_ALLOC:
+ pp_string (buffer, "force_alloc");
+ break;
+ case OMP_CLAUSE_MAP_FORCE_TO:
+ pp_string (buffer, "force_to");
+ break;
+ case OMP_CLAUSE_MAP_FORCE_FROM:
+ pp_string (buffer, "force_from");
+ break;
+ case OMP_CLAUSE_MAP_FORCE_TOFROM:
+ pp_string (buffer, "force_tofrom");
+ break;
+ case OMP_CLAUSE_MAP_FORCE_PRESENT:
+ pp_string (buffer, "force_present");
+ break;
+ case OMP_CLAUSE_MAP_FORCE_DEALLOC:
+ pp_string (buffer, "force_dealloc");
+ break;
+ case OMP_CLAUSE_MAP_FORCE_DEVICEPTR:
+ pp_string (buffer, "force_deviceptr");
+ break;
default:
gcc_unreachable ();
}
@@ -563,6 +590,12 @@ dump_omp_clause (pretty_printer *buffer, tree clause, int spc, int flags)
spc, flags, false);
goto print_clause_size;
+ case OMP_CLAUSE__CACHE_:
+ pp_string (buffer, "(");
+ dump_generic_node (buffer, OMP_CLAUSE_DECL (clause),
+ spc, flags, false);
+ goto print_clause_size;
+
case OMP_CLAUSE_NUM_TEAMS:
pp_string (buffer, "num_teams(");
dump_generic_node (buffer, OMP_CLAUSE_NUM_TEAMS_EXPR (clause),
@@ -636,6 +669,66 @@ dump_omp_clause (pretty_printer *buffer, tree clause, int spc, int flags)
pp_right_paren (buffer);
break;
+ case OMP_CLAUSE_GANG:
+ pp_string (buffer, "gang(");
+ dump_generic_node (buffer, OMP_CLAUSE_GANG_EXPR (clause),
+ spc, flags, false);
+ pp_character(buffer, ')');
+ break;
+
+ case OMP_CLAUSE_ASYNC:
+ pp_string (buffer, "async");
+ if (OMP_CLAUSE_ASYNC_EXPR (clause))
+ {
+ pp_character(buffer, '(');
+ dump_generic_node (buffer, OMP_CLAUSE_ASYNC_EXPR (clause),
+ spc, flags, false);
+ pp_character(buffer, ')');
+ }
+ break;
+
+ case OMP_CLAUSE_WAIT:
+ pp_string (buffer, "wait(");
+ dump_generic_node (buffer, OMP_CLAUSE_WAIT_EXPR (clause),
+ spc, flags, false);
+ pp_character(buffer, ')');
+ break;
+
+ case OMP_CLAUSE_WORKER:
+ pp_string (buffer, "worker(");
+ dump_generic_node (buffer, OMP_CLAUSE_WORKER_EXPR (clause),
+ spc, flags, false);
+ pp_character(buffer, ')');
+ break;
+
+ case OMP_CLAUSE_VECTOR:
+ pp_string (buffer, "vector(");
+ dump_generic_node (buffer, OMP_CLAUSE_VECTOR_EXPR (clause),
+ spc, flags, false);
+ pp_character(buffer, ')');
+ break;
+
+ case OMP_CLAUSE_NUM_GANGS:
+ pp_string (buffer, "num_gangs(");
+ dump_generic_node (buffer, OMP_CLAUSE_NUM_GANGS_EXPR (clause),
+ spc, flags, false);
+ pp_character (buffer, ')');
+ break;
+
+ case OMP_CLAUSE_NUM_WORKERS:
+ pp_string (buffer, "num_workers(");
+ dump_generic_node (buffer, OMP_CLAUSE_NUM_WORKERS_EXPR (clause),
+ spc, flags, false);
+ pp_character (buffer, ')');
+ break;
+
+ case OMP_CLAUSE_VECTOR_LENGTH:
+ pp_string (buffer, "vector_length(");
+ dump_generic_node (buffer, OMP_CLAUSE_VECTOR_LENGTH_EXPR (clause),
+ spc, flags, false);
+ pp_character (buffer, ')');
+ break;
+
case OMP_CLAUSE_INBRANCH:
pp_string (buffer, "inbranch");
break;
@@ -654,6 +747,9 @@ dump_omp_clause (pretty_printer *buffer, tree clause, int spc, int flags)
case OMP_CLAUSE_TASKGROUP:
pp_string (buffer, "taskgroup");
break;
+ case OMP_CLAUSE_INDEPENDENT:
+ pp_string (buffer, "independent");
+ break;
default:
/* Should never happen. */
@@ -2407,6 +2503,51 @@ dump_generic_node (pretty_printer *buffer, tree node, int spc, int flags,
pp_string (buffer, " > ");
break;
+ case OACC_PARALLEL:
+ pp_string (buffer, "#pragma acc parallel");
+ dump_omp_clauses (buffer, OACC_PARALLEL_CLAUSES (node), spc, flags);
+ goto dump_omp_body;
+
+ case OACC_KERNELS:
+ pp_string (buffer, "#pragma acc kernels");
+ dump_omp_clauses (buffer, OACC_KERNELS_CLAUSES (node), spc, flags);
+ goto dump_omp_body;
+
+ case OACC_DATA:
+ pp_string (buffer, "#pragma acc data");
+ dump_omp_clauses (buffer, OACC_DATA_CLAUSES (node), spc, flags);
+ goto dump_omp_body;
+
+ case OACC_HOST_DATA:
+ pp_string (buffer, "#pragma acc host_data");
+ dump_omp_clauses (buffer, OACC_HOST_DATA_CLAUSES (node), spc, flags);
+ goto dump_omp_body;
+
+ case OACC_DECLARE:
+ pp_string (buffer, "#pragma acc declare");
+ dump_omp_clauses (buffer, OACC_DECLARE_CLAUSES (node), spc, flags);
+ break;
+
+ case OACC_UPDATE:
+ pp_string (buffer, "#pragma acc update");
+ dump_omp_clauses (buffer, OACC_UPDATE_CLAUSES (node), spc, flags);
+ break;
+
+ case OACC_ENTER_DATA:
+ pp_string (buffer, "#pragma acc enter data");
+ dump_omp_clauses (buffer, OACC_ENTER_DATA_CLAUSES (node), spc, flags);
+ break;
+
+ case OACC_EXIT_DATA:
+ pp_string (buffer, "#pragma acc exit data");
+ dump_omp_clauses (buffer, OACC_EXIT_DATA_CLAUSES (node), spc, flags);
+ break;
+
+ case OACC_CACHE:
+ pp_string (buffer, "#pragma acc cache");
+ dump_omp_clauses (buffer, OACC_CACHE_CLAUSES (node), spc, flags);
+ break;
+
case OMP_PARALLEL:
pp_string (buffer, "#pragma omp parallel");
dump_omp_clauses (buffer, OMP_PARALLEL_CLAUSES (node), spc, flags);
@@ -2451,6 +2592,10 @@ dump_generic_node (pretty_printer *buffer, tree node, int spc, int flags,
pp_string (buffer, "#pragma omp distribute");
goto dump_omp_loop;
+ case OACC_LOOP:
+ pp_string (buffer, "#pragma acc loop");
+ goto dump_omp_loop;
+
case OMP_TEAMS:
pp_string (buffer, "#pragma omp teams");
dump_omp_clauses (buffer, OMP_TEAMS_CLAUSES (node), spc, flags);
@@ -270,6 +270,12 @@ unsigned const char omp_clause_num_ops[] =
2, /* OMP_CLAUSE_FROM */
2, /* OMP_CLAUSE_TO */
2, /* OMP_CLAUSE_MAP */
+ 2, /* OMP_CLAUSE__CACHE_ */
+ 1, /* OMP_CLAUSE_DEVICE_RESIDENT */
+ 1, /* OMP_CLAUSE_USE_DEVICE */
+ 1, /* OMP_CLAUSE_GANG */
+ 1, /* OMP_CLAUSE_ASYNC */
+ 1, /* OMP_CLAUSE_WAIT */
1, /* OMP_CLAUSE__LOOPTEMP_ */
1, /* OMP_CLAUSE_IF */
1, /* OMP_CLAUSE_NUM_THREADS */
@@ -296,6 +302,12 @@ unsigned const char omp_clause_num_ops[] =
0, /* OMP_CLAUSE_TASKGROUP */
1, /* OMP_CLAUSE__SIMDUID_ */
1, /* OMP_CLAUSE__CILK_FOR_COUNT_ */
+ 0, /* OMP_CLAUSE_INDEPENDENT */
+ 1, /* OMP_CLAUSE_WORKER */
+ 1, /* OMP_CLAUSE_VECTOR */
+ 1, /* OMP_CLAUSE_NUM_GANGS */
+ 1, /* OMP_CLAUSE_NUM_WORKERS */
+ 1, /* OMP_CLAUSE_VECTOR_LENGTH */
};
const char * const omp_clause_code_name[] =
@@ -315,6 +327,12 @@ const char * const omp_clause_code_name[] =
"from",
"to",
"map",
+ "_cache_",
+ "device_resident",
+ "use_device",
+ "gang",
+ "async",
+ "wait",
"_looptemp_",
"if",
"num_threads",
@@ -340,7 +358,13 @@ const char * const omp_clause_code_name[] =
"sections",
"taskgroup",
"_simduid_",
- "_Cilk_for_count_"
+ "_Cilk_for_count_",
+ "independent",
+ "worker",
+ "vector",
+ "num_gangs",
+ "num_workers",
+ "vector_length"
};
@@ -11119,6 +11143,16 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,
case OMP_CLAUSE:
switch (OMP_CLAUSE_CODE (*tp))
{
+ case OMP_CLAUSE_DEVICE_RESIDENT:
+ case OMP_CLAUSE_USE_DEVICE:
+ case OMP_CLAUSE_GANG:
+ case OMP_CLAUSE_ASYNC:
+ case OMP_CLAUSE_WAIT:
+ case OMP_CLAUSE_WORKER:
+ case OMP_CLAUSE_VECTOR:
+ case OMP_CLAUSE_NUM_GANGS:
+ case OMP_CLAUSE_NUM_WORKERS:
+ case OMP_CLAUSE_VECTOR_LENGTH:
case OMP_CLAUSE_PRIVATE:
case OMP_CLAUSE_SHARED:
case OMP_CLAUSE_FIRSTPRIVATE:
@@ -11142,6 +11176,7 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,
WALK_SUBTREE (OMP_CLAUSE_OPERAND (*tp, 0));
/* FALLTHRU */
+ case OMP_CLAUSE_INDEPENDENT:
case OMP_CLAUSE_NOWAIT:
case OMP_CLAUSE_ORDERED:
case OMP_CLAUSE_DEFAULT:
@@ -11179,6 +11214,7 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,
case OMP_CLAUSE_FROM:
case OMP_CLAUSE_TO:
case OMP_CLAUSE_MAP:
+ case OMP_CLAUSE__CACHE_:
WALK_SUBTREE (OMP_CLAUSE_DECL (*tp));
WALK_SUBTREE (OMP_CLAUSE_OPERAND (*tp, 1));
WALK_SUBTREE_TAIL (OMP_CLAUSE_CHAIN (*tp));
@@ -1030,8 +1030,33 @@ DEFTREECODE (TARGET_MEM_REF, "target_mem_ref", tcc_reference, 5)
chain of component references offsetting p by c. */
DEFTREECODE (MEM_REF, "mem_ref", tcc_reference, 2)
-/* The ordering of the codes between OMP_PARALLEL and OMP_CRITICAL is
- exposed to TREE_RANGE_CHECK. */
+/* OpenACC and OpenMP. As it is exposed in TREE_RANGE_CHECK invocations, do
+ not change the ordering of these codes. */
+
+/* OpenACC - #pragma acc parallel [clause1 ... clauseN]
+ Operand 0: OACC_PARALLEL_BODY: Code to be executed in parallel.
+ Operand 1: OACC_PARALLEL_CLAUSES: List of clauses. */
+
+DEFTREECODE (OACC_PARALLEL, "oacc_parallel", tcc_statement, 2)
+
+/* OpenACC - #pragma acc kernels [clause1 ... clauseN]
+ Operand 0: OACC_KERNELS_BODY: Sequence of kernels.
+ Operand 1: OACC_KERNELS_CLAUSES: List of clauses. */
+
+DEFTREECODE (OACC_KERNELS, "oacc_kernels", tcc_statement, 2)
+
+/* OpenACC - #pragma acc data [clause1 ... clauseN]
+ Operand 0: OACC_DATA_BODY: Data construct body.
+ Operand 1: OACC_DATA_CLAUSES: List of clauses. */
+
+DEFTREECODE (OACC_DATA, "oacc_data", tcc_statement, 2)
+
+/* OpenACC - #pragma acc host_data [clause1 ... clauseN]
+ Operand 0: OACC_HOST_DATA_BODY: Host_data construct body.
+ Operand 1: OACC_HOST_DATA_CLAUSES: List of clauses. */
+
+DEFTREECODE (OACC_HOST_DATA, "oacc_host_data", tcc_statement, 2)
+
/* OpenMP - #pragma omp parallel [clause1 ... clauseN]
Operand 0: OMP_PARALLEL_BODY: Code to be executed by all threads.
Operand 1: OMP_PARALLEL_CLAUSES: List of clauses. */
@@ -1062,7 +1087,7 @@ DEFTREECODE (OMP_TASK, "omp_task", tcc_statement, 2)
private. N1, N2 and INCR are required to be loop invariant integer
expressions that are evaluated without any synchronization.
The evaluation order, frequency of evaluation and side-effects are
- unspecified by the standard. */
+ unspecified by the standards. */
DEFTREECODE (OMP_FOR, "omp_for", tcc_statement, 6)
/* OpenMP - #pragma omp simd [clause1 ... clauseN]
@@ -1081,6 +1106,10 @@ DEFTREECODE (CILK_FOR, "cilk_for", tcc_statement, 6)
Operands like for OMP_FOR. */
DEFTREECODE (OMP_DISTRIBUTE, "omp_distribute", tcc_statement, 6)
+/* OpenMP - #pragma acc loop [clause1 ... clauseN]
+ Operands like for OMP_FOR. */
+DEFTREECODE (OACC_LOOP, "oacc_loop", tcc_statement, 6)
+
/* OpenMP - #pragma omp teams [clause1 ... clauseN]
Operand 0: OMP_TEAMS_BODY: Teams body.
Operand 1: OMP_TEAMS_CLAUSES: List of clauses. */
@@ -1127,6 +1156,27 @@ DEFTREECODE (OMP_ORDERED, "omp_ordered", tcc_statement, 1)
Operand 1: OMP_CRITICAL_NAME: Identifier for critical section. */
DEFTREECODE (OMP_CRITICAL, "omp_critical", tcc_statement, 2)
+/* OpenACC - #pragma acc declare [clause1 ... clauseN]
+ Operand 0: OACC_DECLARE_CLAUSES: List of clauses. */
+DEFTREECODE (OACC_DECLARE, "oacc_declare", tcc_statement, 1)
+
+/* OpenACC - #pragma acc update [clause1 ... clauseN]
+ Operand 0: OACC_UPDATE_CLAUSES: List of clauses. */
+DEFTREECODE (OACC_UPDATE, "oacc_update", tcc_statement, 1)
+
+/* OpenACC - #pragma acc enter data [clause1 ... clauseN]
+ Operand 0: OACC_ENTER_DATA_CLAUSES: List of clauses. */
+DEFTREECODE (OACC_ENTER_DATA, "oacc_enter_data", tcc_statement, 1)
+
+/* OpenACC - #pragma acc exit data [clause1 ... clauseN]
+ Operand 0: OACC_EXIT_DATA_CLAUSES: List of clauses. */
+DEFTREECODE (OACC_EXIT_DATA, "oacc_exit_data", tcc_statement, 1)
+
+/* OpenACC - #pragma acc cache (variable1 ... variableN)
+ Operand 0: OACC_CACHE_CLAUSES: List of variables (transformed into
+ OMP_CLAUSE__CACHE_ clauses). */
+DEFTREECODE (OACC_CACHE, "oacc_cache", tcc_statement, 1)
+
/* OpenMP - #pragma omp target update [clause1 ... clauseN]
Operand 0: OMP_TARGET_UPDATE_CLAUSES: List of clauses. */
DEFTREECODE (OMP_TARGET_UPDATE, "omp_target_update", tcc_statement, 1)
@@ -1184,12 +1184,47 @@ extern void protected_set_expr_location (tree, location_t);
#define TRANSACTION_EXPR_RELAXED(NODE) \
(TRANSACTION_EXPR_CHECK (NODE)->base.public_flag)
-/* OpenMP directive and clause accessors. */
+/* OpenMP and OpenACC directive and clause accessors. */
#define OMP_BODY(NODE) \
- TREE_OPERAND (TREE_RANGE_CHECK (NODE, OMP_PARALLEL, OMP_CRITICAL), 0)
+ TREE_OPERAND (TREE_RANGE_CHECK (NODE, OACC_PARALLEL, OMP_CRITICAL), 0)
#define OMP_CLAUSES(NODE) \
- TREE_OPERAND (TREE_RANGE_CHECK (NODE, OMP_PARALLEL, OMP_SINGLE), 1)
+ TREE_OPERAND (TREE_RANGE_CHECK (NODE, OACC_PARALLEL, OMP_SINGLE), 1)
+
+#define OACC_PARALLEL_BODY(NODE) \
+ TREE_OPERAND (OACC_PARALLEL_CHECK (NODE), 0)
+#define OACC_PARALLEL_CLAUSES(NODE) \
+ TREE_OPERAND (OACC_PARALLEL_CHECK (NODE), 1)
+
+#define OACC_KERNELS_BODY(NODE) \
+ TREE_OPERAND (OACC_KERNELS_CHECK(NODE), 0)
+#define OACC_KERNELS_CLAUSES(NODE) \
+ TREE_OPERAND (OACC_KERNELS_CHECK(NODE), 1)
+
+#define OACC_DATA_BODY(NODE) \
+ TREE_OPERAND (OACC_DATA_CHECK (NODE), 0)
+#define OACC_DATA_CLAUSES(NODE) \
+ TREE_OPERAND (OACC_DATA_CHECK (NODE), 1)
+
+#define OACC_HOST_DATA_BODY(NODE) \
+ TREE_OPERAND (OACC_HOST_DATA_CHECK (NODE), 0)
+#define OACC_HOST_DATA_CLAUSES(NODE) \
+ TREE_OPERAND (OACC_HOST_DATA_CHECK (NODE), 1)
+
+#define OACC_DECLARE_CLAUSES(NODE) \
+ TREE_OPERAND (OACC_DECLARE_CHECK (NODE), 0)
+
+#define OACC_ENTER_DATA_CLAUSES(NODE) \
+ TREE_OPERAND (OACC_ENTER_DATA_CHECK (NODE), 0)
+
+#define OACC_EXIT_DATA_CLAUSES(NODE) \
+ TREE_OPERAND (OACC_EXIT_DATA_CHECK (NODE), 0)
+
+#define OACC_UPDATE_CLAUSES(NODE) \
+ TREE_OPERAND (OACC_UPDATE_CHECK (NODE), 0)
+
+#define OACC_CACHE_CLAUSES(NODE) \
+ TREE_OPERAND (OACC_CACHE_CHECK (NODE), 0)
#define OMP_PARALLEL_BODY(NODE) TREE_OPERAND (OMP_PARALLEL_CHECK (NODE), 0)
#define OMP_PARALLEL_CLAUSES(NODE) TREE_OPERAND (OMP_PARALLEL_CHECK (NODE), 1)
@@ -1201,7 +1236,7 @@ extern void protected_set_expr_location (tree, location_t);
#define OMP_TASKREG_BODY(NODE) TREE_OPERAND (OMP_TASKREG_CHECK (NODE), 0)
#define OMP_TASKREG_CLAUSES(NODE) TREE_OPERAND (OMP_TASKREG_CHECK (NODE), 1)
-#define OMP_LOOP_CHECK(NODE) TREE_RANGE_CHECK (NODE, OMP_FOR, OMP_DISTRIBUTE)
+#define OMP_LOOP_CHECK(NODE) TREE_RANGE_CHECK (NODE, OMP_FOR, OACC_LOOP)
#define OMP_FOR_BODY(NODE) TREE_OPERAND (OMP_LOOP_CHECK (NODE), 0)
#define OMP_FOR_CLAUSES(NODE) TREE_OPERAND (OMP_LOOP_CHECK (NODE), 1)
#define OMP_FOR_INIT(NODE) TREE_OPERAND (OMP_LOOP_CHECK (NODE), 2)
@@ -1243,7 +1278,7 @@ extern void protected_set_expr_location (tree, location_t);
#define OMP_CLAUSE_SIZE(NODE) \
OMP_CLAUSE_OPERAND (OMP_CLAUSE_RANGE_CHECK (OMP_CLAUSE_CHECK (NODE), \
OMP_CLAUSE_FROM, \
- OMP_CLAUSE_MAP), 1)
+ OMP_CLAUSE__CACHE_), 1)
#define OMP_CLAUSE_CHAIN(NODE) TREE_CHAIN (OMP_CLAUSE_CHECK (NODE))
#define OMP_CLAUSE_DECL(NODE) \
@@ -1260,6 +1295,15 @@ extern void protected_set_expr_location (tree, location_t);
#define OMP_SECTION_LAST(NODE) \
(OMP_SECTION_CHECK (NODE)->base.private_flag)
+/* True on an OACC_KERNELS statement if is represents combined kernels loop
+ directive. */
+#define OACC_KERNELS_COMBINED(NODE) \
+ (OACC_KERNELS_CHECK (NODE)->base.private_flag)
+
+/* Like OACC_KERNELS_COMBINED, but for parallel loop directive. */
+#define OACC_PARALLEL_COMBINED(NODE) \
+ (OACC_PARALLEL_CHECK (NODE)->base.private_flag)
+
/* True on an OMP_PARALLEL statement if it represents an explicit
combined parallel work-sharing constructs. */
#define OMP_PARALLEL_COMBINED(NODE) \
@@ -1302,6 +1346,32 @@ extern void protected_set_expr_location (tree, location_t);
#define OMP_CLAUSE_SCHEDULE_CHUNK_EXPR(NODE) \
OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_SCHEDULE), 0)
+/* OpenACC clause expressions */
+#define OMP_CLAUSE_GANG_EXPR(NODE) \
+ OMP_CLAUSE_OPERAND ( \
+ OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_GANG), 0)
+#define OMP_CLAUSE_ASYNC_EXPR(NODE) \
+ OMP_CLAUSE_OPERAND ( \
+ OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_ASYNC), 0)
+#define OMP_CLAUSE_WAIT_EXPR(NODE) \
+ OMP_CLAUSE_OPERAND ( \
+ OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_WAIT), 0)
+#define OMP_CLAUSE_VECTOR_EXPR(NODE) \
+ OMP_CLAUSE_OPERAND ( \
+ OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_VECTOR), 0)
+#define OMP_CLAUSE_WORKER_EXPR(NODE) \
+ OMP_CLAUSE_OPERAND ( \
+ OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_WORKER), 0)
+#define OMP_CLAUSE_NUM_GANGS_EXPR(NODE) \
+ OMP_CLAUSE_OPERAND ( \
+ OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_NUM_GANGS), 0)
+#define OMP_CLAUSE_NUM_WORKERS_EXPR(NODE) \
+ OMP_CLAUSE_OPERAND ( \
+ OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_NUM_WORKERS), 0)
+#define OMP_CLAUSE_VECTOR_LENGTH_EXPR(NODE) \
+ OMP_CLAUSE_OPERAND ( \
+ OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_VECTOR_LENGTH), 0)
+
#define OMP_CLAUSE_DEPEND_KIND(NODE) \
(OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_DEPEND)->omp_clause.subcode.depend_kind)