@@ -1,4 +1,41 @@
2014-12-17 Thomas Schwinge <thomas@codesourcery.com>
+ Bernd Schmidt <bernds@codesourcery.com>
+
+ * tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_AUTO, and
+ OMP_CLAUSE_SEQ.
+ * tree.c (omp_clause_num_ops, omp_clause_code_name, walk_tree_1):
+ Update accordingly.
+ * gimplify.c (gimplify_scan_omp_clauses)
+ (gimplify_adjust_omp_clauses): Handle OMP_CLAUSE_AUTO,
+ OMP_CLAUSE_GANG, OMP_CLAUSE_SEQ, OMP_CLAUSE_VECTOR, and
+ OMP_CLAUSE_WORKER.
+ * tree-pretty-print.c (dump_omp_clause): Likewise.
+ * omp-low.c (scan_sharing_clauses): Likewise.
+ (struct omp_context): Add gwv_below, and gwv_this members.
+ (enclosing_target_ctx, oacc_loop_or_target_p): New functions.
+ (scan_omp_for, scan_omp_target): Use those.
+
+2014-12-17 Thomas Schwinge <thomas@codesourcery.com>
+
+ * tree-nested.c (convert_nonlocal_reference_stmt)
+ (convert_local_reference_stmt, convert_tramp_reference_stmt)
+ (convert_gimple_call) <GIMPLE_OMP_TARGET>: Use
+ is_gimple_omp_offloaded.
+
+ * omp-low.c (lower_reduction_clauses): Simplify OpenACC handling.
+
+ * gimple.h: Rename is_gimple_omp_oacc_specifically to
+ is_gimple_omp_oacc. Update all users.
+
+ * gimplify.c (enum omp_region_type): Restore ORT_TARGET_DATA and
+ ORT_TARGET. Update all users.
+
+ * doc/generic.texi (OpenMP): Move OpenACC stuff into...
+ (OpenACC): ... this new subsection.
+
+ * builtin-types.def: Remove BT_FN_VOID_INT_PTR_INT.
+
+ * oacc-builtins.def: Merge into omp-builtins.def. Update all users.
* builtins.c (expand_builtin_acc_on_device): Use
include/gomp-constants.h.
@@ -873,8 +873,7 @@ 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 \
- oacc-builtins.def omp-builtins.def \
+BUILTINS_DEF = builtins.def sync-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)
@@ -372,8 +372,6 @@ DEF_FUNCTION_TYPE_3 (BT_FN_VOID_PTR_INT_SIZE,
BT_VOID, BT_PTR, BT_INT, BT_SIZE)
DEF_FUNCTION_TYPE_3 (BT_FN_VOID_PTR_INT_INT,
BT_VOID, BT_PTR, BT_INT, BT_INT)
-DEF_FUNCTION_TYPE_3 (BT_FN_VOID_INT_PTR_INT,
- BT_VOID, BT_INT, BT_PTR, BT_INT)
DEF_FUNCTION_TYPE_3 (BT_FN_VOID_CONST_PTR_PTR_SIZE,
BT_VOID, BT_CONST_PTR, BT_PTR, BT_SIZE)
DEF_FUNCTION_TYPE_3 (BT_FN_INT_STRING_CONST_STRING_VALIST_ARG,
@@ -146,19 +146,16 @@ 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
+/* Builtin used by the implementation of OpenACC and OpenMP. 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)
+ 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. */
+ flag_openacc, true, true, ATTRS, false, true)
#undef DEF_GOMP_BUILTIN
#define DEF_GOMP_BUILTIN(ENUM, NAME, TYPE, ATTRS) \
DEF_BUILTIN (ENUM, "__builtin_" NAME, BUILT_IN_NORMAL, TYPE, TYPE, \
@@ -905,10 +902,7 @@ 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. */
+/* Offloading and Multi Processing builtins. */
#include "omp-builtins.def"
/* Cilk keywords builtins. */
@@ -1,4 +1,11 @@
2014-12-17 Thomas Schwinge <thomas@codesourcery.com>
+ Bernd Schmidt <bernds@codesourcery.com>
+
+ * c-pragma.h (enum pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_AUTO,
+ PRAGMA_OMP_CLAUSE_GANG, PRAGMA_OMP_CLAUSE_SEQ,
+ PRAGMA_OMP_CLAUSE_VECTOR, and PRAGMA_OMP_CLAUSE_WORKER.
+
+2014-12-17 Thomas Schwinge <thomas@codesourcery.com>
* c-omp.c (c_finish_oacc_wait): Use include/gomp-constants.h.
@@ -1,4 +1,4 @@
-/* This file contains routines to construct GNU OpenMP constructs,
+/* This file contains routines to construct OpenACC and OpenMP constructs,
called from parsing in the C and C++ front ends.
Copyright (C) 2005-2014 Free Software Foundation, Inc.
@@ -1189,7 +1189,7 @@ static const struct omp_pragma_def oacc_pragmas[] = {
{ "loop", PRAGMA_OACC_LOOP },
{ "parallel", PRAGMA_OACC_PARALLEL },
{ "update", PRAGMA_OACC_UPDATE },
- { "wait", PRAGMA_OACC_WAIT },
+ { "wait", PRAGMA_OACC_WAIT }
};
static const struct omp_pragma_def omp_pragmas[] = {
{ "atomic", PRAGMA_OMP_ATOMIC },
@@ -81,6 +81,7 @@ typedef enum pragma_omp_clause {
PRAGMA_OMP_CLAUSE_ALIGNED,
PRAGMA_OMP_CLAUSE_ASYNC,
+ PRAGMA_OMP_CLAUSE_AUTO,
PRAGMA_OMP_CLAUSE_COLLAPSE,
PRAGMA_OMP_CLAUSE_COPY,
PRAGMA_OMP_CLAUSE_COPYIN,
@@ -97,6 +98,7 @@ typedef enum pragma_omp_clause {
PRAGMA_OMP_CLAUSE_FIRSTPRIVATE,
PRAGMA_OMP_CLAUSE_FOR,
PRAGMA_OMP_CLAUSE_FROM,
+ PRAGMA_OMP_CLAUSE_GANG,
PRAGMA_OMP_CLAUSE_HOST,
PRAGMA_OMP_CLAUSE_IF,
PRAGMA_OMP_CLAUSE_INBRANCH,
@@ -124,6 +126,7 @@ typedef enum pragma_omp_clause {
PRAGMA_OMP_CLAUSE_SCHEDULE,
PRAGMA_OMP_CLAUSE_SECTIONS,
PRAGMA_OMP_CLAUSE_SELF,
+ PRAGMA_OMP_CLAUSE_SEQ,
PRAGMA_OMP_CLAUSE_SHARED,
PRAGMA_OMP_CLAUSE_SIMDLEN,
PRAGMA_OMP_CLAUSE_TASKGROUP,
@@ -131,8 +134,10 @@ typedef enum pragma_omp_clause {
PRAGMA_OMP_CLAUSE_TO,
PRAGMA_OMP_CLAUSE_UNIFORM,
PRAGMA_OMP_CLAUSE_UNTIED,
+ PRAGMA_OMP_CLAUSE_VECTOR,
PRAGMA_OMP_CLAUSE_VECTOR_LENGTH,
PRAGMA_OMP_CLAUSE_WAIT,
+ PRAGMA_OMP_CLAUSE_WORKER,
/* Clauses for Cilk Plus SIMD-enabled function. */
PRAGMA_CILK_CLAUSE_NOMASK,
@@ -1,4 +1,15 @@
2014-12-17 Thomas Schwinge <thomas@codesourcery.com>
+ Bernd Schmidt <bernds@codesourcery.com>
+
+ * c-parser.c (c_parser_omp_clause_name): Handle "auto", "gang",
+ "seq", "vector", and "worker".
+ (c_parser_oacc_all_clauses): Handle
+ PRAGMA_OMP_CLAUSE_FIRSTPRIVATE, and PRAGMA_OMP_CLAUSE_PRIVATE.
+ * c-typeck.c (c_finish_omp_clauses): Handle OMP_CLAUSE_AUTO,
+ OMP_CLAUSE_GANG, OMP_CLAUSE_SEQ, OMP_CLAUSE_VECTOR, and
+ OMP_CLAUSE_WORKER.
+
+2014-12-17 Thomas Schwinge <thomas@codesourcery.com>
* c-parser.c (c_parser_oacc_clause_async): Use
include/gomp-constants.h.
@@ -9830,7 +9830,9 @@ c_parser_omp_clause_name (c_parser *parser)
{
pragma_omp_clause result = PRAGMA_OMP_CLAUSE_NONE;
- if (c_parser_next_token_is_keyword (parser, RID_IF))
+ if (c_parser_next_token_is_keyword (parser, RID_AUTO))
+ result = PRAGMA_OMP_CLAUSE_AUTO;
+ else if (c_parser_next_token_is_keyword (parser, RID_IF))
result = PRAGMA_OMP_CLAUSE_IF;
else if (c_parser_next_token_is_keyword (parser, RID_DEFAULT))
result = PRAGMA_OMP_CLAUSE_DEFAULT;
@@ -9882,6 +9884,10 @@ c_parser_omp_clause_name (c_parser *parser)
else if (!strcmp ("from", p))
result = PRAGMA_OMP_CLAUSE_FROM;
break;
+ case 'g':
+ if (!strcmp ("gang", p))
+ result = PRAGMA_OMP_CLAUSE_GANG;
+ break;
case 'h':
if (!strcmp ("host", p))
result = PRAGMA_OMP_CLAUSE_HOST;
@@ -9957,6 +9963,8 @@ c_parser_omp_clause_name (c_parser *parser)
result = PRAGMA_OMP_CLAUSE_SCHEDULE;
else if (!strcmp ("sections", p))
result = PRAGMA_OMP_CLAUSE_SECTIONS;
+ else if (!strcmp ("seq", p))
+ result = PRAGMA_OMP_CLAUSE_SEQ;
else if (!strcmp ("shared", p))
result = PRAGMA_OMP_CLAUSE_SHARED;
else if (!strcmp ("simdlen", p))
@@ -9979,7 +9987,9 @@ c_parser_omp_clause_name (c_parser *parser)
result = PRAGMA_OMP_CLAUSE_UNTIED;
break;
case 'v':
- if (!strcmp ("vector_length", p))
+ if (!strcmp ("vector", p))
+ result = PRAGMA_OMP_CLAUSE_VECTOR;
+ else if (!strcmp ("vector_length", p))
result = PRAGMA_OMP_CLAUSE_VECTOR_LENGTH;
else if (flag_cilkplus && !strcmp ("vectorlength", p))
result = PRAGMA_CILK_CLAUSE_VECTORLENGTH;
@@ -9987,6 +9997,8 @@ c_parser_omp_clause_name (c_parser *parser)
case 'w':
if (!strcmp ("wait", p))
result = PRAGMA_OMP_CLAUSE_WAIT;
+ else if (!strcmp ("worker", p))
+ result = PRAGMA_OMP_CLAUSE_WORKER;
break;
}
}
@@ -10697,9 +10709,7 @@ c_parser_oacc_clause_async (c_parser *parser, tree list)
return list;
}
else
- {
- t = c_fully_fold (t, false, NULL);
- }
+ t = c_fully_fold (t, false, NULL);
check_no_duplicate_clause (list, OMP_CLAUSE_ASYNC, "async");
@@ -11618,6 +11628,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
clauses = c_parser_oacc_data_clause_deviceptr (parser, clauses);
c_name = "deviceptr";
break;
+ case PRAGMA_OMP_CLAUSE_FIRSTPRIVATE:
+ clauses = c_parser_omp_clause_firstprivate (parser, clauses);
+ c_name = "firstprivate";
+ break;
case PRAGMA_OMP_CLAUSE_HOST:
clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "host";
@@ -11654,6 +11668,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "present_or_create";
break;
+ case PRAGMA_OMP_CLAUSE_PRIVATE:
+ clauses = c_parser_omp_clause_private (parser, clauses);
+ c_name = "private";
+ break;
case PRAGMA_OMP_CLAUSE_REDUCTION:
clauses = c_parser_omp_clause_reduction (parser, clauses);
c_name = "reduction";
@@ -11990,8 +12008,8 @@ c_parser_oacc_data (location_t loc, c_parser *parser)
{
tree stmt, clauses, block;
- clauses = c_parser_oacc_all_clauses (parser, OACC_DATA_CLAUSE_MASK,
- "#pragma acc data");
+ clauses = c_parser_oacc_all_clauses (parser, OACC_DATA_CLAUSE_MASK,
+ "#pragma acc data");
block = c_begin_omp_parallel ();
add_stmt (c_parser_omp_structured_block (parser));
@@ -12313,6 +12313,11 @@ c_finish_omp_clauses (tree clauses)
case OMP_CLAUSE_VECTOR_LENGTH:
case OMP_CLAUSE_ASYNC:
case OMP_CLAUSE_WAIT:
+ case OMP_CLAUSE_AUTO:
+ case OMP_CLAUSE_SEQ:
+ case OMP_CLAUSE_GANG:
+ case OMP_CLAUSE_WORKER:
+ case OMP_CLAUSE_VECTOR:
pc = &OMP_CLAUSE_CHAIN (c);
continue;
@@ -1,3 +1,9 @@
+2014-12-17 James Norris <jnorris@codesourcery.com>
+
+ * parser.c (cp_parser_oacc_clause_vector_length)
+ (cp_parser_oacc_wait_list, cp_parser_omp_clause_num_gangs)
+ (cp_parser_omp_clause_num_workers): Revise parsing.
+
2014-12-17 Thomas Schwinge <thomas@codesourcery.com>
* parser.c (cp_parser_oacc_clause_async): Use
@@ -27930,17 +27930,12 @@ cp_parser_oacc_clause_vector_length (cp_parser *parser, tree list)
tree t, c;
location_t location = cp_lexer_peek_token (parser->lexer)->location;
bool error = false;
- HOST_WIDE_INT n;
if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
return list;
t = cp_parser_condition (parser);
- if (t == error_mark_node
- || !INTEGRAL_TYPE_P (TREE_TYPE (t))
- || !tree_fits_shwi_p (t)
- || (n = tree_to_shwi (t)) <= 0
- || (int) n != n)
+ if (t == error_mark_node || !INTEGRAL_TYPE_P (TREE_TYPE (t)))
{
error_at (location, "expected positive integer expression");
error = true;
@@ -27954,8 +27949,8 @@ cp_parser_oacc_clause_vector_length (cp_parser *parser, tree list)
return list;
}
- check_no_duplicate_clause (list, OMP_CLAUSE_VECTOR_LENGTH,
- "vector_length", location);
+ check_no_duplicate_clause (list, OMP_CLAUSE_VECTOR_LENGTH, "vector_length",
+ location);
c = build_omp_clause (location, OMP_CLAUSE_VECTOR_LENGTH);
OMP_CLAUSE_VECTOR_LENGTH_EXPR (c) = t;
@@ -27975,11 +27970,11 @@ cp_parser_oacc_wait_list (cp_parser *parser, location_t clause_loc, tree list)
tree t, args_tree;
args = cp_parser_parenthesized_expression_list (parser, non_attr,
- /*cast_p=*/false,
- /*allow_expansion_p=*/true,
- /*non_constant_p=*/NULL);
+ /*cast_p=*/false,
+ /*allow_expansion_p=*/true,
+ /*non_constant_p=*/NULL);
- if (args == NULL || args->length() == 0)
+ if (args == NULL || args->length () == 0)
{
cp_parser_error (parser, "expected integer expression before ')'");
if (args != NULL)
@@ -27996,22 +27991,18 @@ cp_parser_oacc_wait_list (cp_parser *parser, location_t clause_loc, tree list)
tree targ = TREE_VALUE (t);
if (targ != error_mark_node)
- {
+ {
if (!INTEGRAL_TYPE_P (TREE_TYPE (targ)))
+ error ("%<wait%> expression must be integral");
+ else
{
- error("%<wait%> expression must be integral");
- targ = error_mark_node;
+ tree c = build_omp_clause (clause_loc, OMP_CLAUSE_WAIT);
+
+ mark_rvalue_use (targ);
+ OMP_CLAUSE_DECL (c) = targ;
+ OMP_CLAUSE_CHAIN (c) = list;
+ list = c;
}
- else
- {
- tree c = build_omp_clause (clause_loc, OMP_CLAUSE_WAIT);
-
- mark_rvalue_use (targ);
-
- OMP_CLAUSE_DECL (c) = targ;
- OMP_CLAUSE_CHAIN (c) = list;
- list = c;
- }
}
}
@@ -28230,7 +28221,6 @@ cp_parser_omp_clause_num_gangs (cp_parser *parser, tree list)
{
tree t, c;
location_t location = cp_lexer_peek_token (parser->lexer)->location;
- HOST_WIDE_INT n;
if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
return list;
@@ -28243,10 +28233,7 @@ cp_parser_omp_clause_num_gangs (cp_parser *parser, tree list)
/*or_comma=*/false,
/*consume_paren=*/true);
- if (!INTEGRAL_TYPE_P (TREE_TYPE (t))
- || !tree_fits_shwi_p (t)
- || (n = tree_to_shwi (t)) <= 0
- || (int) n != n)
+ if (!INTEGRAL_TYPE_P (TREE_TYPE (t)))
{
error_at (location, "expected positive integer expression");
return list;
@@ -28300,7 +28287,6 @@ cp_parser_omp_clause_num_workers (cp_parser *parser, tree list)
{
tree t, c;
location_t location = cp_lexer_peek_token (parser->lexer)->location;
- HOST_WIDE_INT n;
if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
return list;
@@ -28313,10 +28299,7 @@ cp_parser_omp_clause_num_workers (cp_parser *parser, tree list)
/*or_comma=*/false,
/*consume_paren=*/true);
- if (!INTEGRAL_TYPE_P (TREE_TYPE (t))
- || !tree_fits_shwi_p (t)
- || (n = tree_to_shwi (t)) <= 0
- || (int) n != n)
+ if (!INTEGRAL_TYPE_P (TREE_TYPE (t)))
{
error_at (location, "expected positive integer expression");
return list;
@@ -29044,7 +29027,7 @@ cp_parser_oacc_clause_async (cp_parser *parser, tree list)
t = cp_parser_expression (parser);
if (t == error_mark_node
- || !cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN))
+ || !cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN))
cp_parser_skip_to_closing_parenthesis (parser, /*recovering=*/true,
/*or_comma=*/false,
/*consume_paren=*/true);
@@ -1819,6 +1819,7 @@ There are also several varieties of complex statements.
* Jumps::
* Cleanups::
* OpenMP::
+* OpenACC::
@end menu
@node Basic Statements
@@ -2049,16 +2050,6 @@ edge. Rethrowing the exception is represented using @code{RESX_EXPR}.
@node OpenMP
@subsection OpenMP
-@tindex OACC_PARALLEL
-@tindex OACC_KERNELS
-@tindex OACC_DATA
-@tindex OACC_HOST_DATA
-@tindex OACC_LOOP
-@tindex OACC_DECLARE
-@tindex OACC_UPDATE
-@tindex OACC_ENTER_DATA
-@tindex OACC_EXIT_DATA
-@tindex OACC_CACHE
@tindex OMP_PARALLEL
@tindex OMP_FOR
@tindex OMP_SECTIONS
@@ -2076,48 +2067,6 @@ All the statements starting with @code{OMP_} represent directives and
clauses used by the OpenMP API @w{@uref{http://www.openmp.org/}}.
@table @code
-@item OACC_PARALLEL
-
-Represents @code{#pragma acc parallel [clause1 @dots{} clauseN]}.
-
-@item OACC_KERNELS
-
-Represents @code{#pragma acc kernels [clause1 @dots{} clauseN]}.
-
-@item OACC_DATA
-
-Represents @code{#pragma acc data [clause1 @dots{} clauseN]}.
-
-@item OACC_HOST_DATA
-
-Represents @code{#pragma acc host_data [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_DECLARE
-
-Represents @code{#pragma acc declare [clause1 @dots{} clauseN]}.
-
-@item OACC_UPDATE
-
-Represents @code{#pragma acc update [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_CACHE
-
-Represents @code{#pragma acc cache (var @dots{})}.
-
@item OMP_PARALLEL
Represents @code{#pragma omp parallel [clause1 @dots{} clauseN]}. It
@@ -2236,7 +2185,7 @@ building code (@code{omp-low.c}).
@item OMP_CONTINUE
Similarly, this instruction does not represent an OpenMP
-directive, it is used by @code{OACC_LOOP}, @code{OMP_FOR} as well as
+directive, it is used by @code{OMP_FOR} (and similar codes) as well as
@code{OMP_SECTIONS} to mark the place where the code needs to
loop to the next iteration, or the next section, respectively.
@@ -2284,6 +2233,67 @@ compilation.
@end table
+@node OpenACC
+@subsection OpenACC
+@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
+
+All the statements starting with @code{OACC_} represent directives and
+clauses used by the OpenACC API @w{@uref{http://www.openacc.org/}}.
+
+@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]}.
+
+@end table
+
@c ---------------------------------------------------------------------
@c Functions
@c ---------------------------------------------------------------------
@@ -1,5 +1,7 @@
2014-12-17 Thomas Schwinge <thomas@codesourcery.com>
+ * types.def: Remove BT_FN_VOID_INT_PTR_INT.
+
* openmp.c (gfc_match_omp_clauses): Use include/gomp-constants.h.
2014-11-20 Cesar Philippidis <cesar@codesourcery.com>
@@ -1195,20 +1195,29 @@ gfc_init_builtin_functions (void)
#undef DEF_GOACC_BUILTIN_COMPILER
#define DEF_GOACC_BUILTIN_COMPILER(code, name, type, attr) \
gfc_define_builtin (name, builtin_types[type], code, name, attr);
-#include "../oacc-builtins.def"
-#undef DEF_GOACC_BUILTIN_COMPILER
+#undef DEF_GOMP_BUILTIN
+#define DEF_GOMP_BUILTIN(code, name, type, attr) /* ignore */
+#include "../omp-builtins.def"
#undef DEF_GOACC_BUILTIN
+#undef DEF_GOACC_BUILTIN_COMPILER
+#undef DEF_GOMP_BUILTIN
}
if (gfc_option.gfc_flag_openmp
|| gfc_option.gfc_flag_openmp_simd
|| flag_tree_parallelize_loops)
{
+#undef DEF_GOACC_BUILTIN
+#define DEF_GOACC_BUILTIN(code, name, type, attr) /* ignore */
+#undef DEF_GOACC_BUILTIN_COMPILER
+#define DEF_GOACC_BUILTIN_COMPILER(code, name, type, attr) /* ignore */
#undef DEF_GOMP_BUILTIN
#define DEF_GOMP_BUILTIN(code, name, type, attr) \
gfc_define_builtin ("__builtin_" name, builtin_types[type], \
code, name, attr);
#include "../omp-builtins.def"
+#undef DEF_GOACC_BUILTIN
+#undef DEF_GOACC_BUILTIN_COMPILER
#undef DEF_GOMP_BUILTIN
}
@@ -145,7 +145,6 @@ DEF_FUNCTION_TYPE_3 (BT_FN_VOID_VPTR_I2_INT, BT_VOID, BT_VOLATILE_PTR, BT_I2, BT
DEF_FUNCTION_TYPE_3 (BT_FN_VOID_VPTR_I4_INT, BT_VOID, BT_VOLATILE_PTR, BT_I4, BT_INT)
DEF_FUNCTION_TYPE_3 (BT_FN_VOID_VPTR_I8_INT, BT_VOID, BT_VOLATILE_PTR, BT_I8, BT_INT)
DEF_FUNCTION_TYPE_3 (BT_FN_VOID_VPTR_I16_INT, BT_VOID, BT_VOLATILE_PTR, BT_I16, BT_INT)
-DEF_FUNCTION_TYPE_3 (BT_FN_VOID_INT_PTR_INT, BT_VOID, BT_INT, BT_PTR, BT_INT)
DEF_FUNCTION_TYPE_4 (BT_FN_VOID_OMPFN_PTR_UINT_UINT,
BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR, BT_UINT, BT_UINT)
@@ -866,7 +866,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 loop construct's clauses.
+ CLAUSES, are any of the construct's clauses.
COLLAPSE is the collapse count.
PRE_BODY is the sequence of statements that are loop invariant. */
@@ -1750,7 +1750,6 @@ gimple_copy (gimple stmt)
break;
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));
@@ -1824,7 +1823,6 @@ 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;
@@ -206,7 +206,8 @@ 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. */
+ exposed by the range check in gimple_omp_subcode(). */
+
/* 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
@@ -356,10 +357,10 @@ 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 {kernels,parallel,data}
+ #pragma acc {kernels,parallel,data,enter data,exit data,update}
#pragma omp target {,data,update}
BODY is the sequence of statements inside the construct
- (NULL for target update).
+ (NULL for some variants).
CLAUSES is an OMP_CLAUSE chain holding the associated clauses.
CHILD_FN is set when outlining the body of the offloaded region.
All the statements in BODY are moved into this newly created
@@ -297,7 +297,7 @@ struct GTY((tag("GSS_CALL")))
};
-/* OpenMP statements (#pragma omp). */
+/* OMP statements. */
struct GTY((tag("GSS_OMP")))
gimple_statement_omp : public gimple_statement_base
@@ -5556,7 +5556,7 @@ gimple_return_set_retbnd (gimple gs, tree retval)
}
-/* Returns true when the gimple statement STMT is any of the OpenMP types. */
+/* Returns true when the gimple statement STMT is any of the OMP types. */
#define CASE_GIMPLE_OMP \
case GIMPLE_OMP_PARALLEL: \
@@ -5589,14 +5589,11 @@ is_gimple_omp (const_gimple stmt)
}
}
-/* Return true if STMT is any of the OpenACC types specifically.
-
- TODO: This function should go away eventually, once all its callers have
- either been fixed, changed into more specific checks, or verified to not
- need any special handling for OpenACC. */
+/* Return true if the OMP gimple statement STMT is any of the OpenACC types
+ specifically. */
static inline bool
-is_gimple_omp_oacc_specifically (const_gimple stmt)
+is_gimple_omp_oacc (const_gimple stmt)
{
gcc_assert (is_gimple_omp (stmt));
switch (gimple_code (stmt))
@@ -5627,7 +5624,7 @@ is_gimple_omp_oacc_specifically (const_gimple stmt)
}
-/* Return true if OMP_* STMT is offloaded. */
+/* Return true if the OMP gimple statement STMT is offloaded. */
static inline bool
is_gimple_omp_offloaded (const_gimple stmt)
@@ -91,8 +91,7 @@ enum gimplify_omp_var_data
GOVD_LINEAR = 2048,
GOVD_ALIGNED = 4096,
- /* Flags for GOVD_MAP. */
- /* Don't copy back. */
+ /* Flag for GOVD_MAP: don't copy back. */
GOVD_MAP_TO_ONLY = 8192,
GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
@@ -110,11 +109,10 @@ enum omp_region_type
ORT_TASK = 4,
ORT_UNTIED_TASK = 5,
ORT_TEAMS = 8,
- ORT_TARGET = 16,
-
- /* Flags for ORT_TARGET. */
- /* Prepare this region for offloading. */
- ORT_TARGET_OFFLOAD = 32
+ /* Data region. */
+ ORT_TARGET_DATA = 16,
+ /* Data region with offloading. */
+ ORT_TARGET = 32
};
/* Gimplify hashtable helper. */
@@ -1783,7 +1781,7 @@ gimplify_var_or_parm_decl (tree *expr_p)
return GS_ERROR;
}
- /* When within an OpenMP context, notice uses of variables. */
+ /* When within an OMP context, notice uses of variables. */
if (gimplify_omp_ctxp && omp_notice_variable (gimplify_omp_ctxp, decl, true))
return GS_ALL_DONE;
@@ -2252,7 +2250,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 inside offloading regsion: it can break code by adding decl
+/* Don't fold inside offloading regions: it can break code by adding decl
references that weren't in the source. We'll do it during omplower pass
instead. */
@@ -2261,8 +2259,7 @@ 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
- && ctx->region_type & ORT_TARGET_OFFLOAD)
+ if (ctx->region_type == ORT_TARGET)
return false;
return fold_stmt (gsi);
}
@@ -5517,13 +5514,11 @@ omp_firstprivatize_variable (struct gimplify_omp_ctx *ctx, tree decl)
else
return;
}
- else if (ctx->region_type & ORT_TARGET)
- {
- if (ctx->region_type & ORT_TARGET_OFFLOAD)
- omp_add_variable (ctx, decl, GOVD_MAP | GOVD_MAP_TO_ONLY);
- }
+ else if (ctx->region_type == ORT_TARGET)
+ 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_SIMD
+ && ctx->region_type != ORT_TARGET_DATA)
omp_add_variable (ctx, decl, GOVD_FIRSTPRIVATE);
ctx = ctx->outer_context;
@@ -5587,7 +5582,7 @@ omp_firstprivatize_type_sizes (struct gimplify_omp_ctx *ctx, tree type)
lang_hooks.types.omp_firstprivatize_type_sizes (ctx, type);
}
-/* Add an entry for DECL in the OpenMP context CTX with FLAGS. */
+/* Add an entry for DECL in the OMP context CTX with FLAGS. */
static void
omp_add_variable (struct gimplify_omp_ctx *ctx, tree decl, unsigned int flags)
@@ -5691,7 +5686,7 @@ omp_add_variable (struct gimplify_omp_ctx *ctx, tree decl, unsigned int flags)
splay_tree_insert (ctx->variables, (splay_tree_key)decl, flags);
}
-/* Notice a threadprivate variable DECL used in OpenMP context CTX.
+/* Notice a threadprivate variable DECL used in OMP context CTX.
This just prints out diagnostics about threadprivate variable uses
in untied tasks. If DECL2 is non-NULL, prevent this warning
on that variable. */
@@ -5704,8 +5699,7 @@ 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)
- && (octx->region_type & ORT_TARGET_OFFLOAD))
+ if (octx->region_type == ORT_TARGET)
{
n = splay_tree_lookup (octx->variables, (splay_tree_key)decl);
if (n == NULL)
@@ -5734,7 +5728,7 @@ omp_notice_threadprivate_variable (struct gimplify_omp_ctx *ctx, tree decl,
return false;
}
-/* Record the fact that DECL was used within the OpenMP context CTX.
+/* Record the fact that DECL was used within the OMP context CTX.
IN_CODE is true when real code uses DECL, and false when we should
merely emit default(none) errors. Return true if DECL is going to
be remapped and thus DECL shouldn't be gimplified into its
@@ -5766,8 +5760,7 @@ 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)
- && (ctx->region_type & ORT_TARGET_OFFLOAD))
+ if (ctx->region_type == ORT_TARGET)
{
ret = lang_hooks.decls.omp_disregard_value_expr (decl, true);
if (n == NULL)
@@ -5798,8 +5791,7 @@ 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)
- && !(ctx->region_type & ORT_TARGET_OFFLOAD)))
+ || ctx->region_type == ORT_TARGET_DATA)
goto do_outer;
/* ??? Some compiler-generated variables (like SAVE_EXPRs) could be
@@ -5852,7 +5844,7 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
{
splay_tree_node n2;
- if (octx->region_type & ORT_TARGET)
+ if ((octx->region_type & (ORT_TARGET_DATA | ORT_TARGET)) != 0)
continue;
n2 = splay_tree_lookup (octx->variables, (splay_tree_key) decl);
if (n2 && (n2->value & GOVD_DATA_SHARE_CLASS) != GOVD_SHARED)
@@ -6005,7 +5997,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)
+ if ((ctx->region_type & (ORT_TARGET | ORT_TARGET_DATA)) != 0)
continue;
n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
@@ -6017,7 +6009,7 @@ 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
+/* Scan the OMP clauses in *LIST_P, installing mappings into a new
and previous omp contexts. */
static void
@@ -6294,22 +6286,26 @@ 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_ASYNC:
+ case OMP_CLAUSE_WAIT:
case OMP_CLAUSE_NUM_GANGS:
case OMP_CLAUSE_NUM_WORKERS:
case OMP_CLAUSE_VECTOR_LENGTH:
+ case OMP_CLAUSE_GANG:
+ case OMP_CLAUSE_WORKER:
+ case OMP_CLAUSE_VECTOR:
if (gimplify_expr (&OMP_CLAUSE_OPERAND (c, 0), pre_p, NULL,
is_gimple_val, fb_rvalue) == GS_ERROR)
remove = true;
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_GANG
+ && gimplify_expr (&OMP_CLAUSE_OPERAND (c, 1), 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;
@@ -6317,6 +6313,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
case OMP_CLAUSE_ORDERED:
case OMP_CLAUSE_UNTIED:
case OMP_CLAUSE_COLLAPSE:
+ case OMP_CLAUSE_AUTO:
+ case OMP_CLAUSE_SEQ:
case OMP_CLAUSE_MERGEABLE:
case OMP_CLAUSE_PROC_BIND:
case OMP_CLAUSE_SAFELEN:
@@ -6437,12 +6435,9 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
OMP_CLAUSE_PRIVATE_OUTER_REF (clause) = 1;
else if (code == OMP_CLAUSE_MAP)
{
- 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;
-
+ OMP_CLAUSE_MAP_KIND (clause) = flags & GOVD_MAP_TO_ONLY
+ ? OMP_CLAUSE_MAP_TO
+ : OMP_CLAUSE_MAP_TOFROM;
if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
{
@@ -6609,9 +6604,7 @@ 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)
- && (ctx->region_type & ORT_TARGET_OFFLOAD)
- && !(n->value & GOVD_SEEN))
+ if (ctx->region_type == ORT_TARGET && !(n->value & GOVD_SEEN))
remove = true;
else if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST
@@ -6697,19 +6690,21 @@ 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_ASYNC:
+ case OMP_CLAUSE_WAIT:
+ case OMP_CLAUSE_DEVICE_RESIDENT:
+ case OMP_CLAUSE_USE_DEVICE:
+ case OMP_CLAUSE_INDEPENDENT:
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:
+ case OMP_CLAUSE_AUTO:
+ case OMP_CLAUSE_SEQ:
+ break;
+
default:
gcc_unreachable ();
}
@@ -6855,7 +6850,6 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
gimple_seq for_body, for_pre_body;
int i;
bool simd;
- enum omp_region_type ort;
bitmap has_decl_expr = NULL;
orig_for_stmt = for_stmt = *expr_p;
@@ -6865,23 +6859,19 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
case OMP_FOR:
case CILK_FOR:
case OMP_DISTRIBUTE:
+ case OACC_LOOP:
simd = false;
- ort = ORT_WORKSHARE;
break;
case OMP_SIMD:
case CILK_SIMD:
simd = true;
- ort = ORT_SIMD;
- break;
- case OACC_LOOP:
- simd = false;
- ort = /* TODO */ ORT_WORKSHARE;
break;
default:
gcc_unreachable ();
}
- gimplify_scan_omp_clauses (&OMP_FOR_CLAUSES (for_stmt), pre_p, ort);
+ gimplify_scan_omp_clauses (&OMP_FOR_CLAUSES (for_stmt), pre_p,
+ simd ? ORT_SIMD : ORT_WORKSHARE);
if (TREE_CODE (for_stmt) == OMP_DISTRIBUTE)
gimplify_omp_ctxp->distribute = true;
@@ -7258,7 +7248,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
return GS_ALL_DONE;
}
-/* Gimplify the gross structure of several OpenACC or OpenMP constructs. */
+/* Gimplify the gross structure of several OMP constructs. */
static void
gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
@@ -7277,12 +7267,12 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
case OACC_KERNELS:
case OACC_PARALLEL:
case OMP_TARGET:
- ort = (enum omp_region_type) (ORT_TARGET | ORT_TARGET_OFFLOAD);
- break;
- case OACC_DATA:
- case OMP_TARGET_DATA:
ort = ORT_TARGET;
break;
+ case OACC_DATA:
+ case OMP_TARGET_DATA:
+ ort = ORT_TARGET_DATA;
+ break;
case OMP_TEAMS:
ort = ORT_TEAMS;
break;
@@ -7290,7 +7280,7 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
gcc_unreachable ();
}
gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ort);
- if (ort & ORT_TARGET)
+ if (ort == ORT_TARGET || ort == ORT_TARGET_DATA)
{
push_gimplify_context ();
gimple g = gimplify_and_return_first (OMP_BODY (expr), &body);
@@ -7298,7 +7288,7 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
pop_gimplify_context (g);
else
pop_gimplify_context (NULL);
- if (!(ort & ORT_TARGET_OFFLOAD))
+ if (ort == ORT_TARGET_DATA)
{
enum built_in_function end_ix;
switch (TREE_CODE (expr))
@@ -7329,7 +7319,7 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
{
case OACC_DATA:
stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_DATA,
- OACC_DATA_CLAUSES (expr));
+ OMP_CLAUSES (expr));
break;
case OACC_KERNELS:
stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_KERNELS,
@@ -7364,8 +7354,8 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
*expr_p = NULL_TREE;
}
-/* Gimplify the gross structure of OpenACC update and OpenMP target update
- constructs. */
+/* Gimplify the gross structure of OpenACC enter/exit data, update, and OpenMP
+ target update constructs. */
static void
gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
@@ -7577,7 +7567,7 @@ gimplify_transaction (tree *expr_p, gimple_seq *pre_p)
int subcode = 0;
/* Wrap the transaction body in a BIND_EXPR so we have a context
- where to put decls for OpenMP. */
+ where to put decls for OMP. */
if (TREE_CODE (tbody) != BIND_EXPR)
{
tree bind = build3 (BIND_EXPR, void_type_node, NULL, tbody, NULL);
@@ -8314,7 +8304,7 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
break;
case RESULT_DECL:
- /* When within an OpenMP context, notice uses of variables. */
+ /* When within an OMP context, notice uses of variables. */
if (gimplify_omp_ctxp)
omp_notice_variable (gimplify_omp_ctxp, *expr_p, true);
ret = GS_ALL_DONE;
@@ -8325,12 +8315,6 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
ret = GS_ALL_DONE;
break;
- case OACC_HOST_DATA:
- case OACC_DECLARE:
- sorry ("directive not yet implemented");
- ret = GS_ALL_DONE;
- break;
-
case OMP_PARALLEL:
gimplify_omp_parallel (expr_p, pre_p);
ret = GS_ALL_DONE;
@@ -8350,6 +8334,17 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
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_HOST_DATA:
+ case OACC_DECLARE:
+ sorry ("directive not yet implemented");
+ ret = GS_ALL_DONE;
+ break;
+
case OACC_KERNELS:
if (OACC_KERNELS_COMBINED (*expr_p))
sorry ("directive not yet implemented");
@@ -8366,11 +8361,6 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
ret = GS_ALL_DONE;
break;
- case OACC_CACHE:
- gimplify_oacc_cache (expr_p, pre_p);
- ret = GS_ALL_DONE;
- break;
-
case OACC_DATA:
case OMP_SECTIONS:
case OMP_SINGLE:
@@ -8381,10 +8371,10 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
ret = GS_ALL_DONE;
break;
- case OACC_UPDATE:
- case OMP_TARGET_UPDATE:
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;
break;
@@ -9007,9 +8997,7 @@ gimplify_body (tree fndecl, bool do_parms)
{
gcc_assert (gimplify_omp_ctxp == NULL);
if (lookup_attribute ("omp declare target", DECL_ATTRIBUTES (fndecl)))
- gimplify_omp_ctxp
- = new_omp_context ((enum omp_region_type) (ORT_TARGET
- | ORT_TARGET_OFFLOAD));
+ gimplify_omp_ctxp = new_omp_context (ORT_TARGET);
}
/* Unshare most shared trees in the body and in that of any nested functions.
deleted file mode 100644
@@ -1,56 +0,0 @@
-/* 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)
@@ -1,5 +1,5 @@
/* This file contains the definitions and documentation for the
- OpenMP builtins used in the GNU compiler.
+ Offloading and Multi Processing builtins used in the GNU compiler.
Copyright (C) 2005-2014 Free Software Foundation, Inc.
This file is part of GCC.
@@ -20,10 +20,44 @@ along with GCC; see the file COPYING3. If not see
/* Before including this file, you should define a macro:
+ DEF_GOACC_BUILTIN (ENUM, NAME, TYPE, ATTRS)
+ DEF_GOACC_BUILTIN_COMPILER (ENUM, NAME, TYPE, ATTRS)
DEF_GOMP_BUILTIN (ENUM, NAME, TYPE, ATTRS)
See builtins.def for details. */
+/* The reason why they aren't in gcc/builtins.def is that the Fortran front end
+ doesn't source those. */
+
+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 (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)
+
+DEF_GOACC_BUILTIN_COMPILER (BUILT_IN_ACC_ON_DEVICE, "acc_on_device",
+ BT_FN_INT_INT, ATTR_CONST_NOTHROW_LEAF_LIST)
+
DEF_GOMP_BUILTIN (BUILT_IN_OMP_GET_THREAD_NUM, "omp_get_thread_num",
BT_FN_INT, ATTR_CONST_NOTHROW_LEAF_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_OMP_GET_NUM_THREADS, "omp_get_num_threads",
@@ -1,6 +1,7 @@
-/* Lowering pass for OpenMP directives. Converts OpenMP directives
- into explicit calls to the runtime library (libgomp) and data
- marshalling to implement data sharing and copying clauses.
+/* Lowering pass for OMP directives. Converts OMP directives into explicit
+ calls to the runtime library (libgomp), data marshalling to implement data
+ sharing and copying clauses, offloading to accelerators, and more.
+
Contributed by Diego Novillo <dnovillo@redhat.com>
Copyright (C) 2005-2014 Free Software Foundation, Inc.
@@ -92,7 +93,7 @@ along with GCC; see the file COPYING3. If not see
#include "gomp-constants.h"
-/* Lowering of OpenMP parallel and workshare constructs proceeds in two
+/* Lowering of OMP parallel and workshare constructs proceeds in two
phases. The first phase scans the function looking for OMP statements
and then for variables that must be replaced to satisfy data sharing
clauses. The second phase expands code for the constructs, as well as
@@ -100,10 +101,10 @@ along with GCC; see the file COPYING3. If not see
expressions.
Final code generation is done by pass_expand_omp. The flowgraph is
- scanned for parallel regions which are then moved to a new
- function, to be invoked by the thread library. */
+ scanned for regions which are then moved to a new
+ function, to be invoked by the thread library, or offloaded. */
-/* Parallel region information. Every parallel and workshare
+/* OMP region information. Every parallel and workshare
directive is enclosed between two markers, the OMP_* directive
and a corresponding OMP_RETURN statement. */
@@ -142,6 +143,12 @@ struct omp_region
bool is_combined_parallel;
};
+/* Levels of parallelism as defined by OpenACC. Increasing numbers
+ correspond to deeper loop nesting levels. */
+#define MASK_GANG 1
+#define MASK_WORKER 2
+#define MASK_VECTOR 4
+
/* Context structure. Used to store information about each parallel
directive in the code. */
@@ -199,6 +206,14 @@ typedef struct omp_context
/* True if this construct can be cancelled. */
bool cancellable;
+
+ /* For OpenACC loops, a mask of gang, worker and vector used at
+ levels below this one. */
+ int gwv_below;
+ /* For OpenACC loops, a mask of gang, worker and vector used at
+ this level and above. For parallel and kernels clauses, a mask
+ indicating which of num_gangs/num_workers/num_vectors was used. */
+ int gwv_this;
} omp_context;
/* A structure holding the elements of:
@@ -350,7 +365,7 @@ static void lower_omp (gimple_seq *, omp_context *);
static tree lookup_decl_in_outer_ctx (tree, omp_context *);
static tree maybe_lookup_decl_in_outer_ctx (tree, omp_context *);
-/* Find an OpenMP clause of type KIND within CLAUSES. */
+/* Find an OMP clause of type KIND within CLAUSES. */
tree
find_omp_clause (tree clauses, enum omp_clause_code kind)
@@ -424,9 +439,9 @@ extract_omp_for_data (gomp_for *for_stmt, struct omp_for_data *fd,
else
fd->loops = &fd->loop;
- fd->have_nowait = (gimple_omp_for_kind (for_stmt) != GF_OMP_FOR_KIND_FOR);
+ fd->have_nowait = distribute || simd;
fd->have_ordered = false;
- fd->sched_kind = /* TODO: OACC_LOOP */ OMP_CLAUSE_SCHEDULE_STATIC;
+ fd->sched_kind = OMP_CLAUSE_SCHEDULE_STATIC;
fd->chunk_size = NULL_TREE;
if (gimple_omp_for_kind (fd->for_stmt) == GF_OMP_FOR_KIND_CILKFOR)
fd->sched_kind = OMP_CLAUSE_SCHEDULE_CILKFOR;
@@ -696,7 +711,6 @@ extract_omp_for_data (gomp_for *for_stmt, struct omp_for_data *fd,
fd->loop.cond_code = LT_EXPR;
}
- //TODO
/* 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. */
@@ -991,8 +1005,6 @@ maybe_lookup_field (tree var, omp_context *ctx)
static inline tree
lookup_oacc_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;
@@ -1020,7 +1032,7 @@ 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));
+ gcc_assert (!is_gimple_omp_oacc (shared_ctx->stmt));
/* ??? Trivially accessible from anywhere. But why would we even
be passing an address in this case? Should we simply assert
@@ -1227,7 +1239,7 @@ install_var_field (tree var, bool by_ref, int mask, omp_context *ctx)
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));
+ || !is_gimple_omp_oacc (ctx->stmt));
type = TREE_TYPE (var);
if (mask & 4)
@@ -1504,7 +1516,6 @@ new_omp_context (gimple stmt, omp_context *outer_ctx)
ctx->cb = outer_ctx->cb;
ctx->cb.block = NULL;
ctx->depth = outer_ctx->depth + 1;
- /* FIXME: handle reductions recursively. */
ctx->reduction_map = outer_ctx->reduction_map;
}
else
@@ -1519,7 +1530,6 @@ new_omp_context (gimple stmt, omp_context *outer_ctx)
ctx->cb.eh_lp_nr = 0;
ctx->cb.transform_call_graph_edges = CB_CGE_MOVE;
ctx->depth = 1;
- //TODO ctx->reduction_map = TODO;
}
ctx->cb.decl_map = new hash_map<tree, tree>;
@@ -1679,7 +1689,6 @@ 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)
@@ -1714,7 +1723,6 @@ 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))
@@ -1722,16 +1730,14 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
/* FALLTHRU */
case OMP_CLAUSE_FIRSTPRIVATE:
- if (is_gimple_omp_oacc_specifically (ctx->stmt))
+ if (is_gimple_omp_oacc (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:
+ case OMP_CLAUSE_LINEAR:
decl = OMP_CLAUSE_DECL (c);
do_private:
if (is_variable_sized (decl))
@@ -1757,10 +1763,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
install_var_field (decl, by_ref, 3, ctx);
}
install_var_local (decl, ctx);
- //TODO
- if (is_gimple_omp_oacc_specifically (ctx->stmt))
- {
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
+ if (is_gimple_omp_oacc (ctx->stmt)
+ && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
{
/* Create a decl for the reduction array. */
tree var = OMP_CLAUSE_DECL (c);
@@ -1780,11 +1784,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
(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);
@@ -1793,18 +1795,17 @@ 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:
@@ -1813,9 +1814,6 @@ 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:
@@ -1825,29 +1823,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
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);
@@ -1858,32 +1834,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)
- {
- 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
- target regions that are not offloaded; there is nothing to map for
+ regions that are not offloaded; there is nothing to map for
those. */
if (!is_gimple_omp_offloaded (ctx->stmt)
&& !POINTER_TYPE_P (TREE_TYPE (decl)))
break;
}
-#if 0
- /* In target regions that are not offloaded, libgomp won't pay
- attention to OMP_CLAUSE_MAP_FORCE_DEVICEPTR -- but I think we need
- to handle it here anyway, in order to create a visible copy of the
- variable. */
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
- {
- if (!is_gimple_omp_offloaded (ctx->stmt))
- break;
- }
-#endif
if (DECL_P (decl))
{
if (DECL_SIZE (decl)
@@ -1899,10 +1860,6 @@ 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)
@@ -1950,23 +1907,19 @@ 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:
- if (is_gimple_omp_oacc_specifically (ctx->stmt))
- {
- sorry ("clause not supported yet");
- break;
- }
- /* FALLTHRU */
- case OMP_CLAUSE_COLLAPSE:
case OMP_CLAUSE_ASYNC:
case OMP_CLAUSE_WAIT:
+ case OMP_CLAUSE_GANG:
+ case OMP_CLAUSE_WORKER:
+ case OMP_CLAUSE_VECTOR:
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)
@@ -1975,11 +1928,10 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
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:
+ case OMP_CLAUSE_AUTO:
+ case OMP_CLAUSE_SEQ:
sorry ("Clause not supported yet");
break;
@@ -1993,7 +1945,6 @@ 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))
@@ -2003,17 +1954,15 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
/* FALLTHRU */
case OMP_CLAUSE_FIRSTPRIVATE:
- if (is_gimple_omp_oacc_specifically (ctx->stmt))
+ if (is_gimple_omp_oacc (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:
case OMP_CLAUSE_PRIVATE:
+ case OMP_CLAUSE_REDUCTION:
+ case OMP_CLAUSE_LINEAR:
decl = OMP_CLAUSE_DECL (c);
if (is_variable_sized (decl))
install_var_local (decl, ctx);
@@ -2029,7 +1978,6 @@ 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;
@@ -2039,26 +1987,15 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
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));
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)
- {
- 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)))
@@ -2070,9 +2007,6 @@ 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);
@@ -2088,6 +2022,7 @@ 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:
@@ -2096,6 +2031,7 @@ 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:
@@ -2107,24 +2043,22 @@ 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_WAIT:
case OMP_CLAUSE_NUM_GANGS:
case OMP_CLAUSE_NUM_WORKERS:
case OMP_CLAUSE_VECTOR_LENGTH:
- case OMP_CLAUSE_WAIT:
+ case OMP_CLAUSE_GANG:
+ case OMP_CLAUSE_WORKER:
+ case OMP_CLAUSE_VECTOR:
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:
+ case OMP_CLAUSE_AUTO:
+ case OMP_CLAUSE_SEQ:
sorry ("Clause not supported yet");
break;
@@ -2133,9 +2067,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
}
}
+ gcc_checking_assert (!scan_array_reductions
+ || !is_gimple_omp_oacc (ctx->stmt));
if (scan_array_reductions)
- {
- gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt));
for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION
&& OMP_CLAUSE_REDUCTION_PLACEHOLDER (c))
@@ -2149,7 +2083,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LINEAR
&& OMP_CLAUSE_LINEAR_GIMPLE_SEQ (c))
scan_omp (&OMP_CLAUSE_LINEAR_GIMPLE_SEQ (c), ctx);
- }
}
/* Create a new name for omp child function. Returns an identifier. If
@@ -2218,8 +2151,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);
+ gcc_checking_assert (!is_gimple_omp_oacc (ctx->stmt)
+ || !task_copy);
if (!task_copy)
ctx->cb.dst_fn = decl;
else
@@ -2575,17 +2508,84 @@ finish_taskreg_scan (omp_context *ctx)
}
+static omp_context *
+enclosing_target_ctx (omp_context *ctx)
+{
+ while (ctx != NULL
+ && gimple_code (ctx->stmt) != GIMPLE_OMP_TARGET)
+ ctx = ctx->outer;
+ gcc_assert (ctx != NULL);
+ return ctx;
+}
+
+static bool
+oacc_loop_or_target_p (gimple stmt)
+{
+ enum gimple_code outer_type = gimple_code (stmt);
+ return ((outer_type == GIMPLE_OMP_TARGET
+ && ((gimple_omp_target_kind (stmt)
+ == GF_OMP_TARGET_KIND_OACC_PARALLEL)
+ || (gimple_omp_target_kind (stmt)
+ == GF_OMP_TARGET_KIND_OACC_KERNELS)))
+ || (outer_type == GIMPLE_OMP_FOR
+ && gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_OACC_LOOP));
+}
+
/* Scan a GIMPLE_OMP_FOR. */
static void
scan_omp_for (gomp_for *stmt, omp_context *outer_ctx)
{
+ enum gimple_code outer_type = GIMPLE_ERROR_MARK;
omp_context *ctx;
size_t i;
+ tree clauses = gimple_omp_for_clauses (stmt);
+
+ if (outer_ctx)
+ outer_type = gimple_code (outer_ctx->stmt);
ctx = new_omp_context (stmt, outer_ctx);
- scan_sharing_clauses (gimple_omp_for_clauses (stmt), ctx);
+ if (is_gimple_omp_oacc (stmt))
+ {
+ if (outer_ctx && outer_type == GIMPLE_OMP_FOR)
+ ctx->gwv_this = outer_ctx->gwv_this;
+ for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+ {
+ int val;
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_GANG)
+ val = MASK_GANG;
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_WORKER)
+ val = MASK_WORKER;
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_VECTOR)
+ val = MASK_VECTOR;
+ else
+ continue;
+ ctx->gwv_this |= val;
+ if (!outer_ctx)
+ {
+ /* Skip; not nested inside a region. */
+ continue;
+ }
+ if (!oacc_loop_or_target_p (outer_ctx->stmt))
+ {
+ /* Skip; not nested inside an OpenACC region. */
+ continue;
+ }
+ if (outer_type == GIMPLE_OMP_FOR)
+ outer_ctx->gwv_below |= val;
+ if (OMP_CLAUSE_OPERAND (c, 0) != NULL_TREE)
+ {
+ omp_context *enclosing = enclosing_target_ctx (outer_ctx);
+ if (gimple_omp_target_kind (enclosing->stmt)
+ == GF_OMP_TARGET_KIND_OACC_PARALLEL)
+ error_at (gimple_location (stmt),
+ "no arguments allowed to gang, worker and vector clauses inside parallel");
+ }
+ }
+ }
+
+ scan_sharing_clauses (clauses, ctx);
scan_omp (gimple_omp_for_pre_body_ptr (stmt), ctx);
for (i = 0; i < gimple_omp_for_collapse (stmt); i++)
@@ -2596,6 +2596,19 @@ scan_omp_for (gomp_for *stmt, omp_context *outer_ctx)
scan_omp_op (gimple_omp_for_incr_ptr (stmt, i), ctx);
}
scan_omp (gimple_omp_body_ptr (stmt), ctx);
+
+ if (is_gimple_omp_oacc (stmt))
+ {
+ if (ctx->gwv_this & ctx->gwv_below)
+ error_at (gimple_location (stmt),
+ "gang, worker and vector may occur only once in a loop nest");
+ else if (ctx->gwv_below != 0
+ && ctx->gwv_this > ctx->gwv_below)
+ error_at (gimple_location (stmt),
+ "gang, worker and vector must occur in this order in a loop nest");
+ if (outer_ctx && outer_type == GIMPLE_OMP_FOR)
+ outer_ctx->gwv_below |= ctx->gwv_below;
+ }
}
/* Scan an OpenMP sections directive. */
@@ -2643,6 +2656,7 @@ scan_omp_target (gomp_target *stmt, omp_context *outer_ctx)
omp_context *ctx;
tree name;
bool offloaded = is_gimple_omp_offloaded (stmt);
+ tree clauses = gimple_omp_target_clauses (stmt);
ctx = new_omp_context (stmt, outer_ctx);
ctx->field_map = splay_tree_new (splay_tree_compare_pointers, 0, 0);
@@ -2656,7 +2670,7 @@ scan_omp_target (gomp_target *stmt, omp_context *outer_ctx)
TYPE_NAME (ctx->record_type) = name;
if (offloaded)
{
- if (is_gimple_omp_oacc_specifically (stmt))
+ if (is_gimple_omp_oacc (stmt))
ctx->reduction_map = splay_tree_new (splay_tree_compare_pointers,
0, 0);
@@ -2664,7 +2678,20 @@ scan_omp_target (gomp_target *stmt, omp_context *outer_ctx)
gimple_omp_target_set_child_fn (stmt, ctx->cb.dst_fn);
}
- scan_sharing_clauses (gimple_omp_target_clauses (stmt), ctx);
+ if (is_gimple_omp_oacc (stmt))
+ {
+ for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+ {
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_NUM_GANGS)
+ ctx->gwv_this |= MASK_GANG;
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_NUM_WORKERS)
+ ctx->gwv_this |= MASK_WORKER;
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_VECTOR_LENGTH)
+ ctx->gwv_this |= MASK_VECTOR;
+ }
+ }
+
+ scan_sharing_clauses (clauses, ctx);
scan_omp (gimple_omp_body_ptr (stmt), ctx);
if (TYPE_FIELDS (ctx->record_type) == NULL)
@@ -2701,16 +2728,14 @@ scan_omp_teams (gomp_teams *stmt, omp_context *outer_ctx)
static bool
check_omp_nesting_restrictions (gimple stmt, omp_context *ctx)
{
- /* TODO: Some OpenACC/OpenMP nesting should be allowed. */
-
/* No nesting of non-OpenACC STMT (that is, an OpenMP one, or a GOMP builtin)
inside an OpenACC CTX. */
if (!(is_gimple_omp (stmt)
- && is_gimple_omp_oacc_specifically (stmt)))
+ && is_gimple_omp_oacc (stmt)))
{
for (omp_context *ctx_ = ctx; ctx_ != NULL; ctx_ = ctx_->outer)
if (is_gimple_omp (ctx_->stmt)
- && is_gimple_omp_oacc_specifically (ctx_->stmt))
+ && is_gimple_omp_oacc (ctx_->stmt))
{
error_at (gimple_location (stmt),
"non-OpenACC construct inside of OpenACC region");
@@ -2982,7 +3007,7 @@ check_omp_nesting_restrictions (gimple stmt, omp_context *ctx)
if (gimple_code (ctx->stmt) != GIMPLE_OMP_TARGET)
{
if (is_gimple_omp (stmt)
- && is_gimple_omp_oacc_specifically (stmt)
+ && is_gimple_omp_oacc (stmt)
&& is_gimple_omp (ctx->stmt))
{
error_at (gimple_location (stmt),
@@ -3016,21 +3041,21 @@ check_omp_nesting_restrictions (gimple stmt, omp_context *ctx)
}
/* OpenACC/OpenMP mismatch? */
- if (is_gimple_omp_oacc_specifically (stmt)
- != is_gimple_omp_oacc_specifically (ctx->stmt))
+ if (is_gimple_omp_oacc (stmt)
+ != is_gimple_omp_oacc (ctx->stmt))
{
error_at (gimple_location (stmt),
"%s %s construct inside of %s %s region",
- (is_gimple_omp_oacc_specifically (stmt)
+ (is_gimple_omp_oacc (stmt)
? "OpenACC" : "OpenMP"), stmt_name,
- (is_gimple_omp_oacc_specifically (ctx->stmt)
+ (is_gimple_omp_oacc (ctx->stmt)
? "OpenACC" : "OpenMP"), ctx_stmt_name);
return false;
}
if (is_gimple_omp_offloaded (ctx->stmt))
{
/* No GIMPLE_OMP_TARGET inside offloaded OpenACC CTX. */
- if (is_gimple_omp_oacc_specifically (ctx->stmt))
+ if (is_gimple_omp_oacc (ctx->stmt))
{
error_at (gimple_location (stmt),
"%s construct inside of %s region",
@@ -3039,7 +3064,7 @@ check_omp_nesting_restrictions (gimple stmt, omp_context *ctx)
}
else
{
- gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
+ gcc_checking_assert (!is_gimple_omp_oacc (stmt));
warning_at (gimple_location (stmt), 0,
"%s construct inside of %s region",
stmt_name, ctx_stmt_name);
@@ -3057,7 +3082,7 @@ check_omp_nesting_restrictions (gimple stmt, omp_context *ctx)
/* Helper function scan_omp.
Callback for walk_tree or operators in walk_gimple_stmt used to
- scan for OpenMP directives in TP. */
+ scan for OMP directives in TP. */
static tree
scan_omp_1_op (tree *tp, int *walk_subtrees, void *data)
@@ -3120,7 +3145,7 @@ setjmp_or_longjmp_p (const_tree fndecl)
/* Helper function for scan_omp.
- Callback for walk_gimple_stmt used to scan for OpenMP directives in
+ Callback for walk_gimple_stmt used to scan for OMP directives in
the current statement in GSI. */
static tree
@@ -3133,7 +3158,7 @@ scan_omp_1_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
if (gimple_has_location (stmt))
input_location = gimple_location (stmt);
- /* Check the OpenMP nesting restrictions. */
+ /* Check the nesting restrictions. */
bool remove = false;
if (is_gimple_omp (stmt))
remove = !check_omp_nesting_restrictions (stmt, ctx);
@@ -3241,7 +3266,7 @@ scan_omp_1_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
/* Scan all the statements starting at the current statement. CTX
- contains context information about the OpenMP directives and
+ contains context information about the OMP directives and
clauses found during the scan. */
static void
@@ -3507,8 +3532,6 @@ 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 ();
@@ -4513,7 +4536,7 @@ lower_reduction_clauses (tree clauses, gimple_seq *stmt_seqp, omp_context *ctx)
return;
/* Initialize thread info for OpenACC. */
- if (is_gimple_omp_oacc_specifically (ctx->stmt))
+ if (is_gimple_omp_oacc (ctx->stmt))
{
/* Get the current thread id. */
tree call = builtin_decl_explicit (BUILT_IN_GOACC_GET_THREAD_NUM);
@@ -4544,10 +4567,14 @@ 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 (ctx->stmt))
+ {
+ gcc_checking_assert (!OMP_CLAUSE_REDUCTION_PLACEHOLDER (c));
+
+ oacc_lower_reduction_var_helper (stmt_seqp, ctx, tid, var, new_var);
+ }
+ else if (count == 1)
{
- if (!is_gimple_omp_oacc_specifically (ctx->stmt))
- {
tree addr = build_fold_addr_expr_loc (clause_loc, ref);
addr = save_expr (addr);
@@ -4556,16 +4583,8 @@ lower_reduction_clauses (tree clauses, gimple_seq *stmt_seqp, omp_context *ctx)
x = build2 (OMP_ATOMIC, void_type_node, addr, x);
gimplify_and_add (x, stmt_seqp);
return;
- }
- else
- {
- oacc_lower_reduction_var_helper (stmt_seqp, ctx, tid, var,
- new_var);
- return;
- }
}
-
- if (OMP_CLAUSE_REDUCTION_PLACEHOLDER (c))
+ else if (OMP_CLAUSE_REDUCTION_PLACEHOLDER (c))
{
tree placeholder = OMP_CLAUSE_REDUCTION_PLACEHOLDER (c);
@@ -4582,21 +4601,13 @@ lower_reduction_clauses (tree clauses, gimple_seq *stmt_seqp, omp_context *ctx)
}
else
{
- if (is_gimple_omp_oacc_specifically (ctx->stmt))
- {
- oacc_lower_reduction_var_helper (stmt_seqp, ctx, tid, var,
- new_var);
- }
- else
- {
- x = build2 (code, TREE_TYPE (ref), ref, new_var);
- ref = build_outer_var_ref (var, ctx);
- gimplify_assign (ref, x, &sub_seq);
- }
+ x = build2 (code, TREE_TYPE (ref), ref, new_var);
+ ref = build_outer_var_ref (var, ctx);
+ gimplify_assign (ref, x, &sub_seq);
}
}
- if (is_gimple_omp_oacc_specifically (ctx->stmt))
+ if (is_gimple_omp_oacc (ctx->stmt))
return;
stmt = gimple_build_call (builtin_decl_explicit (BUILT_IN_GOMP_ATOMIC_START),
@@ -4617,8 +4628,6 @@ 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))
@@ -4669,8 +4678,6 @@ 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))
@@ -4762,8 +4769,6 @@ 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)
@@ -6065,8 +6070,6 @@ 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;
@@ -6136,9 +6139,6 @@ 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;
@@ -6167,9 +6167,6 @@ 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,
@@ -6188,9 +6185,6 @@ 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);
@@ -6510,9 +6504,9 @@ expand_omp_for_static_nochunk (struct omp_region *region,
tree *counts = NULL;
tree n1, n2, step;
- gcc_assert ((gimple_omp_for_kind (fd->for_stmt)
- != GF_OMP_FOR_KIND_OACC_LOOP)
- || !inner_stmt);
+ gcc_checking_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))
@@ -6599,28 +6593,24 @@ expand_omp_for_static_nochunk (struct omp_region *region,
{
case GF_OMP_FOR_KIND_FOR:
nthreads = builtin_decl_explicit (BUILT_IN_OMP_GET_NUM_THREADS);
- nthreads = build_call_expr (nthreads, 0);
threadid = builtin_decl_explicit (BUILT_IN_OMP_GET_THREAD_NUM);
- threadid = build_call_expr (threadid, 0);
break;
case GF_OMP_FOR_KIND_DISTRIBUTE:
nthreads = builtin_decl_explicit (BUILT_IN_OMP_GET_NUM_TEAMS);
- nthreads = build_call_expr (nthreads, 0);
threadid = builtin_decl_explicit (BUILT_IN_OMP_GET_TEAM_NUM);
- threadid = build_call_expr (threadid, 0);
break;
case GF_OMP_FOR_KIND_OACC_LOOP:
nthreads = builtin_decl_explicit (BUILT_IN_GOACC_GET_NUM_THREADS);
- nthreads = build_call_expr (nthreads, 0);
threadid = builtin_decl_explicit (BUILT_IN_GOACC_GET_THREAD_NUM);
- threadid = build_call_expr (threadid, 0);
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);
+ 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);
@@ -6630,9 +6620,6 @@ 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);
@@ -6798,11 +6785,11 @@ expand_omp_for_static_nochunk (struct omp_region *region,
gsi = gsi_last_bb (exit_bb);
if (!gimple_omp_return_nowait_p (gsi_stmt (gsi)))
{
- gcc_assert (gimple_omp_for_kind (fd->for_stmt)
- != GF_OMP_FOR_KIND_OACC_LOOP);
-
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_checking_assert (t == NULL_TREE);
+ else
+ gsi_insert_after (&gsi, build_omp_barrier (t), GSI_SAME_STMT);
}
gsi_remove (&gsi, true);
@@ -6820,9 +6807,6 @@ 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;
}
@@ -6910,9 +6894,9 @@ expand_omp_for_static_chunk (struct omp_region *region,
tree *counts = NULL;
tree n1, n2, step;
- gcc_assert ((gimple_omp_for_kind (fd->for_stmt)
- != GF_OMP_FOR_KIND_OACC_LOOP)
- || !inner_stmt);
+ gcc_checking_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))
@@ -7003,28 +6987,24 @@ expand_omp_for_static_chunk (struct omp_region *region,
{
case GF_OMP_FOR_KIND_FOR:
nthreads = builtin_decl_explicit (BUILT_IN_OMP_GET_NUM_THREADS);
- nthreads = build_call_expr (nthreads, 0);
threadid = builtin_decl_explicit (BUILT_IN_OMP_GET_THREAD_NUM);
- threadid = build_call_expr (threadid, 0);
break;
case GF_OMP_FOR_KIND_DISTRIBUTE:
nthreads = builtin_decl_explicit (BUILT_IN_OMP_GET_NUM_TEAMS);
- nthreads = build_call_expr (nthreads, 0);
threadid = builtin_decl_explicit (BUILT_IN_OMP_GET_TEAM_NUM);
- threadid = build_call_expr (threadid, 0);
break;
case GF_OMP_FOR_KIND_OACC_LOOP:
nthreads = builtin_decl_explicit (BUILT_IN_GOACC_GET_NUM_THREADS);
- nthreads = build_call_expr (nthreads, 0);
threadid = builtin_decl_explicit (BUILT_IN_GOACC_GET_THREAD_NUM);
- threadid = build_call_expr (threadid, 0);
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);
+ 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);
@@ -7034,9 +7014,6 @@ 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);
@@ -7219,11 +7196,11 @@ expand_omp_for_static_chunk (struct omp_region *region,
gsi = gsi_last_bb (exit_bb);
if (!gimple_omp_return_nowait_p (gsi_stmt (gsi)))
{
- gcc_assert (gimple_omp_for_kind (fd->for_stmt)
- != GF_OMP_FOR_KIND_OACC_LOOP);
-
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_checking_assert (t == NULL_TREE);
+ else
+ gsi_insert_after (&gsi, build_omp_barrier (t), GSI_SAME_STMT);
}
gsi_remove (&gsi, true);
@@ -7236,9 +7213,6 @@ 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;
}
@@ -7841,7 +7815,7 @@ expand_omp_simd (struct omp_region *region, struct omp_for_data *fd)
}
-/* Expand the OpenMP loop defined by REGION. */
+/* Expand the OMP loop defined by REGION. */
static void
expand_omp_for (struct omp_region *region, gimple inner_stmt)
@@ -8827,8 +8801,8 @@ expand_omp_target (struct omp_region *region)
/* Supported by expand_omp_taskreg, but not here. */
if (child_cfun != NULL)
- gcc_assert (!child_cfun->cfg);
- gcc_assert (!gimple_in_ssa_p (cfun));
+ gcc_checking_assert (!child_cfun->cfg);
+ gcc_checking_assert (!gimple_in_ssa_p (cfun));
entry_bb = region->entry;
exit_bb = region->exit;
@@ -8913,8 +8887,8 @@ expand_omp_target (struct omp_region *region)
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_code (entry_stmt));
+ gcc_assert (stmt
+ && gimple_code (stmt) == gimple_code (entry_stmt));
gsi_remove (&gsi, true);
e = split_block (entry_bb, stmt);
entry_bb = e->dest;
@@ -9047,9 +9021,9 @@ expand_omp_target (struct omp_region *region)
{
/* 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);
+ gcc_checking_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);
@@ -9526,12 +9500,12 @@ public:
/* opt_pass methods: */
virtual unsigned int execute (function *)
{
- bool gate = ((flag_openacc != 0 || flag_openmp != 0
- || flag_openmp_simd != 0 || flag_cilkplus != 0)
+ bool gate = ((flag_cilkplus != 0 || flag_openacc != 0 || flag_openmp != 0
+ || flag_openmp_simd != 0)
&& !seen_error ());
/* This pass always runs, to provide PROP_gimple_eomp.
- But there is nothing to do unless -fopenmp is given. */
+ But often, there is nothing to do. */
if (!gate)
return 0;
@@ -9587,6 +9561,8 @@ make_pass_expand_omp_ssa (gcc::context *ctxt)
return new pass_expand_omp_ssa (ctxt);
}
+/* Routines to lower OMP directives into OMP-GIMPLE. */
+
/* Helper function to preform, potentially COMPLEX_TYPE, operation and
convert it to gimple. */
static void
@@ -9625,8 +9601,6 @@ oacc_gimple_assign (tree dest, tree_code op, tree src, gimple_seq *seq)
tree i = create_tmp_var (TREE_TYPE (TREE_TYPE (dest)));
tree result;
- gcc_assert (op == PLUS_EXPR || op == MULT_EXPR);
-
if (op == PLUS_EXPR)
{
stmt = gimple_build_assign (r, op, rdest, rsrc);
@@ -9662,6 +9636,8 @@ oacc_gimple_assign (tree dest, tree_code op, tree src, gimple_seq *seq)
stmt = gimple_build_assign (i, PLUS_EXPR, ad, bc);
gimple_seq_add_stmt (seq, stmt);
}
+ else
+ gcc_unreachable ();
result = build2 (COMPLEX_EXPR, TREE_TYPE (dest), r, i);
gimplify_assign (dest, result, seq);
@@ -9670,27 +9646,26 @@ oacc_gimple_assign (tree dest, tree_code op, tree src, gimple_seq *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. FIXME: This function assumes that there are
- vector_length threads in total. */
+ reduction. */
static void
oacc_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;
- /* Find the innermost PARALLEL openmp context. FIXME: OpenACC kernels
- may require extra care unless they are converted to openmp for loops. */
+ /* Find the innermost OpenACC parallel context. */
if (gimple_code (ctx->stmt) == GIMPLE_OMP_TARGET
&& (gimple_omp_target_kind (ctx->stmt)
== GF_OMP_TARGET_KIND_OACC_PARALLEL))
octx = ctx;
else
octx = ctx->outer;
+ gcc_checking_assert (gimple_code (octx->stmt) == GIMPLE_OMP_TARGET
+ && (gimple_omp_target_kind (octx->stmt)
+ == GF_OMP_TARGET_KIND_OACC_PARALLEL));
/* Extract the clauses. */
oc = gimple_omp_target_clauses (octx->stmt);
@@ -9721,10 +9696,7 @@ oacc_initialize_reduction_data (tree clauses, tree nthreads,
size = create_tmp_var (sizetype);
gimplify_assign (size, fold_build1 (NOP_EXPR, sizetype, t), stmt_seqp);
- /* Now allocate memory for it. FIXME: Allocating memory for the
- reduction array may be unnecessary once the final reduction is able
- to be preformed on the accelerator. Instead of allocating memory on
- the host side, it could just be allocated on the accelerator. */
+ /* 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);
@@ -9733,8 +9705,6 @@ oacc_initialize_reduction_data (tree clauses, tree nthreads,
/* Map this array into the accelerator. */
/* Add the reduction array to the list of clauses. */
- /* FIXME: Currently, these variables must be placed in the outer
- most clause so that copy-out works. */
tree x = array;
t = build_omp_clause (gimple_location (ctx->stmt), OMP_CLAUSE_MAP);
OMP_CLAUSE_MAP_KIND (t) = OMP_CLAUSE_MAP_FORCE_FROM;
@@ -9758,8 +9728,6 @@ static void
oacc_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;
@@ -9874,8 +9842,6 @@ static void
oacc_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;
@@ -9901,12 +9867,11 @@ oacc_process_reduction_data (gimple_seq *body, gimple_seq *in_stmt_seqp,
tree clauses, nthreads, t, c, acc_device, acc_device_host, call,
enter, exit;
bool reduction_found = false;
-
+
gimple stmt = gsi_stmt (gsi);
switch (gimple_code (stmt))
{
- /* FIXME: A reduction may also appear in an oacc parallel. */
case GIMPLE_OMP_FOR:
clauses = gimple_omp_for_clauses (stmt);
@@ -9985,8 +9950,6 @@ oacc_process_reduction_data (gimple_seq *body, gimple_seq *in_stmt_seqp,
}
}
-/* Routines to lower OpenMP directives into OMP-GIMPLE. */
-
/* 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
@@ -10558,7 +10521,7 @@ lower_omp_for_lastprivate (struct omp_for_data *fd, gimple_seq *body_p,
}
-/* Lower code for an OpenMP loop directive. */
+/* Lower code for an OMP loop directive. */
static void
lower_omp_for (gimple_stmt_iterator *gsi_p, omp_context *ctx)
@@ -10600,8 +10563,6 @@ 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;
@@ -10791,8 +10752,6 @@ 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;
@@ -11255,7 +11214,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
irlist = NULL;
orlist = NULL;
if (offloaded
- && is_gimple_omp_oacc_specifically (stmt))
+ && is_gimple_omp_oacc (stmt))
oacc_process_reduction_data (&tgt_body, &irlist, &orlist, ctx);
for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
@@ -11284,7 +11243,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
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));
+ gcc_assert (is_gimple_omp_oacc (stmt));
break;
default:
gcc_unreachable ();
@@ -11293,10 +11252,6 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
/* 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))
{
@@ -11323,10 +11278,6 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
{
x = build_receiver_ref (var, true, ctx);
tree new_var = lookup_decl (var, ctx);
- 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)
@@ -11372,7 +11323,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
TREE_STATIC (TREE_VEC_ELT (t, 1)) = 1;
tree tkind_type;
int talign_shift;
- if (is_gimple_omp_oacc_specifically (stmt))
+ if (is_gimple_omp_oacc (stmt))
{
tkind_type = short_unsigned_type_node;
talign_shift = 8;
@@ -11449,14 +11400,10 @@ 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);
- 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_oacc_reduction (var, ctx))
{
- gcc_assert (offloaded
- && is_gimple_omp_oacc_specifically (stmt));
+ gcc_checking_assert (offloaded
+ && is_gimple_omp_oacc (stmt));
gimplify_assign (x, var, &ilist);
}
else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
@@ -11526,10 +11473,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));
+ gcc_checking_assert (tkind
+ < (HOST_WIDE_INT_C (1U) << talign_shift));
talign = ceil_log2 (talign);
tkind |= talign << talign_shift;
- gcc_assert (tkind <= tree_to_uhwi (TYPE_MAX_VALUE (tkind_type)));
+ gcc_checking_assert (tkind
+ <= tree_to_uhwi (TYPE_MAX_VALUE (tkind_type)));
CONSTRUCTOR_APPEND_ELT (vkind, purpose,
build_int_cstu (tkind_type, tkind));
if (nc && nc != c)
@@ -11672,7 +11621,7 @@ lower_omp_teams (gimple_stmt_iterator *gsi_p, omp_context *ctx)
/* Callback for lower_omp_1. Return non-NULL if *tp needs to be
regimplified. If DATA is non-NULL, lower_omp_1 is outside
- of OpenMP context, but with task_shared_vars set. */
+ of OMP context, but with task_shared_vars set. */
static tree
lower_omp_regimplify_p (tree *tp, int *walk_subtrees,
@@ -11712,7 +11661,7 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context *ctx)
memset (&wi, '\0', sizeof (wi));
/* If we have issued syntax errors, avoid doing any heavy lifting.
- Just replace the OpenMP directives with a NOP to avoid
+ Just replace the OMP directives with a NOP to avoid
confusing RTL expansion. */
if (seen_error () && is_gimple_omp (stmt))
{
@@ -11810,8 +11759,6 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context *ctx)
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:
@@ -11920,8 +11867,8 @@ execute_lower_omp (void)
/* This pass always runs, to provide PROP_gimple_lomp.
But often, there is nothing to do. */
- if (flag_openacc == 0 && flag_openmp == 0 && flag_openmp_simd == 0
- && flag_cilkplus == 0)
+ 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,
@@ -12000,8 +11947,8 @@ 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));
+ gcc_checking_assert (!branch_ctx || is_gimple_omp (branch_ctx));
+ gcc_checking_assert (!label_ctx || is_gimple_omp (label_ctx));
if (label_ctx == branch_ctx)
return false;
@@ -12020,16 +11967,16 @@ diagnose_sb_0 (gimple_stmt_iterator *gsi_p,
}
if (flag_openacc)
{
- if ((branch_ctx && is_gimple_omp_oacc_specifically (branch_ctx))
- || (label_ctx && is_gimple_omp_oacc_specifically (label_ctx)))
+ if ((branch_ctx && is_gimple_omp_oacc (branch_ctx))
+ || (label_ctx && is_gimple_omp_oacc (label_ctx)))
{
- gcc_assert (kind == NULL);
+ gcc_checking_assert (kind == NULL);
kind = "OpenACC";
}
}
if (kind == NULL)
{
- gcc_assert (flag_openmp);
+ gcc_checking_assert (flag_openmp);
kind = "OpenMP";
}
@@ -12110,7 +12057,7 @@ diagnose_sb_1 (gimple_stmt_iterator *gsi_p, bool *handled_ops_p,
case GIMPLE_OMP_TARGET:
case GIMPLE_OMP_TEAMS:
case GIMPLE_OMP_TASKGROUP:
- /* The minimal context here is just the current construct. */
+ /* The minimal context here is just the current OMP construct. */
inner_context = stmt;
wi->info = inner_context;
walk_gimple_seq (gimple_omp_body (stmt), diagnose_sb_1, NULL, wi);
@@ -12244,7 +12191,7 @@ diagnose_sb_2 (gimple_stmt_iterator *gsi_p, bool *handled_ops_p,
}
/* Called from tree-cfg.c::make_edges to create cfg edges for all relevant
- GIMPLE codes. */
+ GIMPLE_* codes. */
bool
make_gimple_omp_edges (basic_block bb, struct omp_region **region,
int *region_idx)
@@ -12427,7 +12374,7 @@ public:
/* opt_pass methods: */
virtual bool gate (function *)
{
- return flag_openacc || flag_openmp || flag_cilkplus;
+ return flag_cilkplus || flag_openacc || flag_openmp;
}
virtual unsigned int execute (function *)
{
@@ -1,6 +1,6 @@
extern int i;
-/* TODO: While the OpenACC specification does allow for certain kinds of
+/* While the OpenACC specification does allow for certain kinds of
nesting, we don't support many of these yet. */
void
f_acc_parallel (void)
@@ -19,7 +19,7 @@ f_acc_parallel (void)
}
}
-/* TODO: While the OpenACC specification does allow for certain kinds of
+/* While the OpenACC specification does allow for certain kinds of
nesting, we don't support many of these yet. */
void
f_acc_kernels (void)
@@ -39,7 +39,7 @@ end program test
! { dg-final { scan-tree-dump-times "collapse\\(2\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "independent" 1 "original" } }
-! { dg-final { scan-tree-dump-times "gang\\(3\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "gang\\(num: 3\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "worker\\(3\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "vector\\(5\\)" 1 "original" } }
@@ -222,19 +222,19 @@ enum omp_clause_code {
(c_parser_omp_variable_list). */
OMP_CLAUSE_ERROR = 0,
- /* OpenMP/OpenACC clause: private (variable_list). */
+ /* OpenACC/OpenMP clause: private (variable_list). */
OMP_CLAUSE_PRIVATE,
/* OpenMP clause: shared (variable_list). */
OMP_CLAUSE_SHARED,
- /* OpenMP/OpenACC clause: firstprivate (variable_list). */
+ /* OpenACC/OpenMP clause: firstprivate (variable_list). */
OMP_CLAUSE_FIRSTPRIVATE,
/* OpenMP clause: lastprivate (variable_list). */
OMP_CLAUSE_LASTPRIVATE,
- /* OpenMP/OpenACC 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
@@ -285,9 +285,9 @@ enum omp_clause_code {
/* 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
+ /* 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. */
@@ -296,13 +296,19 @@ enum omp_clause_code {
/* OpenACC clause: async [(integer-expression)]. */
OMP_CLAUSE_ASYNC,
- /* OpenACC clause/directive: wait [(integer-expression-list)]. */
+ /* OpenACC clause: wait [(integer-expression-list)]. */
OMP_CLAUSE_WAIT,
+ /* OpenACC clause: auto. */
+ OMP_CLAUSE_AUTO,
+
+ /* OpenACC clause: seq. */
+ OMP_CLAUSE_SEQ,
+
/* Internal clause: temporary for combined loops expansion. */
OMP_CLAUSE__LOOPTEMP_,
- /* OpenMP/OpenACC clause: if (scalar-expression). */
+ /* OpenACC/OpenMP clause: if (scalar-expression). */
OMP_CLAUSE_IF,
/* OpenMP clause: num_threads (integer-expression). */
@@ -315,13 +321,12 @@ enum omp_clause_code {
OMP_CLAUSE_NOWAIT,
/* OpenMP clause: ordered. */
- /* OpenACC clause: seq. */
OMP_CLAUSE_ORDERED,
/* OpenMP clause: default. */
OMP_CLAUSE_DEFAULT,
- /* OpenMP/OpenACC clause: collapse (constant-integer-expression). */
+ /* OpenACC/OpenMP clause: collapse (constant-integer-expression). */
OMP_CLAUSE_COLLAPSE,
/* OpenMP clause: untied. */
@@ -1453,7 +1453,6 @@ 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),
@@ -1510,7 +1509,6 @@ 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),
@@ -4156,7 +4154,7 @@ estimate_num_insns (gimple stmt, eni_weights *weights)
return (estimate_num_insns_seq (gimple_try_eval (stmt), weights)
+ estimate_num_insns_seq (gimple_try_cleanup (stmt), weights));
- /* OpenMP directives are generally very expensive. */
+ /* OMP directives are generally very expensive. */
case GIMPLE_OMP_RETURN:
case GIMPLE_OMP_SECTIONS_SWITCH:
@@ -627,8 +627,6 @@ 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;
@@ -833,7 +831,7 @@ static void note_nonlocal_vla_type (struct nesting_info *info, tree type);
/* A subroutine of convert_nonlocal_reference_op. Create a local variable
in the nested function with DECL_VALUE_EXPR set to reference the true
variable in the parent function. This is used both for debug info
- and in OpenMP lowering. */
+ and in OMP lowering. */
static tree
get_nonlocal_debug_decl (struct nesting_info *info, tree decl)
@@ -1355,7 +1353,6 @@ 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),
@@ -1383,8 +1380,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)
+ if (!is_gimple_omp_offloaded (stmt))
{
save_suppress = info->suppress_expansion;
convert_nonlocal_omp_clauses (gimple_omp_target_clauses_ptr (stmt),
@@ -1488,7 +1484,7 @@ convert_nonlocal_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
/* A subroutine of convert_local_reference. Create a local variable
in the parent function with DECL_VALUE_EXPR set to reference the
- field in FRAME. This is used both for debug info and in OpenMP
+ field in FRAME. This is used both for debug info and in OMP
lowering. */
static tree
@@ -1926,7 +1922,6 @@ 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),
@@ -1954,8 +1949,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)
+ if (!is_gimple_omp_offloaded (stmt))
{
save_suppress = info->suppress_expansion;
convert_local_omp_clauses (gimple_omp_target_clauses_ptr (stmt), wi);
@@ -2288,8 +2282,7 @@ convert_tramp_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
}
case GIMPLE_OMP_TARGET:
- gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
- if (gimple_omp_target_kind (stmt) != GF_OMP_TARGET_KIND_REGION)
+ if (!is_gimple_omp_offloaded (stmt))
{
*handled_ops_p = false;
return NULL_TREE;
@@ -2388,8 +2381,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)
+ if (!is_gimple_omp_offloaded (stmt))
{
walk_body (convert_gimple_call, NULL, info, gimple_omp_body_ptr (stmt));
break;
@@ -2426,7 +2418,6 @@ 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 */
@@ -2438,7 +2429,6 @@ 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;
@@ -667,10 +667,30 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, int flags)
break;
case OMP_CLAUSE_GANG:
- pp_string (pp, "gang(");
- dump_generic_node (pp, OMP_CLAUSE_GANG_EXPR (clause),
- spc, flags, false);
- pp_character(pp, ')');
+ pp_string (pp, "gang");
+ if (OMP_CLAUSE_GANG_EXPR (clause) != NULL_TREE)
+ {
+ pp_string (pp, "(num: ");
+ dump_generic_node (pp, OMP_CLAUSE_GANG_EXPR (clause),
+ spc, flags, false);
+ }
+ if (OMP_CLAUSE_GANG_STATIC_EXPR (clause) != NULL_TREE)
+ {
+ if (OMP_CLAUSE_GANG_EXPR (clause) == NULL_TREE)
+ pp_left_paren (pp);
+ else
+ pp_space (pp);
+ pp_string (pp, "static:");
+ if (OMP_CLAUSE_GANG_STATIC_EXPR (clause)
+ == integer_minus_one_node)
+ pp_character (pp, '*');
+ else
+ dump_generic_node (pp, OMP_CLAUSE_GANG_STATIC_EXPR (clause),
+ spc, flags, false);
+ }
+ if (OMP_CLAUSE_GANG_EXPR (clause) != NULL_TREE
+ || OMP_CLAUSE_GANG_STATIC_EXPR (clause) != NULL_TREE)
+ pp_right_paren (pp);
break;
case OMP_CLAUSE_ASYNC:
@@ -684,6 +704,11 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, int flags)
}
break;
+ case OMP_CLAUSE_AUTO:
+ case OMP_CLAUSE_SEQ:
+ pp_string (pp, omp_clause_code_name[OMP_CLAUSE_CODE (clause)]);
+ break;
+
case OMP_CLAUSE_WAIT:
pp_string (pp, "wait(");
dump_generic_node (pp, OMP_CLAUSE_WAIT_EXPR (clause),
@@ -692,17 +717,25 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, int flags)
break;
case OMP_CLAUSE_WORKER:
- pp_string (pp, "worker(");
- dump_generic_node (pp, OMP_CLAUSE_WORKER_EXPR (clause),
- spc, flags, false);
- pp_character(pp, ')');
+ pp_string (pp, "worker");
+ if (OMP_CLAUSE_WORKER_EXPR (clause) != NULL_TREE)
+ {
+ pp_left_paren (pp);
+ dump_generic_node (pp, OMP_CLAUSE_WORKER_EXPR (clause),
+ spc, flags, false);
+ pp_right_paren (pp);
+ }
break;
case OMP_CLAUSE_VECTOR:
- pp_string (pp, "vector(");
- dump_generic_node (pp, OMP_CLAUSE_VECTOR_EXPR (clause),
- spc, flags, false);
- pp_character(pp, ')');
+ pp_string (pp, "vector");
+ if (OMP_CLAUSE_VECTOR_EXPR (clause) != NULL_TREE)
+ {
+ pp_left_paren (pp);
+ dump_generic_node (pp, OMP_CLAUSE_VECTOR_EXPR (clause),
+ spc, flags, false);
+ pp_right_paren (pp);
+ }
break;
case OMP_CLAUSE_NUM_GANGS:
@@ -318,9 +318,11 @@ unsigned const char omp_clause_num_ops[] =
2, /* OMP_CLAUSE__CACHE_ */
1, /* OMP_CLAUSE_DEVICE_RESIDENT */
1, /* OMP_CLAUSE_USE_DEVICE */
- 1, /* OMP_CLAUSE_GANG */
+ 2, /* OMP_CLAUSE_GANG */
1, /* OMP_CLAUSE_ASYNC */
1, /* OMP_CLAUSE_WAIT */
+ 0, /* OMP_CLAUSE_AUTO */
+ 0, /* OMP_CLAUSE_SEQ */
1, /* OMP_CLAUSE__LOOPTEMP_ */
1, /* OMP_CLAUSE_IF */
1, /* OMP_CLAUSE_NUM_THREADS */
@@ -378,6 +380,8 @@ const char * const omp_clause_code_name[] =
"gang",
"async",
"wait",
+ "auto",
+ "seq",
"_looptemp_",
"if",
"num_threads",
@@ -11155,9 +11159,12 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,
case OMP_CLAUSE:
switch (OMP_CLAUSE_CODE (*tp))
{
+ case OMP_CLAUSE_GANG:
+ WALK_SUBTREE (OMP_CLAUSE_OPERAND (*tp, 1));
+ /* FALLTHRU */
+
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:
@@ -11201,6 +11208,8 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,
case OMP_CLAUSE_PARALLEL:
case OMP_CLAUSE_SECTIONS:
case OMP_CLAUSE_TASKGROUP:
+ case OMP_CLAUSE_AUTO:
+ case OMP_CLAUSE_SEQ:
WALK_SUBTREE_TAIL (OMP_CLAUSE_CHAIN (*tp));
case OMP_CLAUSE_LASTPRIVATE:
@@ -1156,14 +1156,15 @@ 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 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)
+
/* 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)
@@ -1172,10 +1173,9 @@ DEFTREECODE (OACC_ENTER_DATA, "oacc_enter_data", tcc_statement, 1)
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)
+/* OpenACC - #pragma acc update [clause1 ... clauseN]
+ Operand 0: OACC_UPDATE_CLAUSES: List of clauses. */
+DEFTREECODE (OACC_UPDATE, "oacc_update", tcc_statement, 1)
/* OpenMP - #pragma omp target update [clause1 ... clauseN]
Operand 0: OMP_TARGET_UPDATE_CLAUSES: List of clauses. */
@@ -1231,6 +1231,9 @@ extern void protected_set_expr_location (tree, location_t);
#define OACC_HOST_DATA_CLAUSES(NODE) \
TREE_OPERAND (OACC_HOST_DATA_CHECK (NODE), 1)
+#define OACC_CACHE_CLAUSES(NODE) \
+ TREE_OPERAND (OACC_CACHE_CHECK (NODE), 0)
+
#define OACC_DECLARE_CLAUSES(NODE) \
TREE_OPERAND (OACC_DECLARE_CHECK (NODE), 0)
@@ -1243,9 +1246,6 @@ extern void protected_set_expr_location (tree, location_t);
#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)
@@ -1370,6 +1370,9 @@ extern void protected_set_expr_location (tree, location_t);
#define OMP_CLAUSE_GANG_EXPR(NODE) \
OMP_CLAUSE_OPERAND ( \
OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_GANG), 0)
+#define OMP_CLAUSE_GANG_STATIC_EXPR(NODE) \
+ OMP_CLAUSE_OPERAND ( \
+ OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_GANG), 1)
#define OMP_CLAUSE_ASYNC_EXPR(NODE) \
OMP_CLAUSE_OPERAND ( \
OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_ASYNC), 0)
@@ -771,7 +771,6 @@ __do_global_ctors (void)
#error "What are you doing with crtstuff.c, then?"
#endif
-
#else /* ! CRT_BEGIN && ! CRT_END */
#error "One of CRT_BEGIN or CRT_END must be defined."
#endif
Hi! Committed to gomp-4_0-branch in r218849: commit a6ef9a6d913b6a46872f6f9021d4071e26c6b79c Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4> Date: Wed Dec 17 23:01:26 2014 +0000 Miscellaneous code cleanup and extension. gcc/cp/ * parser.c (cp_parser_oacc_clause_vector_length) (cp_parser_oacc_wait_list, cp_parser_omp_clause_num_gangs) (cp_parser_omp_clause_num_workers): Revise parsing. gcc/c-family/ * c-pragma.h (enum pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_AUTO, PRAGMA_OMP_CLAUSE_GANG, PRAGMA_OMP_CLAUSE_SEQ, PRAGMA_OMP_CLAUSE_VECTOR, and PRAGMA_OMP_CLAUSE_WORKER. gcc/c/ * c-parser.c (c_parser_omp_clause_name): Handle "auto", "gang", "seq", "vector", and "worker". (c_parser_oacc_all_clauses): Handle PRAGMA_OMP_CLAUSE_FIRSTPRIVATE, and PRAGMA_OMP_CLAUSE_PRIVATE. * c-typeck.c (c_finish_omp_clauses): Handle OMP_CLAUSE_AUTO, OMP_CLAUSE_GANG, OMP_CLAUSE_SEQ, OMP_CLAUSE_VECTOR, and OMP_CLAUSE_WORKER. gcc/ * tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_AUTO, and OMP_CLAUSE_SEQ. * tree.c (omp_clause_num_ops, omp_clause_code_name, walk_tree_1): Update accordingly. * gimplify.c (gimplify_scan_omp_clauses) (gimplify_adjust_omp_clauses): Handle OMP_CLAUSE_AUTO, OMP_CLAUSE_GANG, OMP_CLAUSE_SEQ, OMP_CLAUSE_VECTOR, and OMP_CLAUSE_WORKER. * tree-pretty-print.c (dump_omp_clause): Likewise. * omp-low.c (scan_sharing_clauses): Likewise. (struct omp_context): Add gwv_below, and gwv_this members. (enclosing_target_ctx, oacc_loop_or_target_p): New functions. (scan_omp_for, scan_omp_target): Use those. gcc/ * tree-nested.c (convert_nonlocal_reference_stmt) (convert_local_reference_stmt, convert_tramp_reference_stmt) (convert_gimple_call) <GIMPLE_OMP_TARGET>: Use is_gimple_omp_offloaded. gcc/ * omp-low.c (lower_reduction_clauses): Simplify OpenACC handling. gcc/ * gimple.h: Rename is_gimple_omp_oacc_specifically to is_gimple_omp_oacc. Update all users. gcc/ * gimplify.c (enum omp_region_type): Restore ORT_TARGET_DATA and ORT_TARGET. Update all users. gcc/ * doc/generic.texi (OpenMP): Move OpenACC stuff into... (OpenACC): ... this new subsection. gcc/ * oacc-builtins.def: Merge into omp-builtins.def. Update all users. gcc/ * builtin-types.def: Remove BT_FN_VOID_INT_PTR_INT. gcc/fortran/ * types.def: Remove BT_FN_VOID_INT_PTR_INT. Throughout, remove FIXME/TODO/XXX markers, reflect libgomp renaming, remove lots of assertion checks. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@218849 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/ChangeLog.gomp | 37 ++ gcc/Makefile.in | 3 +- gcc/builtin-types.def | 2 - gcc/builtins.def | 14 +- gcc/c-family/ChangeLog.gomp | 7 + gcc/c-family/c-omp.c | 2 +- gcc/c-family/c-pragma.c | 2 +- gcc/c-family/c-pragma.h | 5 + gcc/c/ChangeLog.gomp | 11 + gcc/c/c-parser.c | 32 +- gcc/c/c-typeck.c | 5 + gcc/cp/ChangeLog.gomp | 6 + gcc/cp/parser.c | 55 +-- gcc/doc/generic.texi | 116 ++--- gcc/fortran/ChangeLog.gomp | 2 + gcc/fortran/f95-lang.c | 13 +- gcc/fortran/types.def | 1 - gcc/gimple.c | 4 +- gcc/gimple.def | 7 +- gcc/gimple.h | 15 +- gcc/gimplify.c | 160 +++---- gcc/oacc-builtins.def | 56 --- gcc/omp-builtins.def | 36 +- gcc/omp-low.c | 533 ++++++++++------------ gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c | 4 +- gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90 | 2 +- gcc/tree-core.h | 25 +- gcc/tree-inline.c | 4 +- gcc/tree-nested.c | 22 +- gcc/tree-pretty-print.c | 57 ++- gcc/tree.c | 13 +- gcc/tree.def | 16 +- gcc/tree.h | 9 +- libgcc/crtstuff.c | 1 - 34 files changed, 653 insertions(+), 624 deletions(-) Grüße, Thomas