Hi,
This patch removes the device-specific filtering (for NVidia PTX) from
the parsing stages of the host compiler (for the device_type clause --
separately for C, C++ and Fortran) in favour of fully parsing the
device_type clauses, but not actually implementing anything for them
(device_type support is a feature that we're not planning to implement
just yet: the existing "support" is something of a red herring).
With this patch, the parsed device_type clauses will be ready at OMP
lowering time whenever we choose to do something with them (e.g.
transforming them into a representation that can be streamed out and
re-read by the appropriate offload compiler). The representation is
more-or-less the same for all supported languages, modulo
clause ordering.
I've altered the dtype-*.* tests to account for the new behaviour (and
to not use e.g. mixed-case "nVidia" or "acc_device_nvidia" names, which
are contrary to the recommendations in the spec).
OK to apply, or any comments?
Thanks,
Julian
ChangeLog
gcc/
* gimplify.c (gimplify_scan_omp_clauses): Handle
OMP_CLAUSE_DEVICE_TYPE.
(gimplify_adjust_omp_clauses): Likewise.
* omp-low.c (scan_sharing_clauses): Likewise.
(expand_omp_target): Add "sorry" for device_type support.
* tree-pretty-print.c (dump_omp_clause): Add device_type support.
* tree.c (walk_tree_1): Likewise.
gcc/c/
* c-parser.c (c_parser_oacc_all_clauses): Don't call
c_oacc_filter_device_types.
* c-typeck.c (c_finish_omp_clauses): Handle OMP_CLAUSE_DEVICE_TYPE.
gcc/cp/
* parser.c (cp_parser_oacc_all_clauses): Don't call
c_oacc_filter_device_types.
* pt.c (tsubst_omp_clauses): Handle OMP_CLAUSE_DEVICE_TYPE.
* semantics.c (finish_omp_clauses): Likewise.
gcc/fortran/
* gfortran.h (gfc_omp_clauses): Change "dtype" int field to
"device_types" gfc_expr_list.
* openmp.c (gfc_match_omp_clauses): Remove scan_dtype variable (add
OMP_CLAUSE_DEVICE_TYPE directly to appropriate bitmasks). Parse all
device_type clauses without filtering.
(OACC_LOOP_CLAUSE_DEVICE_TYPE_MASK)
(OACC_KERNELS_CLAUSE_DEVICE_TYPE_MASK)
(OACC_PARALLEL_CLAUSE_DEVICE_TYPE_MASK)
(OACC_ROUTINE_CLAUSE_DEVICE_TYPE_MASK)
(OACC_UPDATE_CLAUSE_DEVICE_TYPE_MASK): Add OMP_CLAUSE_DEVICE_TYPE.
* trans-openmp.c (gfc_trans_omp_clauses): Translate device_type
clauses, and split old body into...
(gfc_trans_omp_clauses_1): New function.
gcc/testsuite/
* c-c++-common/goacc/dtype-1.c: Update test for new behaviour.
* c-c++-common/goacc/dtype-2.c: Likewise.
* c-c++-common/goacc/dtype-3.c: Likewise.
* c-c++-common/goacc/dtype-4.c: Likewise.
* gfortran.dg/goacc/dtype-1.f95: Likewise.
* gfortran.dg/goacc/dtype-2.f95: Likewise.
* gfortran.dg/goacc/dtype-3.f: Likewise.commit 123298186bb8ce87f84b6a3a72743939d4fdae11
Author: Julian Brown <jul...@codesourcery.com>
Date: Thu Jul 16 08:06:01 2015 -0700
Fix device_type parsing, add sorry() for missing implementation of remainder.
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 1c65abf..d90c18e 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -12439,10 +12439,7 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
c_parser_skip_to_pragma_eol (parser);
if (finish_p)
- {
- clauses = c_oacc_filter_device_types (clauses);
- return c_finish_omp_clauses (clauses, true);
- }
+ return c_finish_omp_clauses (clauses, true);
return clauses;
}
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index 98b8e3d..dcc246c 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -12568,6 +12568,10 @@ c_finish_omp_clauses (tree clauses, bool oacc)
pc = &OMP_CLAUSE_CHAIN (c);
continue;
+ case OMP_CLAUSE_DEVICE_TYPE:
+ pc = &OMP_CLAUSE_DEVICE_TYPE_CLAUSES (c);
+ continue;
+
case OMP_CLAUSE_INBRANCH:
case OMP_CLAUSE_NOTINBRANCH:
if (branch_seen)
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 28f0048..80aabed 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -29879,10 +29879,7 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
cp_parser_skip_to_pragma_eol (parser, pragma_tok);
if (finish_p)
- {
- clauses = c_oacc_filter_device_types (clauses);
- return finish_omp_clauses (clauses, true);
- }
+ return finish_omp_clauses (clauses, true);
return clauses;
}
diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c
index 205dc30..056b2c1 100644
--- a/gcc/cp/pt.c
+++ b/gcc/cp/pt.c
@@ -13666,6 +13666,7 @@ tsubst_omp_clauses (tree clauses, bool declare_simd,
case OMP_CLAUSE_AUTO:
case OMP_CLAUSE_SEQ:
case OMP_CLAUSE_TILE:
+ case OMP_CLAUSE_DEVICE_TYPE:
break;
default:
gcc_unreachable ();
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 8935eb6..1ce1dfa 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -5951,6 +5951,7 @@ finish_omp_clauses (tree clauses, bool oacc)
case OMP_CLAUSE_BIND:
case OMP_CLAUSE_NOHOST:
case OMP_CLAUSE_TILE:
+ case OMP_CLAUSE_DEVICE_TYPE:
break;
case OMP_CLAUSE_INBRANCH:
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index 89f6816..12d46a9 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -1267,7 +1267,7 @@ typedef struct gfc_omp_clauses
struct gfc_expr *num_workers_expr;
struct gfc_expr *vector_length_expr;
struct gfc_symbol *routine_bind;
- int dtype;
+ gfc_expr_list *device_types;
struct gfc_omp_clauses *dtype_clauses;
gfc_expr_list *wait_list;
gfc_expr_list *tile_list;
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index c3d3ccf..4be3417 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -507,7 +507,6 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
{
gfc_omp_clauses *base_clauses, *c = gfc_get_omp_clauses ();
locus old_loc;
- bool scan_dtype = false;
base_clauses = c;
@@ -1154,39 +1153,50 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
if ((mask & OMP_CLAUSE_DEVICE) && c->device == NULL
&& gfc_match ("device ( %e )", &c->device) == MATCH_YES)
continue;
- if (((mask & OMP_CLAUSE_DEVICE_TYPE) || scan_dtype)
+ if ((mask & OMP_CLAUSE_DEVICE_TYPE)
&& (gfc_match ("device_type ( ") == MATCH_YES
|| gfc_match ("dtype ( ") == MATCH_YES))
{
- int device = GOMP_DEVICE_NONE;
gfc_omp_clauses *t = gfc_get_omp_clauses ();
+ gfc_expr_list *p = NULL, *head, *tail;
- c->dtype_clauses = t;
- c = t;
+ head = tail = NULL;
if (gfc_match (" * ") == MATCH_YES)
- device = GOMP_DEVICE_DEFAULT;
+ {
+ head = p = gfc_get_expr_list ();
+ p->expr
+ = gfc_get_character_expr (gfc_default_character_kind,
+ &gfc_current_locus, "*", 1);
+ }
else
{
char n[GFC_MAX_SYMBOL_LEN + 1];
- do {
- if (gfc_match (" %n ", n) == MATCH_YES)
- {
- if (!strcasecmp ("nvidia", n))
- device = GOMP_DEVICE_NVIDIA_PTX;
- else
- {
- /* The OpenACC technical committee advises compilers
- to silently ignore unknown devices. */
- }
- }
- else
- {
- gfc_error ("missing device_type argument");
- continue;
- }
- } while (gfc_match (" , ") == MATCH_YES);
+ do
+ {
+ p = gfc_get_expr_list ();
+
+ if (head == NULL)
+ head = tail = p;
+ else
+ {
+ tail->next = p;
+ tail = p;
+ }
+
+ if (gfc_match (" %n ", n) == MATCH_YES)
+ p->expr
+ = gfc_get_character_expr (gfc_default_character_kind,
+ &gfc_current_locus, n,
+ strlen (n));
+ else
+ {
+ gfc_error ("missing device_type argument");
+ continue;
+ }
+ }
+ while (gfc_match (" , ") == MATCH_YES);
}
/* Consume the trailing ')'. */
@@ -1196,9 +1206,12 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
continue;
}
- c->dtype = device;
+ /* Move to chained pointer for parsing remaining clauses. */
+ c->device_types = head;
+ c->dtype_clauses = t;
+ c = t;
+
mask = dtype_mask;
- scan_dtype = true;
continue;
}
if ((mask & OMP_CLAUSE_THREAD_LIMIT) && c->thread_limit == NULL
@@ -1259,69 +1272,6 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
return MATCH_ERROR;
}
- /* Filter out the device_type clauses. */
- if (base_clauses->dtype_clauses)
- {
- gfc_omp_clauses *t;
- gfc_omp_clauses *seen_default = NULL;
- gfc_omp_clauses *seen_nvidia = NULL;
-
- /* Scan for device_type clauses. */
- c = base_clauses->dtype_clauses;
- while (c)
- {
- if (c->dtype == GOMP_DEVICE_DEFAULT)
- {
- if (seen_default)
- gfc_error ("duplicate device_type (*)");
- else
- seen_default = c;
- }
- else if (c->dtype == GOMP_DEVICE_NVIDIA_PTX)
- {
- if (seen_nvidia)
- gfc_error ("duplicate device_type (nvidia)");
- else
- seen_nvidia = c;
- }
- c = c->dtype_clauses;
- }
-
- /* Update the clauses in the original set of clauses. */
- c = seen_nvidia ? seen_nvidia : seen_default;
- if (c)
- {
-#define acc_clause0(mask) do if (c->mask) { base_clauses->mask = 1; } while (0)
-#define acc_clause1(mask, expr, type) do if (c->mask) { type t; \
- base_clauses->mask = 1; t = base_clauses->expr; \
- base_clauses->expr = c->expr; c->expr = t; } while (0)
-
- acc_clause1 (acc_collapse, collapse, int);
- acc_clause1 (gang, gang_expr, gfc_expr *);
- acc_clause1 (worker, worker_expr, gfc_expr *);
- acc_clause1 (vector, vector_expr, gfc_expr *);
- acc_clause0 (par_auto);
- acc_clause0 (independent);
- acc_clause0 (seq);
- acc_clause1 (tile, tile_list, gfc_expr_list *);
- acc_clause1 (async, async_expr, gfc_expr *);
- acc_clause1 (wait, wait_list, gfc_expr_list *);
- acc_clause1 (num_gangs, num_gangs_expr, gfc_expr *);
- acc_clause1 (num_workers, num_workers_expr, gfc_expr *);
- acc_clause1 (vector_length, vector_length_expr, gfc_expr *);
- acc_clause1 (bind, routine_bind, gfc_symbol *);
- }
-
- /* Remove the device_type clauses. */
- c = base_clauses->dtype_clauses;
- while (c)
- {
- t = c->dtype_clauses;
- gfc_free_omp_clauses (c);
- c = t;
- }
- }
-
*cp = base_clauses;
return MATCH_YES;
}
@@ -1384,17 +1334,18 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
#define OACC_LOOP_CLAUSE_DEVICE_TYPE_MASK \
(OMP_CLAUSE_COLLAPSE | OMP_CLAUSE_GANG | OMP_CLAUSE_WORKER \
- | OMP_CLAUSE_VECTOR | OMP_CLAUSE_AUTO | OMP_CLAUSE_SEQ | OMP_CLAUSE_TILE)
+ | OMP_CLAUSE_VECTOR | OMP_CLAUSE_AUTO | OMP_CLAUSE_SEQ | OMP_CLAUSE_TILE \
+ | OMP_CLAUSE_DEVICE_TYPE)
#define OACC_KERNELS_CLAUSE_DEVICE_TYPE_MASK \
- (OMP_CLAUSE_ASYNC | OMP_CLAUSE_WAIT)
+ (OMP_CLAUSE_ASYNC | OMP_CLAUSE_WAIT | OMP_CLAUSE_DEVICE_TYPE)
#define OACC_PARALLEL_CLAUSE_DEVICE_TYPE_MASK \
(OMP_CLAUSE_ASYNC | OMP_CLAUSE_NUM_GANGS | OMP_CLAUSE_NUM_WORKERS \
- | OMP_CLAUSE_VECTOR_LENGTH | OMP_CLAUSE_WAIT)
+ | OMP_CLAUSE_VECTOR_LENGTH | OMP_CLAUSE_WAIT | OMP_CLAUSE_DEVICE_TYPE)
#define OACC_ROUTINE_CLAUSE_DEVICE_TYPE_MASK \
(OMP_CLAUSE_GANG | OMP_CLAUSE_WORKER | OMP_CLAUSE_VECTOR \
- | OMP_CLAUSE_SEQ | OMP_CLAUSE_BIND)
+ | OMP_CLAUSE_SEQ | OMP_CLAUSE_BIND | OMP_CLAUSE_DEVICE_TYPE)
#define OACC_UPDATE_CLAUSE_DEVICE_TYPE_MASK \
- (OMP_CLAUSE_ASYNC | OMP_CLAUSE_WAIT)
+ (OMP_CLAUSE_ASYNC | OMP_CLAUSE_WAIT | OMP_CLAUSE_DEVICE_TYPE)
match
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index 56e65ec..20a1e65 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -1730,8 +1730,8 @@ gfc_convert_expr_to_tree (stmtblock_t *block, gfc_expr *expr)
}
static tree
-gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
- locus where, bool declare_simd = false)
+gfc_trans_omp_clauses_1 (stmtblock_t *block, gfc_omp_clauses *clauses,
+ locus where, bool declare_simd = false)
{
tree omp_clauses = NULL_TREE, chunk_size, c;
int list;
@@ -2661,6 +2661,45 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
return nreverse (omp_clauses);
}
+static tree
+gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
+ locus where, bool declare_simd = false)
+{
+ tree omp_clauses = gfc_trans_omp_clauses_1 (block, clauses, where,
+ declare_simd);
+
+ for (; clauses->device_types; clauses = clauses->dtype_clauses)
+ {
+ tree c, following_clauses = NULL_TREE, dev_list = NULL_TREE;
+
+ if (clauses->dtype_clauses)
+ {
+ gfc_expr_list *p;
+
+ following_clauses
+ = gfc_trans_omp_clauses_1 (block, clauses->dtype_clauses,
+ where, declare_simd);
+
+ for (p = clauses->device_types; p; p = p->next)
+ {
+ tree dev = gfc_conv_constant_to_tree (p->expr);
+ dev = get_identifier (TREE_STRING_POINTER (dev));
+ if (dev_list)
+ dev_list = chainon (dev_list, dev);
+ else
+ dev_list = dev;
+ }
+
+ c = build_omp_clause (where.lb->location, OMP_CLAUSE_DEVICE_TYPE);
+ OMP_CLAUSE_DEVICE_TYPE_CLAUSES (c) = following_clauses;
+ OMP_CLAUSE_DEVICE_TYPE_DEVICES (c) = dev_list;
+ omp_clauses = gfc_trans_add_clause (c, omp_clauses);
+ }
+ }
+
+ return omp_clauses;
+}
+
/* Like gfc_trans_code, but force creation of a BIND_EXPR around it. */
static tree
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index efae2e5..f5ec04a 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -6616,6 +6616,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
case OMP_CLAUSE_PROC_BIND:
case OMP_CLAUSE_SAFELEN:
case OMP_CLAUSE_TILE:
+ case OMP_CLAUSE_DEVICE_TYPE:
break;
case OMP_CLAUSE_ALIGNED:
@@ -7035,6 +7036,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, tree *list_p)
case OMP_CLAUSE_AUTO:
case OMP_CLAUSE_SEQ:
case OMP_CLAUSE_TILE:
+ case OMP_CLAUSE_DEVICE_TYPE:
break;
default:
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 0419dcd..37b853f 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -2028,6 +2028,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_AUTO:
case OMP_CLAUSE_SEQ:
case OMP_CLAUSE_TILE:
+ case OMP_CLAUSE_DEVICE_TYPE:
break;
case OMP_CLAUSE_ALIGNED:
@@ -2163,6 +2164,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_AUTO:
case OMP_CLAUSE_SEQ:
case OMP_CLAUSE_TILE:
+ case OMP_CLAUSE_DEVICE_TYPE:
break;
case OMP_CLAUSE_DEVICE_RESIDENT:
@@ -9774,6 +9776,10 @@ expand_omp_target (struct omp_region *region)
tree t_async;
int t_wait_idx;
+ c = find_omp_clause (clauses, OMP_CLAUSE_DEVICE_TYPE);
+ if (c)
+ sorry ("device_type clause is not supported yet");
+
/* Default values for t_async. */
t_async = fold_convert_loc (gimple_location (entry_stmt),
integer_type_node,
diff --git a/gcc/testsuite/c-c++-common/goacc/dtype-1.c b/gcc/testsuite/c-c++-common/goacc/dtype-1.c
index e6d6e54..665f21e 100644
--- a/gcc/testsuite/c-c++-common/goacc/dtype-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/dtype-1.c
@@ -1,5 +1,6 @@
/* { dg-do compile } */
/* { dg-options "-fopenacc -fdump-tree-omplower" } */
+/* { dg-prune-output "sorry, unimplemented: device_type clause is not supported yet" } */
void
test ()
@@ -8,7 +9,7 @@ test ()
/* ACC PARALLEL DEVICE_TYPE: */
-#pragma acc parallel device_type (nVidia) async (1) num_gangs (100) num_workers (100) vector_length (32) wait (1)
+#pragma acc parallel device_type (nvidia) async (1) num_gangs (100) num_workers (100) vector_length (32) wait (1)
{
}
@@ -45,49 +46,33 @@ test ()
/* ACC LOOP DEVICE_TYPE: */
#pragma acc parallel
-#pragma acc loop dtype (nVidia) gang tile (1)
+#pragma acc loop dtype (nvidia) gang tile (1)
for (i1 = 1; i1 < 10; i1++)
-#pragma acc loop device_type (nVidia) worker collapse (1)
+#pragma acc loop device_type (nvidia) worker collapse (1)
for (i2 = 1; i2 < 10; i2++)
-#pragma acc loop device_type (nVidia) vector
+#pragma acc loop device_type (nvidia) vector
for (i3 = 1; i3 < 10; i3++)
-#pragma acc loop dtype (nVidia) auto
+#pragma acc loop dtype (nvidia) auto
for (i4 = 1; i4 < 10; i4++)
-#pragma acc loop dtype (nVidia)
+#pragma acc loop dtype (nvidia)
for (i5 = 1; i5 < 10; i5++)
-#pragma acc loop device_type (nVidia) seq
+#pragma acc loop device_type (nvidia) seq
for (i6 = 1; i6 < 10; i6++)
{
}
#pragma acc parallel
-#pragma acc loop device_type (nVidia) gang tile (1) dtype (*) seq
+#pragma acc loop device_type (nvidia) gang tile (1) dtype (*) seq
for (i1 = 1; i1 < 10; i1++)
-#pragma acc loop dtype (nVidia) worker collapse (1) device_type (*) seq
+#pragma acc loop dtype (nvidia) worker collapse (1) device_type (*) seq
for (i2 = 1; i2 < 10; i2++)
-#pragma acc loop device_type (nVidia) vector dtype (*) seq
+#pragma acc loop device_type (nvidia) vector dtype (*) seq
for (i3 = 1; i3 < 10; i3++)
-#pragma acc loop dtype (nVidia) auto device_type (*) seq
+#pragma acc loop dtype (nvidia) auto device_type (*) seq
for (i4 = 1; i4 < 10; i4++)
-#pragma acc loop device_type (nVidia) device_type (*) seq
+#pragma acc loop device_type (nvidia) device_type (*) seq
for (i5 = 1; i5 < 10; i5++)
-#pragma acc loop device_type (nVidia) seq
- for (i6 = 1; i6 < 10; i6++)
- {
- }
-
-#pragma acc parallel
-#pragma acc loop dtype (nVidiaGPU) gang tile (1) device_type (*) seq
- for (i1 = 1; i1 < 10; i1++)
-#pragma acc loop device_type (nVidiaGPU) worker collapse (1) dtype (*) seq
- for (i2 = 1; i2 < 10; i2++)
-#pragma acc loop dtype (nVidiaGPU) vector device_type (*) seq
- for (i3 = 1; i3 < 10; i3++)
-#pragma acc loop device_type (nVidiaGPU) auto device_type (*) seq
- for (i4 = 1; i4 < 10; i4++)
-#pragma acc loop dtype (nVidiaGPU) dtype (*) seq
- for (i5 = 1; i5 < 10; i5++)
-#pragma acc loop device_type (nVidiaGPU) seq device_type (*) seq
+#pragma acc loop device_type (nvidia) seq
for (i6 = 1; i6 < 10; i6++)
{
}
@@ -123,36 +108,36 @@ test ()
#pragma acc routine (foo14) dtype (gpu) seq dtype (*) worker
#pragma acc routine (foo15) dtype (gpu) bind (foo) dtype (*) seq
-/* { dg-final { scan-tree-dump-times "oacc_parallel wait\\(1\\) vector_length\\(32\\) num_workers\\(100\\) num_gangs\\(100\\) async\\(1\\)" 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "oacc_parallel device_type\\(nvidia\\) \\\[ wait\\(1\\) vector_length\\(32\\) num_workers\\(100\\) num_gangs\\(100\\) async\\(1\\) \\\]" 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "oacc_parallel wait\\(1\\) vector_length\\(1\\) num_workers\\(1\\) num_gangs\\(1\\) async\\(1\\) wait\\(2\\) vector_length\\(64\\) num_workers\\(200\\) num_gangs\\(200\\) async\\(2\\)" 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "oacc_parallel device_type\\(nvidia\\) \\\[ wait\\(2\\) vector_length\\(64\\) num_workers\\(200\\) num_gangs\\(200\\) async\\(2\\) \\\] wait\\(1\\) vector_length\\(1\\) num_workers\\(1\\) num_gangs\\(1\\) async\\(1\\)" 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "acc_parallel wait\\(1\\) vector_length\\(1\\) num_workers\\(1\\) num_gangs\\(1\\) async\\(1\\) wait\\(3\\) vector_length\\(128\\) num_workers\\(300\\) num_gangs\\(300\\) async\\(3" 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "oacc_parallel device_type\\(\\*\\) \\\[ wait\\(10\\) vector_length\\(10\\) num_workers\\(10\\) num_gangs\\(10\\) async\\(10\\) \\\] device_type\\(nvidia\\) \\\[ wait\\(3\\) vector_length\\(128\\) num_workers\\(300\\) num_gangs\\(300\\) async\\(3\\) \\\] wait\\(1\\) vector_length\\(1\\) num_workers\\(1\\) num_gangs\\(1\\) async\\(1\\)" 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "oacc_kernels async\\(-1\\)" 4 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "oacc_kernels device_type\\(nvidia\\) \\\[ async\\(-1\\) \\\]" 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "oacc_kernels async\\(-1\\) wait\\(2\\) async\\(2\\)" 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "oacc_kernels device_type\\(nvidia\\) \\\[ wait\\(1\\) async\\(1\\) \\\] async\\(-1\\)" 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "oacc_kernels async\\(-1\\) wait\\(0\\) async\\(0\\)" 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "oacc_kernels device_type\\(\\*\\) \\\[ wait\\(0\\) async\\(0\\) \\\] device_type\\(nvidia\\) \\\[ wait\\(2\\) async\\(2\\) \\\] async\\(-1\\)" 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "acc loop tile\\(1\\) gang private\\(i1\\.0\\) private\\(i1\\)" 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "acc loop device_type\\(nvidia\\) \\\[ tile\\(1\\) gang \\\] private\\(i1\\.0\\) private\\(i1\\)" 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "acc loop tile\\(1\\) gang private\\(i1\\.1\\) private\\(i1\\)" 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "acc loop device_type\\(\\*\\) \\\[ seq \\\] device_type\\(nvidia\\) \\\[ tile\\(1\\) gang \\\] private\\(i1\\.1\\) private\\(i1\\)" 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "acc loop seq private\\(i1\\.2\\) private\\(i1\\)" 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "acc loop device_type\\(nvidia\\) \\\[ collapse\\(1\\) worker \\\] private\\(i2\\)" 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "acc loop collapse\\(1\\) worker private\\(i2\\)" 2 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "acc loop device_type\\(nvidia\\) \\\[ vector \\\] private\\(i3\\)" 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "acc loop vector private\\(i3\\)" 2 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "acc loop device_type\\(nvidia\\) \\\[ auto \\\] private\\(i4\\)" 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "acc loop auto private\\(i4\\)" 2 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "acc loop device_type\\(nvidia\\) \\\[ \\\] private\\(i5\\)" 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "acc loop private\\(i5\\)" 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "acc loop device_type\\(nvidia\\) \\\[ seq \\\] private\\(i6\\)" 2 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "acc loop seq private\\(i6\\)" 3 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "acc loop device_type\\(\\*\\) \\\[ seq \\\] device_type\\(nvidia\\) \\\[ collapse\\(1\\) worker \\\] private\\(i2\\)" 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "acc loop seq private\\(i2\\)" 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "acc loop device_type\\(\\*\\) \\\[ seq \\\] device_type\\(nvidia\\) \\\[ vector \\\] private\\(i3\\)" 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "acc loop seq private\\(i4\\)" 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "acc loop device_type\\(\\*\\) \\\[ seq \\\] device_type\\(nvidia\\) \\\[ auto \\\] private\\(i4\\)" 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "acc loop seq private\\(i5\\)" 2 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "acc loop device_type\\(\\*\\) \\\[ seq \\\] device_type\\(nvidia\\) \\\[ \\\] private\\(i5\\)" 1 "omplower" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/dtype-2.c b/gcc/testsuite/c-c++-common/goacc/dtype-2.c
index 6fa48b9..96acab0 100644
--- a/gcc/testsuite/c-c++-common/goacc/dtype-2.c
+++ b/gcc/testsuite/c-c++-common/goacc/dtype-2.c
@@ -1,4 +1,5 @@
/* { dg-do compile } */
+/* { dg-prune-output "sorry, unimplemented: device_type clause is not supported yet" } */
void
test ()
@@ -7,7 +8,7 @@ test ()
/* ACC PARALLEL DEVICE_TYPE: */
-#pragma acc parallel dtype (nVidia) async (1) num_gangs (100) num_workers (100) vector_length (32) wait (1) copy (i1) /* { dg-error "not valid" } */
+#pragma acc parallel dtype (nvidia) async (1) num_gangs (100) num_workers (100) vector_length (32) wait (1) copy (i1) /* { dg-error "not valid" } */
{
}
@@ -20,7 +21,7 @@ test ()
/* ACC LOOP DEVICE_TYPE: */
#pragma acc parallel
-#pragma acc loop device_type (nVidia) gang tile (1) private (i2) /* { dg-error "not valid" } */
+#pragma acc loop device_type (nvidia) gang tile (1) private (i2) /* { dg-error "not valid" } */
for (i1 = 1; i1 < 10; i1++)
{
}
diff --git a/gcc/testsuite/c-c++-common/goacc/dtype-3.c b/gcc/testsuite/c-c++-common/goacc/dtype-3.c
index 53ab94c..bfb06c4 100644
--- a/gcc/testsuite/c-c++-common/goacc/dtype-3.c
+++ b/gcc/testsuite/c-c++-common/goacc/dtype-3.c
@@ -1,4 +1,5 @@
-/* { dg-do compile } */
+/* { dg-do compile { xfail *-*-* } } */
+/* { dg-prune-output "sorry, unimplemented: device_type clause is not supported yet" } */
float b;
#pragma acc declare link (b)
@@ -14,15 +15,15 @@ main (int argc, char **argv)
{
}
-#pragma acc parallel device_type (acc_device_nvidia) num_gangs (1)
+#pragma acc parallel device_type (nvidia) num_gangs (1)
{
}
-#pragma acc parallel device_type (acc_device_host, acc_device_nvidia) num_gangs (1)
+#pragma acc parallel device_type (host, nvidia) num_gangs (1)
{
}
-#pragma acc parallel device_type (acc_device_host) num_gangs (1)
+#pragma acc parallel device_type (host) num_gangs (1)
{
}
diff --git a/gcc/testsuite/c-c++-common/goacc/dtype-4.c b/gcc/testsuite/c-c++-common/goacc/dtype-4.c
index f49d522..8328783 100644
--- a/gcc/testsuite/c-c++-common/goacc/dtype-4.c
+++ b/gcc/testsuite/c-c++-common/goacc/dtype-4.c
@@ -1,4 +1,4 @@
-/* { dg-do compile } */
+/* { dg-do compile { xfail *-*-* } } */
int
main (int argc, char **argv)
@@ -8,19 +8,19 @@ main (int argc, char **argv)
a = 2.0;
b = 0.0;
- #pragma acc parallel copy (a, b) device_type (acc_device_host) num_gangs (1) device_type (acc_device_nvidia) num_gangs (2)
+ #pragma acc parallel copy (a, b) device_type (host) num_gangs (1) device_type (nvidia) num_gangs (2) /* { dg-message "sorry, unimplemented: device_type clause is not supported yet" } */
{
}
- #pragma acc parallel copy (a, b) num_gangs (3) device_type (acc_device_host) num_gangs (1) device_type (acc_device_nvidia) num_gangs (2)
+ #pragma acc parallel copy (a, b) num_gangs (3) device_type (host) num_gangs (1) device_type (nvidia) num_gangs (2) /* { dg-message "sorry, unimplemented: device_type clause is not supported yet" } */
{
}
-#pragma acc parallel copy (a, b) device_type (acc_device_host) num_gangs (1) device_type (acc_device_nvidia) num_gangs (2) device_type (acc_device_host) num_gangs (60) /* { dg-error "duplicate device_type" } */
+#pragma acc parallel copy (a, b) device_type (host) num_gangs (1) device_type (nvidia) num_gangs (2) device_type (host) num_gangs (60) /* { dg-message "sorry, unimplemented: device_type clause is not supported yet" } */
{
}
-#pragma acc parallel copy (a, b) num_gangs (3) device_type (nvidia) num_gangs (1) device_type (nvidia) num_gangs (2) /* { dg-error "duplicate device_type" } */
+#pragma acc parallel copy (a, b) num_gangs (3) device_type (nvidia) num_gangs (1) device_type (nvidia) num_gangs (2) /* { dg-message "sorry, unimplemented: device_type clause is not supported yet" } */
{
}
diff --git a/gcc/testsuite/gfortran.dg/goacc/dtype-1.f95 b/gcc/testsuite/gfortran.dg/goacc/dtype-1.f95
index 8f68bbe..5919ae4 100644
--- a/gcc/testsuite/gfortran.dg/goacc/dtype-1.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/dtype-1.f95
@@ -1,12 +1,13 @@
! { dg-do compile }
! { dg-options "-fopenacc -fdump-tree-omplower" }
+! { dg-prune-output "sorry, unimplemented: device_type clause is not supported yet" }
program dtype
integer i1, i2, i3, i4, i5, i6
!! ACC PARALLEL DEVICE_TYPE:
-!$acc parallel dtype (nVidia) async (1) num_gangs (100) &
+!$acc parallel dtype (nvidia) async (1) num_gangs (100) &
!$acc& num_workers (100) vector_length (32) wait (1)
!$acc end parallel
@@ -46,17 +47,17 @@ program dtype
!! ACC LOOP DEVICE_TYPE:
!$acc parallel
-!$acc loop device_type (nVidia) gang tile (1)
+!$acc loop device_type (nvidia) gang tile (1)
do i1 = 1, 10
- !$acc loop dtype (nVidia) worker collapse (1)
+ !$acc loop dtype (nvidia) worker collapse (1)
do i2 = 1, 10
- !$acc loop device_type (nVidia) vector
+ !$acc loop device_type (nvidia) vector
do i3 = 1, 10
- !$acc loop device_type (nVidia) auto
+ !$acc loop device_type (nvidia) auto
do i4 = 1, 10
- !$acc loop dtype (nVidia)
+ !$acc loop dtype (nvidia)
do i5 = 1, 10
- !$acc loop dtype (nVidia) seq
+ !$acc loop dtype (nvidia) seq
do i6 = 1, 10
end do
end do
@@ -67,42 +68,19 @@ program dtype
!$acc end parallel
!$acc parallel
-!$acc loop dtype (nVidia) gang tile (1) dtype (*) seq
+!$acc loop dtype (nvidia) gang tile (1) dtype (*) seq
do i1 = 1, 10
- !$acc loop device_type (nVidia) worker collapse (1) &
+ !$acc loop device_type (nvidia) worker collapse (1) &
!$acc& device_type (*) seq
do i2 = 1, 10
- !$acc loop device_type (nVidia) vector dtype (*) seq
+ !$acc loop device_type (nvidia) vector dtype (*) seq
do i3 = 1, 10
- !$acc loop device_type (nVidia) auto dtype (*) seq
+ !$acc loop device_type (nvidia) auto dtype (*) seq
do i4 = 1, 10
- !$acc loop dtype (nVidia) &
+ !$acc loop dtype (nvidia) &
!$acc& dtype (*) seq
do i5 = 1, 10
- !$acc loop device_type (nVidia) seq
- do i6 = 1, 10
- end do
- end do
- end do
- end do
- end do
- end do
-!$acc end parallel
-
-!$acc parallel
-!$acc loop dtype (nVidiaGPU) gang tile (1) dtype (*) seq
- do i1 = 1, 10
- !$acc loop dtype (nVidiaGPU) worker collapse (1) &
- !$acc& device_type (*) seq
- do i2 = 1, 10
- !$acc loop device_type (nVidiaGPU) vector device_type (*) seq
- do i3 = 1, 10
- !$acc loop dtype (nVidiaGPU) auto device_type (*) seq
- do i4 = 1, 10
- !$acc loop dtype (nVidiaGPU) &
- !$acc& dtype (*) seq
- do i5 = 1, 10
- !$acc loop dtype (nVidiaGPU) seq device_type (*) seq
+ !$acc loop device_type (nvidia) seq
do i6 = 1, 10
end do
end do
@@ -189,42 +167,38 @@ subroutine sr5b ()
!$acc routine dtype (gpu) bind (foo) device_type (*) seq
end subroutine sr5b
-! { dg-final { scan-tree-dump-times "oacc_parallel async\\(1\\) wait\\(1\\) num_gangs\\(100\\) num_workers\\(100\\) vector_length\\(32\\)" 1 "omplower" } }
-
-! { dg-final { scan-tree-dump-times "oacc_parallel async\\(2\\) wait\\(2\\) num_gangs\\(200\\) num_workers\\(200\\) vector_length\\(64\\)" 1 "omplower" } }
-
-! { dg-final { scan-tree-dump-times "oacc_parallel async\\(3\\) wait\\(3\\) num_gangs\\(300\\) num_workers\\(300\\) vector_length\\(128\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "oacc_parallel device_type\\(nvidia\\) \\\[ async\\(1\\) wait\\(1\\) num_gangs\\(100\\) num_workers\\(100\\) vector_length\\(32\\) \\\]" 1 "omplower" } }
-! { dg-final { scan-tree-dump-times "oacc_parallel async\\(10\\) wait\\(10\\) num_gangs\\(10\\) num_workers\\(10\\) vector_length\\(10\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "oacc_parallel device_type\\(nvidia\\) \\\[ async\\(2\\) wait\\(2\\) num_gangs\\(200\\) num_workers\\(200\\) vector_length\\(64\\) \\\] async\\(1\\) wait\\(1\\) num_gangs\\(1\\) num_workers\\(1\\) vector_length\\(1\\)" 1 "omplower" } }
-! { dg-final { scan-tree-dump-times "oacc_kernels async\\(-1\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "oacc_parallel device_type\\(\\*\\) \\\[ async\\(10\\) wait\\(10\\) num_gangs\\(10\\) num_workers\\(10\\) vector_length\\(10\\) \\\] device_type\\(nvidia\\) \\\[ async\\(3\\) wait\\(3\\) num_gangs\\(300\\) num_workers\\(300\\) vector_length\\(128\\) \\\] async\\(1\\) wait\\(1\\) num_gangs\\(1\\) num_workers\\(1\\) vector_length\\(1\\)" 1 "omplower" } }
-! { dg-final { scan-tree-dump-times "oacc_kernels async\\(1\\) wait\\(1\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "oacc_parallel device_type\\(\\*\\) \\\[ async\\(10\\) wait\\(10\\) num_gangs\\(10\\) num_workers\\(10\\) vector_length\\(10\\) \\\] device_type\\(nvidia_ptx\\) \\\[ async\\(3\\) wait\\(3\\) num_gangs\\(300\\) num_workers\\(300\\) vector_length\\(128\\) \\\] async\\(1\\) wait\\(1\\) num_gangs\\(1\\) num_workers\\(1\\) vector_length\\(1\\)" 1 "omplower" } }
-! { dg-final { scan-tree-dump-times "oacc_kernels async\\(2\\) wait\\(2\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "oacc_kernels device_type\\(nvidia\\) \\\[ async\\(-1\\) \\\]" 1 "omplower" } }
-! { dg-final { scan-tree-dump-times "oacc_kernels async\\(0\\) wait\\(0\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "oacc_kernels device_type\\(nvidia\\) \\\[ async\\(1\\) wait\\(1\\) \\\]" 1 "omplower" } }
-! { dg-final { scan-tree-dump-times "acc loop private\\(i1\\) tile\\(1\\) gang private\\(i1\\.1\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "oacc_kernels device_type\\(\\*\\) \\\[ async\\(0\\) wait\\(0\\) \\\] device_type\\(nvidia\\) \\\[ async\\(2\\) wait\\(2\\) \\\]" 1 "omplower" } }
-! { dg-final { scan-tree-dump-times "acc loop private\\(i1\\) tile\\(1\\) gang private\\(i1\\.2\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "oacc_kernels device_type\\(\\*\\) \\\[ async\\(0\\) wait\\(0\\) \\\] device_type\\(nvidia_ptx\\) \\\[ async\\(1\\) wait\\(1\\) \\\] async\\(-1\\)" 1 "omplower" } }
-! { dg-final { scan-tree-dump-times "acc loop private\\(i1\\) seq private\\(i1\\.3\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "acc loop device_type\\(nvidia\\) \\\[ tile\\(1\\) gang \\\] private\\(i1\\) private\\(i1\\.1\\)" 1 "omplower" } }
-! { dg-final { scan-tree-dump-times "acc loop private\\(i2\\) collapse\\(1\\) worker" 2 "omplower" } }
+! { dg-final { scan-tree-dump-times "acc loop device_type\\(\\*\\) \\\[ seq \\\] device_type\\(nvidia\\) \\\[ tile\\(1\\) gang \\\] private\\(i1\\) private\\(i1\\.2\\)" 1 "omplower" } }
-! { dg-final { scan-tree-dump-times "acc loop private\\(i3\\) vector" 2 "omplower" } }
+! { dg-final { scan-tree-dump-times "acc loop device_type\\(nvidia\\) \\\[ collapse\\(1\\) worker \\\] private\\(i2\\)" 1 "omplower" } }
-! { dg-final { scan-tree-dump-times "acc loop private\\(i4\\) auto" 2 "omplower" } }
+! { dg-final { scan-tree-dump-times "acc loop device_type\\(nvidia\\) \\\[ vector \\\] private\\(i3\\)" 1 "omplower" } }
-! { dg-final { scan-tree-dump-times "acc loop private\\(i4\\)" 3 "omplower" } }
+! { dg-final { scan-tree-dump-times "acc loop device_type\\(nvidia\\) \\\[ auto \\\] private\\(i4\\)" 1 "omplower" } }
-! { dg-final { scan-tree-dump-times "acc loop private\\(i5\\)" 3 "omplower" } }
+! { dg-final { scan-tree-dump-times "acc loop device_type\\(nvidia\\) \\\[ \\\] private\\(i5\\)" 1 "omplower" } }
-! { dg-final { scan-tree-dump-times "acc loop private\\(i6\\) seq" 3 "omplower" } }
+! { dg-final { scan-tree-dump-times "acc loop device_type\\(nvidia\\) \\\[ seq \\\] private\\(i6\\)" 2 "omplower" } }
-! { dg-final { scan-tree-dump-times "acc loop private\\(i2\\) seq" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "acc loop device_type\\(\\*\\) \\\[ seq \\\] device_type\\(nvidia\\) \\\[ collapse\\(1\\) worker \\\] private\\(i2\\)" 1 "omplower" } }
-! { dg-final { scan-tree-dump-times "acc loop private\\(i4\\) seq" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "acc loop device_type\\(\\*\\) \\\[ seq \\\] device_type\\(nvidia\\) \\\[ auto \\\] private\\(i4\\)" 1 "omplower" } }
-! { dg-final { scan-tree-dump-times "acc loop private\\(i5\\) seq" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "acc loop device_type\\(\\*\\) \\\[ seq \\\] device_type\\(nvidia\\) \\\[ \\\] private\\(i5\\)" 1 "omplower" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/dtype-2.f95 b/gcc/testsuite/gfortran.dg/goacc/dtype-2.f95
index 0d96e37..6405749 100644
--- a/gcc/testsuite/gfortran.dg/goacc/dtype-2.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/dtype-2.f95
@@ -1,11 +1,12 @@
! { dg-do compile }
+! { dg-prune-output "sorry, unimplemented: device_type clause is not supported yet" }
program dtype
integer i1, i2, i3, i4, i5, i6
!! ACC PARALLEL DEVICE_TYPE:
-!$acc parallel device_type (nVidia) async (1) num_gangs (100) &
+!$acc parallel device_type (nvidia) async (1) num_gangs (100) &
!$acc& num_workers (100) vector_length (32) wait (1) copy (i1)
!$acc end parallel
@@ -17,7 +18,7 @@ program dtype
!! ACC LOOP DEVICE_TYPE:
!$acc parallel
-!$acc loop dtype (nVidia) gang tile (1) private (i1)
+!$acc loop dtype (nvidia) gang tile (1) private (i1)
do i1 = 1, 10
end do
!$acc end parallel
@@ -28,12 +29,12 @@ program dtype
end program dtype
-! { dg-error "Unclassifiable OpenACC directive" "" { target *-*-* } 8 }
-! { dg-error "Unexpected" "" { target *-*-* } 10 }
+! { dg-error "Unclassifiable OpenACC directive" "" { target *-*-* } 9 }
+! { dg-error "Unexpected" "" { target *-*-* } 11 }
-! { dg-error "Unclassifiable OpenACC directive" "" { target *-*-* } 14 }
-! { dg-error "Unexpected" "" { target *-*-* } 15 }
+! { dg-error "Unclassifiable OpenACC directive" "" { target *-*-* } 15 }
+! { dg-error "Unexpected" "" { target *-*-* } 16 }
-! { dg-error "Unclassifiable OpenACC directive" "" { target *-*-* } 20 }
+! { dg-error "Unclassifiable OpenACC directive" "" { target *-*-* } 21 }
-! { dg-error "Unclassifiable OpenACC directive" "" { target *-*-* } 27 }
+! { dg-error "Unclassifiable OpenACC directive" "" { target *-*-* } 28 }
diff --git a/gcc/testsuite/gfortran.dg/goacc/dtype-3.f b/gcc/testsuite/gfortran.dg/goacc/dtype-3.f
index 2b2d45f..dd82c78 100644
--- a/gcc/testsuite/gfortran.dg/goacc/dtype-3.f
+++ b/gcc/testsuite/gfortran.dg/goacc/dtype-3.f
@@ -1,4 +1,5 @@
! { dg-do compile }
+! { dg-prune-output "sorry, unimplemented: device_type clause is not supported yet" }
IMPLICIT NONE
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index 4b410a3..00fcee5 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -798,6 +798,15 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, int flags)
spc, flags, false);
pp_right_paren (pp);
break;
+ case OMP_CLAUSE_DEVICE_TYPE:
+ pp_string (pp, "device_type(");
+ dump_generic_node (pp, OMP_CLAUSE_DEVICE_TYPE_DEVICES (clause),
+ spc, flags, false);
+ pp_string (pp, ") [");
+ dump_omp_clauses (pp, OMP_CLAUSE_DEVICE_TYPE_CLAUSES (clause),
+ spc, flags);
+ pp_string (pp, " ]");
+ break;
default:
/* Should never happen. */
diff --git a/gcc/tree.c b/gcc/tree.c
index e67315a..ff533a3 100644
--- a/gcc/tree.c
+++ b/gcc/tree.c
@@ -11401,6 +11401,11 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,
case OMP_CLAUSE_TILE:
WALK_SUBTREE_TAIL (OMP_CLAUSE_CHAIN (*tp));
+ case OMP_CLAUSE_DEVICE_TYPE:
+ WALK_SUBTREE (OMP_CLAUSE_DEVICE_TYPE_DEVICES (*tp));
+ WALK_SUBTREE (OMP_CLAUSE_DEVICE_TYPE_CLAUSES (*tp));
+ WALK_SUBTREE_TAIL (OMP_CLAUSE_CHAIN (*tp));
+
case OMP_CLAUSE_LASTPRIVATE:
WALK_SUBTREE (OMP_CLAUSE_DECL (*tp));
WALK_SUBTREE (OMP_CLAUSE_LASTPRIVATE_STMT (*tp));
