From: Sebastian Galindo <[email protected]>
This patch add OpenACC init,set,shutdown directives support for the C
and C++ front-ends.
An unexpected runtime error was fixed: When the device_type clause had
the argument as *, the set directive was passing an invalid argument to
the runtime library.
The bug 126507 was resolved (dtype alias for device_type).
gcc/c-family/ChangeLog:
* c-common.h (c_finish_oacc_executable_directive): Declare.
* c-omp.cc (oacc_eval_device_type): New.
(c_finish_oacc_executable_directive): New. Lower init, set and
shutdown to GOACC builtins; map set's device_type(*) to
GOMP_DEVICE_DEFAULT.
* c-pragma.cc (oacc_pragmas): Register init, set and shutdown.
* c-pragma.h (enum pragma_kind): Add PRAGMA_OACC_INIT,
PRAGMA_OACC_SET and PRAGMA_OACC_SHUTDOWN.
(enum pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_DEVICE_NUM and
PRAGMA_OACC_CLAUSE_DEVICE_TYPE alias.
gcc/c/ChangeLog:
* c-parser.cc (c_parser_omp_clause_name): Recognize device_num and
dtype as an alias of device_type.
(c_parser_oacc_single_int_clause): Allow device_num 0; clamp
non-positive device_num to 0.
(c_parser_oacc_clause_device_type): New. Parse OpenACC
device_type/dtype arguments.
(c_parser_oacc_all_clauses): Handle DEVICE_NUM and OpenACC
DEVICE_TYPE.
(OACC_INIT_CLAUSE_MASK, OACC_SHUTDOWN_CLAUSE_MASK,
OACC_SET_CLAUSE_MASK): New.
(c_parser_oacc_executable_directive): New.
(c_parser_omp_construct): Handle init, set and shutdown.
* c-typeck.cc (c_finish_omp_clauses): Accept OMP_CLAUSE_DEVICE_NUM.
gcc/cp/ChangeLog:
* parser.cc (cp_parser_omp_clause_name): Recognize device_num and
dtype as an alias of device_type.
(cp_parser_oacc_single_int_clause): Allow device_num 0; clamp
non-positive device_num to 0.
(cp_parser_oacc_clause_device_type): New. Parse OpenACC
device_type/dtype arguments.
(cp_parser_oacc_all_clauses): Handle DEVICE_NUM and OpenACC
DEVICE_TYPE.
(OACC_INIT_CLAUSE_MASK, OACC_SHUTDOWN_CLAUSE_MASK,
OACC_SET_CLAUSE_MASK): New.
(cp_parser_oacc_executable_directive): New.
(cp_parser_omp_construct): Handle init, set and shutdown.
(cp_parser_pragma): Allow init, set and shutdown in statement
context.
* semantics.cc (finish_omp_clauses): Handle OMP_CLAUSE_DEVICE_NUM
like other integral OpenACC clauses, including template
instantiation; require a non-negative value.
gcc/fortran/ChangeLog:
* openmp.cc (gfc_match_omp_clauses): Accept dtype as an alias of
OpenACC device_type.
* trans-openmp.cc (gfc_trans_oacc_executable_directive): For set,
map device_type(*) to GOMP_DEVICE_DEFAULT.
gcc/ChangeLog:
* tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_DEVICE_NUM.
(enum oacc_clause_device_type_kind): New.
(struct tree_omp_clause): Add oacc_device_type_kind.
* tree-pretty-print.cc (dump_omp_clause): Dump DEVICE_NUM and
OpenACC DEVICE_TYPE kinds.
* tree.cc (omp_clause_num_ops, omp_clause_code_name): Handle
DEVICE_NUM.
* tree.h (OMP_CLAUSE_OACC_DEVICE_TYPE_KIND): Define.
gcc/testsuite/ChangeLog:
* c-c++-common/goacc/uninit-if-clause.c: Test uninitialized if on
init.
* c-c++-common/goacc/update-if_present-2.c: Expect errors for
if_present on init and shutdown.
* gfortran.dg/goacc/acc-set-1.f90: Test set with device_type(*).
* c-c++-common/goacc/acc-dtype-1.c: New test.
* c-c++-common/goacc/acc-init-1.c: New test.
* c-c++-common/goacc/acc-init-clauses-1.c: New test.
* c-c++-common/goacc/acc-set-1.c: New test.
* c-c++-common/goacc/acc-set-clauses-1.c: New test.
* c-c++-common/goacc/acc-shutdown-1.c: New test.
* c-c++-common/goacc/acc-shutdown-clauses-1.c: New test.
* gfortran.dg/goacc/acc-dtype-1.f90: New test.
Signed-off-by: Sebastian Galindo <[email protected]>
---
This is a forge pull request published on the gcc-patches mailing list mailing
list
as requested by Sebastian Galindo via Sourceware Forge
<[email protected]>.
Forge discussion: https://forge.sourceware.org/gcc/gcc/pulls/225
Get it locally using:
```
git fetch forge-upstream "+refs/versioned_pull/225/*:refs/versioned_pull/225/*"
git switch -c "pr-225-v1" "refs/versioned_pull/225/1/head"
```
Or, download the patch at: https://forge.sourceware.org/gcc/gcc/pulls/225.diff
Created on: 2026-08-25 14:50:28+00:00
Latest update: 2026-08-27 03:11:30+00:00
Changes: 25 changed files, 939 additions, 26 deletions
Head revision: galindos/gcc-fork-galindos ref
talkysafe/wip/init_set_shutdown/c_cpp commit
4d1ae1abe734ff7cd607af9a1d5b4141b7b906f4
Base revision: gcc/gcc ref trunk commit
b085ae4b49b41e146c0a6a6c16c8c588438b9528 r17-3637-gb085ae4b49b41e
Merge base: b085ae4b49b41e146c0a6a6c16c8c588438b9528
Requested Reviewers:
Changed files:
- A: gcc/testsuite/c-c++-common/goacc/acc-dtype-1.c
- A: gcc/testsuite/c-c++-common/goacc/acc-init-1.c
- A: gcc/testsuite/c-c++-common/goacc/acc-init-clauses-1.c
- A: gcc/testsuite/c-c++-common/goacc/acc-set-1.c
- A: gcc/testsuite/c-c++-common/goacc/acc-set-clauses-1.c
- A: gcc/testsuite/c-c++-common/goacc/acc-shutdown-1.c
- A: gcc/testsuite/c-c++-common/goacc/acc-shutdown-clauses-1.c
- A: gcc/testsuite/gfortran.dg/goacc/acc-dtype-1.f90
- M: gcc/c-family/c-common.h
- M: gcc/c-family/c-omp.cc
- M: gcc/c-family/c-pragma.cc
- M: gcc/c-family/c-pragma.h
- M: gcc/c/c-parser.cc
- M: gcc/c/c-typeck.cc
- M: gcc/cp/parser.cc
- M: gcc/cp/semantics.cc
- M: gcc/fortran/openmp.cc
- M: gcc/fortran/trans-openmp.cc
- M: gcc/testsuite/c-c++-common/goacc/uninit-if-clause.c
- M: gcc/testsuite/c-c++-common/goacc/update-if_present-2.c
- M: gcc/testsuite/gfortran.dg/goacc/acc-set-1.f90
- M: gcc/tree-core.h
- M: gcc/tree-pretty-print.cc
- M: gcc/tree.cc
- M: gcc/tree.h
gcc/c-family/c-common.h | 1 +
gcc/c-family/c-omp.cc | 94 ++++++++++
gcc/c-family/c-pragma.cc | 3 +
gcc/c-family/c-pragma.h | 5 +
gcc/c/c-parser.cc | 163 +++++++++++++++++-
gcc/c/c-typeck.cc | 1 +
gcc/cp/parser.cc | 152 +++++++++++++++-
gcc/cp/semantics.cc | 35 +++-
gcc/fortran/openmp.cc | 8 +-
gcc/fortran/trans-openmp.cc | 7 +-
.../c-c++-common/goacc/acc-dtype-1.c | 40 +++++
gcc/testsuite/c-c++-common/goacc/acc-init-1.c | 62 +++++++
.../c-c++-common/goacc/acc-init-clauses-1.c | 35 ++++
gcc/testsuite/c-c++-common/goacc/acc-set-1.c | 53 ++++++
.../c-c++-common/goacc/acc-set-clauses-1.c | 83 +++++++++
.../c-c++-common/goacc/acc-shutdown-1.c | 62 +++++++
.../goacc/acc-shutdown-clauses-1.c | 53 ++++++
.../c-c++-common/goacc/uninit-if-clause.c | 12 +-
.../c-c++-common/goacc/update-if_present-2.c | 4 +-
.../gfortran.dg/goacc/acc-dtype-1.f90 | 36 ++++
gcc/testsuite/gfortran.dg/goacc/acc-set-1.f90 | 7 +
gcc/tree-core.h | 16 +-
gcc/tree-pretty-print.cc | 25 ++-
gcc/tree.cc | 4 +-
gcc/tree.h | 4 +
25 files changed, 939 insertions(+), 26 deletions(-)
create mode 100644 gcc/testsuite/c-c++-common/goacc/acc-dtype-1.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/acc-init-1.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/acc-init-clauses-1.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/acc-set-1.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/acc-set-clauses-1.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/acc-shutdown-1.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/acc-shutdown-clauses-1.c
create mode 100644 gcc/testsuite/gfortran.dg/goacc/acc-dtype-1.f90
diff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h
index b89a8477e7cef..cc076d3651a47 100644
--- a/gcc/c-family/c-common.h
+++ b/gcc/c-family/c-common.h
@@ -1348,6 +1348,7 @@ extern bool c_omp_check_loop_iv_exprs (location_t, enum
tree_code, tree, int,
tree, tree, tree, walk_tree_lh);
extern bool c_omp_check_loop_binding_exprs (tree, vec<tree> *);
extern tree c_finish_oacc_wait (location_t, tree, tree);
+extern tree c_finish_oacc_executable_directive (location_t, tree, int);
extern tree c_oacc_split_loop_clauses (tree, tree *, bool);
extern void c_omp_split_clauses (location_t, enum tree_code, omp_clause_mask,
tree, tree *);
diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc
index 006e54d1e3d98..81ccfbc560d66 100644
--- a/gcc/c-family/c-omp.cc
+++ b/gcc/c-family/c-omp.cc
@@ -81,6 +81,100 @@ c_finish_oacc_wait (location_t loc, tree parms, tree
clauses)
return stmt;
}
+/* Evals the oacc_clause_device_type_kind into the corresponding
+ device declaration in include/gomp-constants.h. */
+int
+oacc_eval_device_type (oacc_clause_device_type_kind kind)
+{
+ switch (kind)
+ {
+ case OACC_CLAUSE_DEVICE_TYPE_NONE:
+ return GOMP_DEVICE_NONE;
+ case OACC_CLAUSE_DEVICE_TYPE_HOST:
+ return GOMP_DEVICE_HOST;
+ case OACC_CLAUSE_DEVICE_TYPE_NVIDIA:
+ return GOMP_DEVICE_NVIDIA_PTX;
+ case OACC_CLAUSE_DEVICE_TYPE_RADEON:
+ return GOMP_DEVICE_GCN;
+ default:
+ gcc_unreachable ();
+ }
+ return -1;
+}
+
+/* Complete a #pragma acc {init,set,shutdown} contruct.
+ LOC is the location of the #pragma. */
+tree
+c_finish_oacc_executable_directive (location_t loc, tree clauses,
+ int oacc_pragma_kind)
+{
+ tree stmt, arg, dev_num, dev_type;
+ built_in_function builtin_kind;
+
+ switch ((pragma_kind) oacc_pragma_kind)
+ {
+ case PRAGMA_OACC_INIT:
+ builtin_kind = BUILT_IN_GOACC_INIT;
+ break;
+ case PRAGMA_OACC_SHUTDOWN:
+ builtin_kind = BUILT_IN_GOACC_SHUTDOWN;
+ break;
+ case PRAGMA_OACC_SET:
+ builtin_kind = BUILT_IN_GOACC_SET_DEVICE;
+ break;
+ default:
+ gcc_unreachable ();
+ break;
+ }
+
+ stmt = builtin_decl_explicit (builtin_kind);
+
+ arg = omp_find_clause (clauses, OMP_CLAUSE_DEVICE_NUM);
+
+ if (arg)
+ dev_num = OMP_CLAUSE_OPERAND (arg, 0);
+ else
+ /* no 'device_num' clause specified by the user, we don't modify the value
+ of ICV 'acc-current-device-num-var' or we do not take any action in init
+ and shutdown directive using -1 value. */
+ dev_num = build_int_cst (integer_type_node, -1);
+
+ arg = omp_find_clause (clauses, OMP_CLAUSE_DEVICE_TYPE);
+
+ /* GOMP_DEVICE_NONE is used to make the operation
+ in all the devices.
+
+ GOMP_DEVICE_DEFAULT is used in set directive
+ to do nothing if the clause do not appear. */
+ int device_type_int = oacc_pragma_kind == PRAGMA_OACC_SET
+ ? GOMP_DEVICE_DEFAULT
+ : GOMP_DEVICE_NONE;
+
+ if (arg)
+ {
+ device_type_int
+ = oacc_eval_device_type (OMP_CLAUSE_OACC_DEVICE_TYPE_KIND (arg));
+ /* For set, device_type (*) means the default device, not "all". */
+ if (device_type_int == GOMP_DEVICE_NONE
+ && oacc_pragma_kind == PRAGMA_OACC_SET)
+ device_type_int = GOMP_DEVICE_DEFAULT;
+ }
+
+ dev_type = build_int_cst (integer_type_node, device_type_int);
+
+ stmt = build_call_expr_loc (loc, stmt, 2, dev_num, dev_type);
+
+ /* If clause handling. */
+ arg = omp_find_clause (clauses, OMP_CLAUSE_IF);
+ if (arg)
+ stmt = build3_loc (input_location, COND_EXPR, void_type_node,
+ OMP_CLAUSE_IF_EXPR (arg), stmt, NULL_TREE);
+
+ add_stmt (stmt);
+
+ return stmt;
+}
+
/* Complete a #pragma omp master construct. STMT is the structured-block
that follows the pragma. LOC is the location of the #pragma. */
diff --git a/gcc/c-family/c-pragma.cc b/gcc/c-family/c-pragma.cc
index d0d4ff69fd7c9..9b447d0f0e6cf 100644
--- a/gcc/c-family/c-pragma.cc
+++ b/gcc/c-family/c-pragma.cc
@@ -1544,11 +1544,14 @@ static const struct omp_pragma_def oacc_pragmas[] = {
{ "enter", PRAGMA_OACC_ENTER_DATA },
{ "exit", PRAGMA_OACC_EXIT_DATA },
{ "host_data", PRAGMA_OACC_HOST_DATA },
+ { "init", PRAGMA_OACC_INIT },
{ "kernels", PRAGMA_OACC_KERNELS },
{ "loop", PRAGMA_OACC_LOOP },
{ "parallel", PRAGMA_OACC_PARALLEL },
{ "routine", PRAGMA_OACC_ROUTINE },
{ "serial", PRAGMA_OACC_SERIAL },
+ { "set", PRAGMA_OACC_SET },
+ { "shutdown", PRAGMA_OACC_SHUTDOWN },
{ "update", PRAGMA_OACC_UPDATE },
{ "wait", PRAGMA_OACC_WAIT }
};
diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index 5d28cb45c669e..5cd6b5eb641c2 100644
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -34,11 +34,14 @@ enum pragma_kind {
PRAGMA_OACC_ENTER_DATA,
PRAGMA_OACC_EXIT_DATA,
PRAGMA_OACC_HOST_DATA,
+ PRAGMA_OACC_INIT,
PRAGMA_OACC_KERNELS,
PRAGMA_OACC_LOOP,
PRAGMA_OACC_PARALLEL,
PRAGMA_OACC_ROUTINE,
PRAGMA_OACC_SERIAL,
+ PRAGMA_OACC_SET,
+ PRAGMA_OACC_SHUTDOWN,
PRAGMA_OACC_UPDATE,
PRAGMA_OACC_WAIT,
@@ -187,6 +190,7 @@ enum pragma_omp_clause {
PRAGMA_OACC_CLAUSE_CREATE,
PRAGMA_OACC_CLAUSE_DELETE,
PRAGMA_OACC_CLAUSE_DEVICEPTR,
+ PRAGMA_OACC_CLAUSE_DEVICE_NUM,
PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT,
PRAGMA_OACC_CLAUSE_FINALIZE,
PRAGMA_OACC_CLAUSE_GANG,
@@ -208,6 +212,7 @@ enum pragma_omp_clause {
PRAGMA_OACC_CLAUSE_COLLAPSE = PRAGMA_OMP_CLAUSE_COLLAPSE,
PRAGMA_OACC_CLAUSE_COPYIN = PRAGMA_OMP_CLAUSE_COPYIN,
PRAGMA_OACC_CLAUSE_DEVICE = PRAGMA_OMP_CLAUSE_DEVICE,
+ PRAGMA_OACC_CLAUSE_DEVICE_TYPE = PRAGMA_OMP_CLAUSE_DEVICE_TYPE,
PRAGMA_OACC_CLAUSE_DEFAULT = PRAGMA_OMP_CLAUSE_DEFAULT,
PRAGMA_OACC_CLAUSE_DETACH = PRAGMA_OMP_CLAUSE_DETACH,
PRAGMA_OACC_CLAUSE_FIRSTPRIVATE = PRAGMA_OMP_CLAUSE_FIRSTPRIVATE,
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 9dc5ba1e3a13f..6744306d51be0 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -1894,6 +1894,9 @@ static struct c_expr c_parser_has_attribute_expression
(c_parser *);
static void c_parser_oacc_declare (c_parser *);
static void c_parser_oacc_enter_exit_data (c_parser *, bool);
static void c_parser_oacc_update (c_parser *);
+static tree c_parser_oacc_executable_directive (location_t, c_parser *,
+ pragma_kind, omp_clause_mask,
+ const char *);
static void c_parser_omp_construct (c_parser *, bool *);
static void c_parser_omp_groupprivate (c_parser *);
static void c_parser_omp_threadprivate (c_parser *);
@@ -16524,9 +16527,11 @@ c_parser_omp_clause_name (c_parser *parser)
result = PRAGMA_OMP_CLAUSE_DEVICE;
else if (!strcmp ("deviceptr", p))
result = PRAGMA_OACC_CLAUSE_DEVICEPTR;
+ else if (!strcmp ("device_num", p))
+ result = PRAGMA_OACC_CLAUSE_DEVICE_NUM;
else if (!strcmp ("device_resident", p))
result = PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT;
- else if (!strcmp ("device_type", p))
+ else if (!strcmp ("device_type", p) || !strcmp ("dtype", p))
result = PRAGMA_OMP_CLAUSE_DEVICE_TYPE;
else if (!strcmp ("dist_schedule", p))
result = PRAGMA_OMP_CLAUSE_DIST_SCHEDULE;
@@ -18370,7 +18375,8 @@ c_parser_omp_clause_is_device_ptr (c_parser *parser,
tree list)
/* OpenACC:
num_gangs ( expression )
num_workers ( expression )
- vector_length ( expression ) */
+ vector_length (expression)
+ device_num (expression) */
static tree
c_parser_oacc_single_int_clause (c_parser *parser, omp_clause_code code,
@@ -18399,16 +18405,23 @@ c_parser_oacc_single_int_clause (c_parser *parser,
omp_clause_code code,
return list;
}
- /* Attempt to statically determine when the number isn't positive. */
- c = fold_build2_loc (expr_loc, LE_EXPR, boolean_type_node, t,
- build_int_cst (TREE_TYPE (t), 0));
+ bool nonneg_p = code == OMP_CLAUSE_DEVICE_NUM;
+
+ /* Attempt to statically determine when the number isn't positive.
+ device_num accepts 0 as argument. */
+ c = fold_build2_loc (expr_loc,
+ (nonneg_p ? LT_EXPR : LE_EXPR),
+ boolean_type_node, t, build_int_cst (TREE_TYPE (t), 0));
protected_set_expr_location (c, expr_loc);
if (c == boolean_true_node)
{
warning_at (expr_loc, 0,
- "%qs value must be positive",
+ (nonneg_p ? G_ ("%qs value must be non-negative")
+ : G_ ("%qs value must be positive")),
omp_clause_code_name[code]);
t = integer_one_node;
+ if (nonneg_p)
+ t = integer_zero_node;
}
check_no_duplicate_clause (list, code, omp_clause_code_name[code]);
@@ -21315,6 +21328,58 @@ c_parser_omp_clause_device_type (c_parser *parser,
tree list)
return list;
}
+/* OpenACC 2.5:
+ device_type (host | nvidia | radeon | *) */
+
+static tree
+c_parser_oacc_clause_device_type (c_parser *parser, tree list)
+{
+ location_t clause_loc = c_parser_peek_token (parser)->location;
+ enum oacc_clause_device_type_kind kind;
+
+ tree c;
+ matching_parens parens;
+
+ if (!parens.require_open (parser))
+ return c;
+
+ if (c_parser_next_token_is (parser, CPP_NAME))
+ {
+ const char *name
+ = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
+
+ if (!strcmp ("host", name))
+ kind = OACC_CLAUSE_DEVICE_TYPE_HOST;
+ else if (!strcmp ("nvidia", name))
+ kind = OACC_CLAUSE_DEVICE_TYPE_NVIDIA;
+ else if (!strcmp ("radeon", name))
+ kind = OACC_CLAUSE_DEVICE_TYPE_RADEON;
+ else
+ goto invalid_kind;
+ }
+ else if (c_parser_next_token_is (parser, CPP_MULT))
+ {
+ kind = OACC_CLAUSE_DEVICE_TYPE_NONE;
+ }
+ else
+ goto invalid_kind;
+
+ check_no_duplicate_clause (list, OMP_CLAUSE_DEVICE_TYPE, "device_type");
+ c_parser_consume_token (parser);
+
+ parens.skip_until_found_close (parser);
+ c = build_omp_clause (clause_loc, OMP_CLAUSE_DEVICE_TYPE);
+ OMP_CLAUSE_OACC_DEVICE_TYPE_KIND (c) = kind;
+ OMP_CLAUSE_CHAIN (c) = list;
+
+ return c;
+
+invalid_kind:
+ c_parser_error (parser, "expected %<host%>, %<nvidia%>, %<radeon%> or
%<*%>");
+ parens.skip_until_found_close (parser);
+ return list;
+}
+
/* OpenMP 4.0:
from ( variable-list )
to ( variable-list )
@@ -22169,10 +22234,20 @@ 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_OACC_CLAUSE_DEVICE_NUM:
+ clauses
+ = c_parser_oacc_single_int_clause (parser, OMP_CLAUSE_DEVICE_NUM,
+ clauses);
+ c_name = "device_num";
+ break;
case PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT:
clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "device_resident";
break;
+ case PRAGMA_OACC_CLAUSE_DEVICE_TYPE:
+ clauses = c_parser_oacc_clause_device_type (parser, clauses);
+ c_name = "device_type";
+ break;
case PRAGMA_OACC_CLAUSE_FINALIZE:
clauses = c_parser_oacc_simple_clause (here, OMP_CLAUSE_FINALIZE,
clauses);
@@ -23508,6 +23583,61 @@ c_parser_oacc_wait (location_t loc, c_parser *parser,
char *p_name)
return stmt;
}
+/* OpenACC 2.5:
+
+ #pragma acc init oacc-init-clause[optseq] new-line
+*/
+
+#define OACC_INIT_CLAUSE_MASK \
+ ((OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_NUM) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_TYPE))
+
+/* OpenACC 2.5:
+
+ #pragma acc shutdown oacc-shutdown-clause[optseq] new-line
+*/
+
+#define OACC_SHUTDOWN_CLAUSE_MASK \
+ ((OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_NUM) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_TYPE))
+
+
+/* OpenACC 2.5:
+
+ #pragma acc set oacc-set-clause[optseq] new-line
+*/
+
+#define OACC_SET_CLAUSE_MASK \
+ ((OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_NUM) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_TYPE))
+
+
+static tree
+c_parser_oacc_executable_directive (location_t loc, c_parser *parser,
+ pragma_kind oacc_pragma_kind,
+ omp_clause_mask clause_mask,
+ const char *where)
+{
+ tree clauses;
+ clauses = c_parser_oacc_all_clauses (parser, clause_mask, where);
+
+ if (oacc_pragma_kind == PRAGMA_OACC_SET
+ && !omp_find_clause (clauses, OMP_CLAUSE_DEVICE_NUM)
+ && !omp_find_clause (clauses, OMP_CLAUSE_DEVICE_TYPE))
+ {
+ error_at (loc,
+ "%<#pragma acc set%> must contain at least one "
+ "%<device_type%> or %<device_num%> clause");
+ return NULL;
+ }
+
+ return c_finish_oacc_executable_directive (loc, clauses,
+ (int) oacc_pragma_kind);
+}
+
struct c_omp_loc_tree
{
location_t loc;
@@ -31098,6 +31228,27 @@ c_parser_omp_construct (c_parser *parser, bool *if_p)
case PRAGMA_OACC_HOST_DATA:
stmt = c_parser_oacc_host_data (loc, parser, if_p);
break;
+ case PRAGMA_OACC_INIT:
+ strcpy (p_name, "#pragma acc");
+ stmt = c_parser_oacc_executable_directive (loc, parser, PRAGMA_OACC_INIT,
+ OACC_INIT_CLAUSE_MASK,
+ "#pragma acc init");
+ break;
+ case PRAGMA_OACC_SET:
+ strcpy (p_name, "#pragma acc");
+ stmt = c_parser_oacc_executable_directive (loc, parser, PRAGMA_OACC_SET,
+ OACC_SET_CLAUSE_MASK,
+ "#pragma acc set");
+ break;
+
+ case PRAGMA_OACC_SHUTDOWN:
+ strcpy (p_name, "#pragma acc");
+ stmt
+ = c_parser_oacc_executable_directive (loc, parser, PRAGMA_OACC_SHUTDOWN,
+ OACC_SHUTDOWN_CLAUSE_MASK,
+ "#pragma acc shutdown");
+ break;
+
case PRAGMA_OACC_KERNELS:
case PRAGMA_OACC_PARALLEL:
case PRAGMA_OACC_SERIAL:
diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc
index ed234f075ca84..fb5ead434a7d0 100644
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -18164,6 +18164,7 @@ c_finish_omp_clauses (tree clauses, enum
c_omp_region_type ort)
case OMP_CLAUSE_TASKGROUP:
case OMP_CLAUSE_PROC_BIND:
case OMP_CLAUSE_DEVICE_TYPE:
+ case OMP_CLAUSE_DEVICE_NUM:
case OMP_CLAUSE_PRIORITY:
case OMP_CLAUSE_THREADS:
case OMP_CLAUSE_SIMD:
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index d3de58f182feb..63bb78c71667d 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -41495,9 +41495,11 @@ cp_parser_omp_clause_name (cp_parser *parser)
result = PRAGMA_OMP_CLAUSE_DEVICE;
else if (!strcmp ("deviceptr", p))
result = PRAGMA_OACC_CLAUSE_DEVICEPTR;
+ else if (!strcmp ("device_num", p))
+ result = PRAGMA_OACC_CLAUSE_DEVICE_NUM;
else if (!strcmp ("device_resident", p))
result = PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT;
- else if (!strcmp ("device_type", p))
+ else if (!strcmp ("device_type", p) || !strcmp ("dtype", p))
result = PRAGMA_OMP_CLAUSE_DEVICE_TYPE;
else if (!strcmp ("dist_schedule", p))
result = PRAGMA_OMP_CLAUSE_DIST_SCHEDULE;
@@ -42305,7 +42307,8 @@ cp_parser_oacc_simple_clause (location_t loc, enum
omp_clause_code code,
/* OpenACC:
num_gangs ( expression )
num_workers ( expression )
- vector_length ( expression ) */
+ vector_length (expression)
+ device_num (expression) */
static tree
cp_parser_oacc_single_int_clause (cp_parser *parser, omp_clause_code code,
@@ -46274,6 +46277,65 @@ cp_parser_omp_clause_device_type (cp_parser *parser,
tree list,
return list;
}
+/* OpenACC 2.5:
+ device_type (host | nvidia | radeon | *) */
+
+static tree
+cp_parser_oacc_clause_device_type (cp_parser *parser, tree list,
+ location_t clause_loc)
+{
+ enum oacc_clause_device_type_kind kind;
+ tree c;
+
+ if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
+ return list;
+
+ if (cp_lexer_next_token_is (parser->lexer, CPP_NAME))
+ {
+ tree id = cp_lexer_peek_token (parser->lexer)->u.value;
+
+ const char *name = IDENTIFIER_POINTER (id);
+
+ if (!strcmp ("host", name))
+ kind = OACC_CLAUSE_DEVICE_TYPE_HOST;
+ else if (!strcmp ("nvidia", name))
+ kind = OACC_CLAUSE_DEVICE_TYPE_NVIDIA;
+ else if (!strcmp ("radeon", name))
+ kind = OACC_CLAUSE_DEVICE_TYPE_RADEON;
+ else
+ goto invalid_kind;
+ }
+ else if (cp_lexer_next_token_is (parser->lexer, CPP_MULT))
+ {
+ kind = OACC_CLAUSE_DEVICE_TYPE_NONE;
+ }
+ else
+ goto invalid_kind;
+
+ cp_lexer_consume_token (parser->lexer);
+
+ if (!cp_parser_require (parser, CPP_CLOSE_PAREN, RT_COMMA_CLOSE_PAREN))
+ goto resync_fail;
+
+ check_no_duplicate_clause (list, OMP_CLAUSE_DEVICE_TYPE, "device_type",
+ clause_loc);
+ c = build_omp_clause (clause_loc, OMP_CLAUSE_DEVICE_TYPE);
+
+ OMP_CLAUSE_OACC_DEVICE_TYPE_KIND (c) = kind;
+ OMP_CLAUSE_CHAIN (c) = list;
+
+ return c;
+
+invalid_kind:
+ cp_parser_error (parser,
+ "expected %<host%>, %<nvidia%>, %<radeon%> or %<*%>");
+resync_fail:
+ cp_parser_skip_to_closing_parenthesis (parser, /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+ return list;
+}
+
/* OpenMP 5.1:
prefer_type ( const-int-expr-or-string-literal-list )
@@ -46850,6 +46912,16 @@ cp_parser_oacc_all_clauses (cp_parser *parser,
omp_clause_mask mask,
clauses = cp_parser_oacc_data_clause_deviceptr (parser, clauses);
c_name = "deviceptr";
break;
+ case PRAGMA_OACC_CLAUSE_DEVICE_NUM:
+ code = OMP_CLAUSE_DEVICE_NUM;
+ c_name = "device_num";
+ clauses
+ = cp_parser_oacc_single_int_clause (parser, code, c_name, clauses);
+ break;
+ case PRAGMA_OACC_CLAUSE_DEVICE_TYPE:
+ c_name = "device_type";
+ clauses = cp_parser_oacc_clause_device_type (parser, clauses, here);
+ break;
case PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT:
clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "device_resident";
@@ -53157,6 +53229,61 @@ cp_parser_oacc_wait (cp_parser *parser, cp_token
*pragma_tok)
return stmt;
}
+/* OpenACC 2.5:
+
+ #pragma acc init oacc-init-clause[optseq] new-line
+*/
+
+#define OACC_INIT_CLAUSE_MASK \
+ ((OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_NUM) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_TYPE))
+
+/* OpenACC 2.5:
+
+ #pragma acc shutdown oacc-shutdown-clause[optseq] new-line
+*/
+
+#define OACC_SHUTDOWN_CLAUSE_MASK \
+ ((OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_NUM) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_TYPE))
+
+
+/* OpenACC 2.5:
+
+ #pragma acc set oacc-set-clause[optseq] new-line
+*/
+
+#define OACC_SET_CLAUSE_MASK \
+ ((OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_NUM) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_TYPE))
+
+static tree
+cp_parser_oacc_executable_directive (cp_parser *parser,
+ pragma_kind oacc_pragma_kind,
+ omp_clause_mask clause_mask,
+ const char *where, cp_token *pragma_tok)
+{
+ location_t loc = cp_lexer_peek_token (parser->lexer)->location;
+ tree clauses;
+ clauses = cp_parser_oacc_all_clauses (parser, clause_mask, where,
pragma_tok);
+
+ if (oacc_pragma_kind == PRAGMA_OACC_SET
+ && !omp_find_clause (clauses, OMP_CLAUSE_DEVICE_NUM)
+ && !omp_find_clause (clauses, OMP_CLAUSE_DEVICE_TYPE))
+ {
+ error_at (loc,
+ "%<#pragma acc set%> must contain at least one "
+ "%<device_type%> or %<device_num%> clause");
+ return NULL;
+ }
+
+ return c_finish_oacc_executable_directive (loc, clauses,
+ (int) oacc_pragma_kind);
+}
+
/* OpenMP 4.0:
# pragma omp declare simd declare-simd-clauses[optseq] new-line */
@@ -56877,12 +57004,30 @@ cp_parser_omp_construct (cp_parser *parser, cp_token
*pragma_tok, bool *if_p)
case PRAGMA_OACC_HOST_DATA:
stmt = cp_parser_oacc_host_data (parser, pragma_tok, if_p);
break;
+ case PRAGMA_OACC_INIT:
+ stmt
+ = cp_parser_oacc_executable_directive (parser, PRAGMA_OACC_INIT,
+ OACC_INIT_CLAUSE_MASK,
+ "#pragma acc init", pragma_tok);
+ break;
case PRAGMA_OACC_KERNELS:
case PRAGMA_OACC_PARALLEL:
case PRAGMA_OACC_SERIAL:
strcpy (p_name, "#pragma acc");
stmt = cp_parser_oacc_compute (parser, pragma_tok, p_name, if_p);
break;
+ case PRAGMA_OACC_SET:
+ stmt
+ = cp_parser_oacc_executable_directive (parser, PRAGMA_OACC_SET,
+ OACC_SET_CLAUSE_MASK,
+ "#pragma acc set", pragma_tok);
+ break;
+ case PRAGMA_OACC_SHUTDOWN:
+ stmt = cp_parser_oacc_executable_directive (parser, PRAGMA_OACC_SHUTDOWN,
+ OACC_SHUTDOWN_CLAUSE_MASK,
+ "#pragma acc shutdown",
+ pragma_tok);
+ break;
case PRAGMA_OACC_LOOP:
strcpy (p_name, "#pragma acc");
stmt = cp_parser_oacc_loop (parser, pragma_tok, p_name, mask, NULL,
@@ -57616,6 +57761,9 @@ cp_parser_pragma (cp_parser *parser, enum
pragma_context context, bool *if_p)
case PRAGMA_OACC_LOOP:
case PRAGMA_OACC_PARALLEL:
case PRAGMA_OACC_SERIAL:
+ case PRAGMA_OACC_INIT:
+ case PRAGMA_OACC_SET:
+ case PRAGMA_OACC_SHUTDOWN:
case PRAGMA_OMP_ASSUME:
case PRAGMA_OMP_ATOMIC:
case PRAGMA_OMP_CRITICAL:
diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index 7907668da36eb..f8b80fa5577be 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -7969,6 +7969,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type
ort)
{
bool remove = false;
bool field_ok = false;
+ bool non_negative = false;
/* We've reached the end of a list of expanded nodes. Reset the group
start pointer. */
@@ -8529,11 +8530,12 @@ finish_omp_clauses (tree clauses, enum
c_omp_region_type ort)
num_tasks_seen = true;
/* FALLTHRU */
+ case OMP_CLAUSE_DYN_GROUPPRIVATE:
+ case OMP_CLAUSE_DEVICE_NUM:
case OMP_CLAUSE_NUM_TEAMS:
case OMP_CLAUSE_NUM_THREADS:
case OMP_CLAUSE_NUM_GANGS:
case OMP_CLAUSE_NUM_WORKERS:
- case OMP_CLAUSE_DYN_GROUPPRIVATE:
case OMP_CLAUSE_VECTOR_LENGTH:
t = OMP_CLAUSE_OPERAND (c, 0);
if (t == error_mark_node)
@@ -8554,6 +8556,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type
ort)
error_at (OMP_CLAUSE_LOCATION (c),
"%<worker%> num expression must be integral");
break;
+ case OMP_CLAUSE_DEVICE_NUM:
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%<device%> num expression must be integral");
+ break;
default:
error_at (OMP_CLAUSE_LOCATION (c),
"%qs expression must be integral",
@@ -8563,12 +8569,13 @@ finish_omp_clauses (tree clauses, enum
c_omp_region_type ort)
}
else
{
+ non_negative = (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DYN_GROUPPRIVATE
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEVICE_NUM);
t = mark_rvalue_use (t);
if (!processing_template_decl)
{
t = maybe_constant_value (t);
- if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DYN_GROUPPRIVATE
- && TREE_CODE (t) == INTEGER_CST
+ if (!non_negative && TREE_CODE (t) == INTEGER_CST
&& tree_int_cst_sgn (t) != 1)
{
switch (OMP_CLAUSE_CODE (c))
@@ -8597,13 +8604,25 @@ finish_omp_clauses (tree clauses, enum
c_omp_region_type ort)
}
t = integer_one_node;
}
- else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DYN_GROUPPRIVATE
- && TREE_CODE (t) == INTEGER_CST
+ else if (non_negative && TREE_CODE (t) == INTEGER_CST
&& tree_int_cst_sgn (t) < 0)
{
- warning_at (OMP_CLAUSE_LOCATION (c), OPT_Wopenmp,
- "%<dyn_groupprivate%> value must be "
- "non-negative");
+ switch (OMP_CLAUSE_CODE (c))
+ {
+ case OMP_CLAUSE_DEVICE_NUM:
+ warning_at (OMP_CLAUSE_LOCATION (c), 0,
+ "%<device_num%> value must be "
+ "non-negative");
+ break;
+ case OMP_CLAUSE_DYN_GROUPPRIVATE:
+
+ warning_at (OMP_CLAUSE_LOCATION (c), OPT_Wopenmp,
+ "%<dyn_groupprivate%> value must be "
+ "non-negative");
+ break;
+ default:
+ gcc_unreachable ();
+ }
t = integer_zero_node;
}
t = fold_build_cleanup_point_expr (TREE_TYPE (t), t);
diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index e559a2db0db9f..f28d48a8f32a5 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -3226,8 +3226,12 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const
omp_mask mask,
allow_derived))
continue;
if ((mask & OMP_CLAUSE_DEVICE_TYPE) && openacc
- && gfc_match_dupl_check (!c->oacc_device_type_present,
- "device_type", true) == MATCH_YES
+ && (gfc_match_dupl_check (!c->oacc_device_type_present,
+ "device_type", true)
+ == MATCH_YES
+ || gfc_match_dupl_check (!c->oacc_device_type_present,
+ "dtype", true)
+ == MATCH_YES)
&& match_oacc_device_type (c) == MATCH_YES)
continue;
if ((mask & OMP_CLAUSE_DEVICE_TYPE) && !openacc
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index 3e4375f6d6010..930027a2741f8 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -6568,7 +6568,12 @@ builtin_oacc_exec_directive:
GOMP_DEVICE_DEFAULT :
GOMP_DEVICE_NONE;
if (clauses->oacc_device_type_present)
- device_type = clauses->oacc_device_type;
+ {
+ device_type = clauses->oacc_device_type;
+ /* For set, device_type (*) means the default device, not "all". */
+ if (device_type == GOMP_DEVICE_NONE && code->op == EXEC_OACC_SET)
+ device_type = GOMP_DEVICE_DEFAULT;
+ }
tree d_type = build_int_cst (integer_type_node, device_type);
diff --git a/gcc/testsuite/c-c++-common/goacc/acc-dtype-1.c
b/gcc/testsuite/c-c++-common/goacc/acc-dtype-1.c
new file mode 100644
index 0000000000000..75b4229cbe102
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/acc-dtype-1.c
@@ -0,0 +1,40 @@
+/* Test OpenACC dtype abbreviation for device_type on init/set/shutdown. */
+
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-original" } */
+
+void
+dtype_init (void)
+{
+ #pragma acc init dtype(host)
+}
+/* { dg-final { scan-tree-dump-times "__builtin_GOACC_init \\(-1, 2\\)" 1
"original" } } */
+
+void
+dtype_init_star (void)
+{
+ #pragma acc init dtype(*)
+}
+/* { dg-final { scan-tree-dump-times "__builtin_GOACC_init \\(-1, 0\\)" 1
"original" } } */
+
+void
+dtype_set_host (void)
+{
+ #pragma acc set dtype(nvidia)
+}
+/* { dg-final { scan-tree-dump-times "__builtin_GOACC_set_device \\(-1, 5\\)"
1 "original" } } */
+
+void
+dtype_set_star (void)
+{
+ /* dtype(*) on set maps to GOMP_DEVICE_DEFAULT. */
+ #pragma acc set dtype(*)
+}
+/* { dg-final { scan-tree-dump-times "__builtin_GOACC_set_device \\(-1, 1\\)"
1 "original" } } */
+
+void
+dtype_shutdown (void)
+{
+ #pragma acc shutdown dtype(radeon) device_num(1)
+}
+/* { dg-final { scan-tree-dump-times "__builtin_GOACC_shutdown \\(1, 8\\)" 1
"original" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/acc-init-1.c
b/gcc/testsuite/c-c++-common/goacc/acc-init-1.c
new file mode 100644
index 0000000000000..098ba7eda1128
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/acc-init-1.c
@@ -0,0 +1,62 @@
+// Test parsing and lowering of the OpenACC init directive.
+
+// { dg-do compile }
+// { dg-additional-options "-fdump-tree-original" }
+
+void
+init0 (void)
+{
+ #pragma acc init
+}
+// { dg-final { scan-tree-dump-times "__builtin_GOACC_init \\(-1, 0\\)" 2
"original" { target c } } }
+// { dg-final { scan-tree-dump-times "__builtin_GOACC_init \\(-1, 0\\)" 1
"original" { target c++ } } }
+
+
+void
+init1 (void)
+{
+ #pragma acc init device_type(host)
+}
+// { dg-final { scan-tree-dump-times "__builtin_GOACC_init \\(-1, 2\\)" 1
"original" } }
+
+
+void
+init2 (void)
+{
+ #pragma acc init device_type(nvidia)
+}
+// { dg-final { scan-tree-dump-times "__builtin_GOACC_init \\(-1, 5\\)" 1
"original" } }
+
+
+void
+init3 (void)
+{
+ #pragma acc init device_num(0)
+}
+// { dg-final { scan-tree-dump-times "__builtin_GOACC_init \\(0, 0\\)" 1
"original" } }
+
+
+void
+init4 (void)
+{
+ #pragma acc init device_type(host) device_num(0)
+}
+// { dg-final { scan-tree-dump-times "__builtin_GOACC_init \\(0, 2\\)" 1
"original" } }
+
+
+void
+init5 (void)
+{
+ #pragma acc init if(false)
+}
+// { dg-final { scan-tree-dump-times "if \\(0\\)\[\\n\\r\]+ *\{\[\\n\\r\]+
*__builtin_GOACC_init \\(-1, 0\\);" 1 "original" { target c } } }
+// { dg-final { scan-tree-dump-not "if \\(0\\)\[\\n\\r\]+ *\{\[\\n\\r\]+
*__builtin_GOACC_init \\(-1, 0\\);" "original" { target c++ } } }
+
+
+void
+init6 (bool l)
+{
+ #pragma acc init if(l) device_type(radeon) device_num(1)
+}
+// { dg-final { scan-tree-dump-times "if \\(l != 0\\)\[\\n\\r\]+
*\{\[\\n\\r\]+ *__builtin_GOACC_init \\(1, 8\\);" 1 "original" { target c } } }
+// { dg-final { scan-tree-dump-times "if \\(l\\)\[\\n\\r\]+ *\{\[\\n\\r\]+
*__builtin_GOACC_init \\(1, 8\\);" 1 "original" { target c++ } } }
diff --git a/gcc/testsuite/c-c++-common/goacc/acc-init-clauses-1.c
b/gcc/testsuite/c-c++-common/goacc/acc-init-clauses-1.c
new file mode 100644
index 0000000000000..12f327388061f
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/acc-init-clauses-1.c
@@ -0,0 +1,35 @@
+/* Test invalid clauses on the OpenACC init directive. */
+
+/* { dg-do compile } */
+
+void
+bad1 (void)
+{
+#pragma acc init device_type(doesnt_exist) /* { dg-error "expected 'host',
'nvidia', 'radeon' or '\\*'" } */
+}
+
+void
+bad2 (void)
+{
+#pragma acc init device_type(nvidia, host)
+ /* { dg-error "expected '\\\)' before ',' token" "" { target c } .-1 } */
+ /* { dg-error "expected '\\\,' or '\\\)'" "" { target c++ } .-2 } */
+}
+
+void
+bad3 (void)
+{
+#pragma acc init if_present /* { dg-error "'if_present' is not valid" } */
+}
+
+void
+bad4 (void)
+{
+#pragma acc init async /* { dg-error "'async' is not valid" } */
+}
+
+void
+bad5 (void)
+{
+#pragma acc init wait /* { dg-error "'wait' is not valid" } */
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/acc-set-1.c
b/gcc/testsuite/c-c++-common/goacc/acc-set-1.c
new file mode 100644
index 0000000000000..5fc8cf2ea1ae7
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/acc-set-1.c
@@ -0,0 +1,53 @@
+// Test parsing and lowering of the OpenACC set directive.
+
+// { dg-do compile }
+// { dg-additional-options "-fdump-tree-original" }
+
+void
+set0 (void)
+{
+ #pragma acc set device_type(host)
+}
+// { dg-final { scan-tree-dump-times "__builtin_GOACC_set_device \\(-1, 2\\)"
1 "original" } }
+
+
+void
+set1 (void)
+{
+ #pragma acc set device_type(nvidia)
+}
+// { dg-final { scan-tree-dump-times "__builtin_GOACC_set_device \\(-1, 5\\)"
1 "original" } }
+
+
+void
+set2 (void)
+{
+ #pragma acc set device_num(0)
+}
+// { dg-final { scan-tree-dump-times "__builtin_GOACC_set_device \\(0, 1\\)" 1
"original" } }
+
+
+void
+set3 (void)
+{
+ #pragma acc set device_type(host) device_num(0)
+}
+// { dg-final { scan-tree-dump-times "__builtin_GOACC_set_device \\(0, 2\\)" 1
"original" } }
+
+
+void
+set4 (bool l)
+{
+ #pragma acc set if(l) device_type(radeon) device_num(1)
+}
+// { dg-final { scan-tree-dump-times "if \\(l != 0\\)\[\\n\\r\]+
*\{\[\\n\\r\]+ *__builtin_GOACC_set_device \\(1, 8\\);" 1 "original" { target c
} } }
+// { dg-final { scan-tree-dump-times "if \\(l\\)\[\\n\\r\]+ *\{\[\\n\\r\]+
*__builtin_GOACC_set_device \\(1, 8\\);" 1 "original" { target c++ } } }
+
+
+void
+set5 (void)
+{
+ /* device_type(*) maps to GOMP_DEVICE_DEFAULT for set. */
+ #pragma acc set device_type(*)
+}
+// { dg-final { scan-tree-dump-times "__builtin_GOACC_set_device \\(-1, 1\\)"
1 "original" } }
diff --git a/gcc/testsuite/c-c++-common/goacc/acc-set-clauses-1.c
b/gcc/testsuite/c-c++-common/goacc/acc-set-clauses-1.c
new file mode 100644
index 0000000000000..ef92ae53b4421
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/acc-set-clauses-1.c
@@ -0,0 +1,83 @@
+/* Test invalid clauses on the OpenACC set directive. */
+
+/* { dg-do compile } */
+
+void
+bad1 (void)
+{
+#pragma acc set device_type(doesnt_exist)
+ /* { dg-error "expected 'host', 'nvidia', 'radeon' or '\\*'" "" { target
*-*-* } .-1 } */
+ /* { dg-error "'#pragma acc set' must contain at least one 'device_type' or
'device_num' clause" "" { target *-*-* } .-2 } */
+}
+
+void
+bad2 (void)
+{
+#pragma acc set device_type(nvidia, host)
+ /* { dg-error "expected '\\\)' before ',' token" "" { target c } .-1 } */
+ /* { dg-error "expected '\\\,' or '\\\)'" "" { target c++ } .-2 } */
+ /* { dg-error "'#pragma acc set' must contain at least one 'device_type' or
'device_num' clause" "" { target c++ } .-3 } */
+}
+
+void
+bad3 (void)
+{
+#pragma acc set /* { dg-error "'#pragma acc set' must contain at least one
'device_type' or 'device_num' clause" } */
+}
+
+void
+bad4 (void)
+{
+#pragma acc set if(1) /* { dg-error "'#pragma acc set' must contain at least
one 'device_type' or 'device_num' clause" } */
+}
+
+void
+bad5 (void)
+{
+#pragma acc set default_async(0)
+ /* { dg-error "expected an OpenACC clause before 'default_async'" "" {
target *-*-* } .-1 } */
+ /* { dg-error "'#pragma acc set' must contain at least one 'device_type' or
'device_num' clause" "" { target *-*-* } .-2 } */
+}
+
+void
+bad6 (void)
+{
+#pragma acc set if_present
+ /* { dg-error "'if_present' is not valid" "" { target *-*-* } .-1 } */
+ /* { dg-error "'#pragma acc set' must contain at least one 'device_type' or
'device_num' clause" "" { target *-*-* } .-2 } */
+}
+
+void
+bad7 (void)
+{
+#pragma acc set async
+ /* { dg-error "'async' is not valid" "" { target *-*-* } .-1 } */
+ /* { dg-error "'#pragma acc set' must contain at least one 'device_type' or
'device_num' clause" "" { target *-*-* } .-2 } */
+}
+
+void
+bad8 (void)
+{
+#pragma acc set wait
+ /* { dg-error "'wait' is not valid" "" { target *-*-* } .-1 } */
+ /* { dg-error "'#pragma acc set' must contain at least one 'device_type' or
'device_num' clause" "" { target *-*-* } .-2 } */
+}
+
+void
+bad9 (void)
+{
+#pragma acc set device_type(host) device_type(nvidia) /* { dg-error "too many
'device_type' clauses" } */
+}
+
+void
+bad10 (void)
+{
+#pragma acc set device_num(0) device_num(1) /* { dg-error "too many
'device_num' clauses" } */
+}
+
+void
+bad11 (void)
+{
+#pragma acc set if(0) if(1) /* { dg-error "too many 'if' clauses" } */
+ /* { dg-error "'#pragma acc set' must contain at least one 'device_type' or
'device_num' clause" "" { target *-*-* } .-1 } */
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/acc-shutdown-1.c
b/gcc/testsuite/c-c++-common/goacc/acc-shutdown-1.c
new file mode 100644
index 0000000000000..e88a34d0e5b1f
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/acc-shutdown-1.c
@@ -0,0 +1,62 @@
+// Test parsing and lowering of the OpenACC shutdown directive.
+
+// { dg-do compile }
+// { dg-additional-options "-fdump-tree-original" }
+
+void
+shutdown0 (void)
+{
+ #pragma acc shutdown
+}
+// { dg-final { scan-tree-dump-times "__builtin_GOACC_shutdown \\(-1, 0\\)" 2
"original" { target c } } }
+// { dg-final { scan-tree-dump-times "__builtin_GOACC_shutdown \\(-1, 0\\)" 1
"original" { target c++ } } }
+
+
+void
+shutdown1 (void)
+{
+ #pragma acc shutdown device_type(host)
+}
+// { dg-final { scan-tree-dump-times "__builtin_GOACC_shutdown \\(-1, 2\\)" 1
"original" } }
+
+
+void
+shutdown2 (void)
+{
+ #pragma acc shutdown device_type(nvidia)
+}
+// { dg-final { scan-tree-dump-times "__builtin_GOACC_shutdown \\(-1, 5\\)" 1
"original" } }
+
+
+void
+shutdown3 (void)
+{
+ #pragma acc shutdown device_num(0)
+}
+// { dg-final { scan-tree-dump-times "__builtin_GOACC_shutdown \\(0, 0\\)" 1
"original" } }
+
+
+void
+shutdown4 (void)
+{
+ #pragma acc shutdown device_type(host) device_num(0)
+}
+// { dg-final { scan-tree-dump-times "__builtin_GOACC_shutdown \\(0, 2\\)" 1
"original" } }
+
+
+void
+shutdown5 (void)
+{
+ #pragma acc shutdown if(false)
+}
+// { dg-final { scan-tree-dump-times "if \\(0\\)\[\\n\\r\]+ *\{\[\\n\\r\]+
*__builtin_GOACC_shutdown \\(-1, 0\\);" 1 "original" { target c } } }
+// { dg-final { scan-tree-dump-not "if \\(0\\)\[\\n\\r\]+ *\{\[\\n\\r\]+
*__builtin_GOACC_shutdown \\(-1, 0\\);" "original" { target c++ } } }
+
+
+void
+shutdown6 (bool l)
+{
+ #pragma acc shutdown if(l) device_type(radeon) device_num(1)
+}
+// { dg-final { scan-tree-dump-times "if \\(l != 0\\)\[\\n\\r\]+
*\{\[\\n\\r\]+ *__builtin_GOACC_shutdown \\(1, 8\\);" 1 "original" { target c }
} }
+// { dg-final { scan-tree-dump-times "if \\(l\\)\[\\n\\r\]+ *\{\[\\n\\r\]+
*__builtin_GOACC_shutdown \\(1, 8\\);" 1 "original" { target c++ } } }
diff --git a/gcc/testsuite/c-c++-common/goacc/acc-shutdown-clauses-1.c
b/gcc/testsuite/c-c++-common/goacc/acc-shutdown-clauses-1.c
new file mode 100644
index 0000000000000..d1fdfd542cfac
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/acc-shutdown-clauses-1.c
@@ -0,0 +1,53 @@
+/* Test invalid clauses on the OpenACC shutdown directive. */
+
+/* { dg-do compile } */
+
+void
+bad1 (void)
+{
+#pragma acc shutdown device_type(doesnt_exist) /* { dg-error "expected 'host',
'nvidia', 'radeon' or '\\*'" } */
+}
+
+void
+bad2 (void)
+{
+#pragma acc shutdown device_type(nvidia, host)
+ /* { dg-error "expected '\\\)' before ',' token" "" { target c } .-1 } */
+ /* { dg-error "expected '\\\,' or '\\\)'" "" { target c++ } .-2 } */
+}
+
+void
+bad3 (void)
+{
+#pragma acc shutdown if_present /* { dg-error "'if_present' is not valid" } */
+}
+
+void
+bad4 (void)
+{
+#pragma acc shutdown async /* { dg-error "'async' is not valid" } */
+}
+
+void
+bad5 (void)
+{
+#pragma acc shutdown wait /* { dg-error "'wait' is not valid" } */
+}
+
+void
+bad6 (void)
+{
+#pragma acc shutdown device_type(host) device_type(nvidia) /* { dg-error "too
many 'device_type' clauses" } */
+}
+
+void
+bad7 (void)
+{
+#pragma acc shutdown device_num(0) device_num(1) /* { dg-error "too many
'device_num' clauses" } */
+}
+
+void
+bad8 (void)
+{
+#pragma acc shutdown if(0) if(1) /* { dg-error "too many 'if' clauses" } */
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/uninit-if-clause.c
b/gcc/testsuite/c-c++-common/goacc/uninit-if-clause.c
index d4738b9644254..904d868d3ebc1 100644
--- a/gcc/testsuite/c-c++-common/goacc/uninit-if-clause.c
+++ b/gcc/testsuite/c-c++-common/goacc/uninit-if-clause.c
@@ -6,18 +6,20 @@
int
main (void)
{
- int l, l2, ls, l3, l4;
+ int l, l2, ls, l3, l4, l5;
/* { dg-note {'l' was declared here} {} { target *-*-* } .-1 } */
/* { dg-note {'l2' was declared here} {} { target *-*-* } .-2 } */
/* { dg-note {'ls' was declared here} {} { target *-*-* } .-3 } */
/* { dg-note {'l3' was declared here} {} { target *-*-* } .-4 } */
/* { dg-note {'l4' was declared here} {} { target *-*-* } .-5 } */
- bool b, b2, bs, b3, b4;
+ /* { dg-note {'l5' was declared here} {} { target *-*-* } .-6 } */
+ bool b, b2, bs, b3, b4, b5;
/* { dg-note {'b' was declared here} {} { target *-*-* } .-1 } */
/* { dg-note {'b2' was declared here} {} { target *-*-* } .-2 } */
/* { dg-note {'bs' was declared here} {} { target *-*-* } .-3 } */
/* { dg-note {'b3' was declared here} {} { target *-*-* } .-4 } */
/* { dg-note {'b4' was declared here} {} { target *-*-* } .-5 } */
+ /* { dg-note {'b5' was declared here} {} { target *-*-* } .-6 } */
int i, i2;
#pragma acc parallel if(l) /* { dg-warning "is used uninitialized" } */
@@ -50,4 +52,10 @@ main (void)
#pragma acc update if(b4) self(i2) /* { dg-warning "is used uninitialized" }
*/
;
+ #pragma acc init if(l5) /* { dg-warning "is used uninitialized" } */
+ ;
+
+ #pragma acc init if(b5) /* { dg-warning "is used uninitialized" } */
+ ;
+
}
diff --git a/gcc/testsuite/c-c++-common/goacc/update-if_present-2.c
b/gcc/testsuite/c-c++-common/goacc/update-if_present-2.c
index 0fd1505ace8f0..3edc073420c5d 100644
--- a/gcc/testsuite/c-c++-common/goacc/update-if_present-2.c
+++ b/gcc/testsuite/c-c++-common/goacc/update-if_present-2.c
@@ -15,8 +15,8 @@ t1 ()
#pragma acc declare create(c) if_present /* { dg-error "'if_present' is not
valid" } */
-#pragma acc init if_present
-#pragma acc shutdown if_present
+#pragma acc init if_present /* { dg-error "'if_present' is not valid" } */
+#pragma acc shutdown if_present /* { dg-error "'if_present' is not valid" } */
}
void
diff --git a/gcc/testsuite/gfortran.dg/goacc/acc-dtype-1.f90
b/gcc/testsuite/gfortran.dg/goacc/acc-dtype-1.f90
new file mode 100644
index 0000000000000..7688155476fa4
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/acc-dtype-1.f90
@@ -0,0 +1,36 @@
+! Test OpenACC dtype abbreviation for device_type on init/set/shutdown.
+
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-original" }
+
+subroutine dtype_init
+ implicit none
+ !$acc init dtype(host)
+end subroutine dtype_init
+
+subroutine dtype_init_star
+ implicit none
+ !$acc init dtype(*)
+end subroutine dtype_init_star
+
+subroutine dtype_set_host
+ implicit none
+ !$acc set dtype(nvidia)
+end subroutine dtype_set_host
+
+subroutine dtype_set_star
+ implicit none
+ ! dtype(*) on set maps to GOMP_DEVICE_DEFAULT.
+ !$acc set dtype(*)
+end subroutine dtype_set_star
+
+subroutine dtype_shutdown
+ implicit none
+ !$acc shutdown dtype(radeon) device_num(1)
+end subroutine dtype_shutdown
+
+! { dg-final { scan-tree-dump-times "__builtin_GOACC_init \\(-1, 2\\)" 1
"original" } }
+! { dg-final { scan-tree-dump-times "__builtin_GOACC_init \\(-1, 0\\)" 1
"original" } }
+! { dg-final { scan-tree-dump-times "__builtin_GOACC_set_device \\(-1, 5\\)" 1
"original" } }
+! { dg-final { scan-tree-dump-times "__builtin_GOACC_set_device \\(-1, 1\\)" 1
"original" } }
+! { dg-final { scan-tree-dump-times "__builtin_GOACC_shutdown \\(1, 8\\)" 1
"original" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/acc-set-1.f90
b/gcc/testsuite/gfortran.dg/goacc/acc-set-1.f90
index 013c301c775f4..5c9e3e3817d7d 100644
--- a/gcc/testsuite/gfortran.dg/goacc/acc-set-1.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/acc-set-1.f90
@@ -29,8 +29,15 @@ subroutine set4(l)
!$acc set if(l) device_type(radeon) device_num(1)
end subroutine set4
+subroutine set5
+ implicit none
+ ! device_type(*) maps to GOMP_DEVICE_DEFAULT for set.
+ !$acc set device_type(*)
+end subroutine set5
+
! { dg-final { scan-tree-dump-times "__builtin_GOACC_set_device \\(-1, 2\\)" 1
"original" } }
! { dg-final { scan-tree-dump-times "__builtin_GOACC_set_device \\(-1, 5\\)" 1
"original" } }
! { dg-final { scan-tree-dump-times "__builtin_GOACC_set_device \\(0, 1\\)" 1
"original" } }
! { dg-final { scan-tree-dump-times "__builtin_GOACC_set_device \\(0, 2\\)" 1
"original" } }
! { dg-final { scan-tree-dump-times "__builtin_GOACC_set_device \\(1, 8\\)" 1
"original" } }
+! { dg-final { scan-tree-dump-times "__builtin_GOACC_set_device \\(-1, 1\\)" 1
"original" } }
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index 1687b1f946724..150d6f06282d7 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -608,7 +608,10 @@ enum omp_clause_code {
stored in OMP_CLAUSE_MESSAGE_LEN.
Note that the 'severity' clause is handled as flag to 'message':
OMP_CLAUSE_MESSAGE_SEVERITY_WARN. */
- OMP_CLAUSE_MESSAGE
+ OMP_CLAUSE_MESSAGE,
+
+ /* OpenACC clause: device_num (integer-expression). */
+ OMP_CLAUSE_DEVICE_NUM,
};
#undef DEFTREESTRUCT
@@ -1693,6 +1696,16 @@ enum omp_clause_device_type_kind
OMP_CLAUSE_DEVICE_TYPE_ANY = 3
};
+// OpenACC device_type shares code with OpenMP device_type, this numeration
+// ensure propper debug output.
+enum oacc_clause_device_type_kind
+{
+ OACC_CLAUSE_DEVICE_TYPE_NONE = 4,
+ OACC_CLAUSE_DEVICE_TYPE_HOST = 5,
+ OACC_CLAUSE_DEVICE_TYPE_NVIDIA = 6,
+ OACC_CLAUSE_DEVICE_TYPE_RADEON = 7,
+};
+
enum omp_clause_linear_kind
{
OMP_CLAUSE_LINEAR_DEFAULT,
@@ -1773,6 +1786,7 @@ struct GTY(()) tree_omp_clause {
enum omp_clause_bind_kind bind_kind;
enum omp_clause_device_type_kind device_type_kind;
enum omp_clause_fallback_kind fallback_kind;
+ enum oacc_clause_device_type_kind oacc_device_type_kind;
} GTY ((skip)) subcode;
/* The gimplification of OMP_CLAUSE_REDUCTION_{INIT,MERGE} for omp-low's
diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc
index 316945c00ede3..0759fa456aa75 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -1377,6 +1377,12 @@ dump_omp_clause (pretty_printer *pp, tree clause, int
spc, dump_flags_t flags)
pp_right_paren (pp);
break;
+ case OMP_CLAUSE_DEVICE_NUM:
+ pp_string (pp, "device_num (");
+ dump_generic_node (pp, OMP_CLAUSE_OPERAND (clause, 0), spc, flags,
false);
+ pp_character (pp, ')');
+ break;
+
case OMP_CLAUSE_DEVICE_TYPE:
pp_string (pp, "device_type(");
switch (OMP_CLAUSE_DEVICE_TYPE_KIND (clause))
@@ -1391,7 +1397,24 @@ dump_omp_clause (pretty_printer *pp, tree clause, int
spc, dump_flags_t flags)
pp_string (pp, "any");
break;
default:
- gcc_unreachable ();
+ // In this case, should be the OpenACC variant.
+ switch (OMP_CLAUSE_OACC_DEVICE_TYPE_KIND (clause))
+ {
+ case OACC_CLAUSE_DEVICE_TYPE_NONE:
+ pp_string (pp, "*");
+ break;
+ case OACC_CLAUSE_DEVICE_TYPE_HOST:
+ pp_string (pp, "host");
+ break;
+ case OACC_CLAUSE_DEVICE_TYPE_NVIDIA:
+ pp_string (pp, "nvidia");
+ break;
+ case OACC_CLAUSE_DEVICE_TYPE_RADEON:
+ pp_string (pp, "radeon");
+ break;
+ default:
+ gcc_unreachable ();
+ }
}
pp_right_paren (pp);
break;
diff --git a/gcc/tree.cc b/gcc/tree.cc
index 69a7c7e6d34d4..1d43c3d35d5f5 100644
--- a/gcc/tree.cc
+++ b/gcc/tree.cc
@@ -401,6 +401,7 @@ unsigned const char omp_clause_num_ops[] =
1, /* OMP_CLAUSE_DYN_GROUPPRIVATE */
3, /* OMP_CLAUSE_USES_ALLOCATORS */
2, /* OMP_CLAUSE_MESSAGE */
+ 1, /* OMP_CLAUSE_DEVICE_NUM. */
};
const char * const omp_clause_code_name[] =
@@ -506,7 +507,8 @@ const char * const omp_clause_code_name[] =
"nocontext",
"dyn_groupprivate",
"uses_allocators",
- "message"
+ "message",
+ "device_num",
};
/* Unless specific to OpenACC, we tend to internally maintain OpenMP-centric
diff --git a/gcc/tree.h b/gcc/tree.h
index 3fb643ea0fb48..06c3287ccc506 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -2015,6 +2015,10 @@ class auto_suppress_location_wrappers
#define OMP_CLAUSE_DEVICE_TYPE_KIND(NODE) \
(OMP_CLAUSE_SUBCODE_CHECK (NODE,
OMP_CLAUSE_DEVICE_TYPE)->omp_clause.subcode.device_type_kind)
+#define OMP_CLAUSE_OACC_DEVICE_TYPE_KIND(NODE) \
+ (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_DEVICE_TYPE) \
+ ->omp_clause.subcode.oacc_device_type_kind)
+
#define OMP_CLAUSE_INDIRECT_EXPR(NODE) \
OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_INDIRECT), 0)
--
2.54.0
