This patch teaches all of the front ends how to error on duplicate
device_type entries such as
#pragma acc parallel device_type (nvidia) num_gangs (1) \
device_type (nvidia) num_gangs (20)
Before such clauses were silently ignored. I also fixed how invalid
device_type claues are reported in the fortran front end. It used report
"Invalid character" and now it states "Unclassifiable OpenACC directive".
I applied this patch to gomp-4_0-branch.
Cesar
2015-06-12 Cesar Philippidis <ce...@codesourcery.com>
gcc/c-family/
* c-omp.c (oacc_extract_device_id): Recognize GOMP_DEVICE_DEFAULT.
(struct identifier_hasher): New struct declaration.
(oacc_filter_device_types): Report errors on duplicate device_type
entries.
gcc/c/
* c-parser.c (c_parser_oacc_clause_device_type): Switch
OMP_CLAUSE_DEVICE_TYPE_DEVICES to tree instead of an int.
gcc/cp/
* parser.c (cp_parser_oacc_clause_device_type): Likewise.
gcc/fortran/
* openmp.c (gfc_match_omp_clauses): Report errors on duplicate
device_type entries.
gcc/testsuite/
* c-c++-common/goacc/dtype-4.c: New test.
* gfortran.dg/goacc/dtype-2.f95: Update error messages.
* gfortran.dg/goacc/dtype-3.f: New test.
diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c
index c30e0d8..bcb6ff4 100644
--- a/gcc/c-family/c-omp.c
+++ b/gcc/c-family/c-omp.c
@@ -41,6 +41,7 @@ along with GCC; see the file COPYING3. If not see
#include "langhooks.h"
#include "omp-low.h"
#include "gomp-constants.h"
+#include "tree-hasher.h"
/* Complete a #pragma oacc wait construct. LOC is the location of
@@ -1097,9 +1098,20 @@ oacc_extract_device_id (const char *device)
{
if (!strcasecmp (device, "nvidia"))
return GOMP_DEVICE_NVIDIA_PTX;
+ else if (!strcmp (device, "*"))
+ return GOMP_DEVICE_DEFAULT;
return GOMP_DEVICE_NONE;
}
+struct identifier_hasher : ggc_cache_hasher<tree>
+{
+ static hashval_t hash (tree t) { return htab_hash_pointer (t); }
+ static bool equal (tree a, tree b)
+ {
+ return !strcmp(IDENTIFIER_POINTER (a), IDENTIFIER_POINTER (b));
+ }
+};
+
/* Filter out the list of unsupported OpenACC device_types. */
tree
@@ -1109,56 +1121,55 @@ oacc_filter_device_types (tree clauses)
tree dtype = NULL_TREE;
tree seen_nvidia = NULL_TREE;
tree seen_default = NULL_TREE;
- int device = 0;
+ hash_table<identifier_hasher> *dt_htab
+ = hash_table<identifier_hasher>::create_ggc (10);
/* First scan for all device_type clauses. */
for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
{
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEVICE_TYPE)
{
- int code = TREE_INT_CST_LOW (OMP_CLAUSE_DEVICE_TYPE_DEVICES (c));
+ tree t;
- if (code == GOMP_DEVICE_DEFAULT)
+ for (t = OMP_CLAUSE_DEVICE_TYPE_DEVICES (c); t; t = TREE_CHAIN (t))
{
- if (device & (1 << GOMP_DEVICE_DEFAULT))
+ if (dt_htab->find (t))
{
- seen_default = NULL_TREE;
error_at (OMP_CLAUSE_LOCATION (c),
- "duplicate device_type (*)");
- goto filter_error;
+ "duplicate device_type (%s)",
+ IDENTIFIER_POINTER (t));
+ goto filter_dtype;
}
- seen_default = OMP_CLAUSE_DEVICE_TYPE_CLAUSES (c);
- }
- else if (code & (1 << GOMP_DEVICE_NVIDIA_PTX))
- {
- if (device & code)
- {
- seen_nvidia = NULL_TREE;
- error_at (OMP_CLAUSE_LOCATION (c),
- "duplicate device_type (nvidia)");
- goto filter_error;
- }
+ int code = oacc_extract_device_id (IDENTIFIER_POINTER (t));
- seen_nvidia = OMP_CLAUSE_DEVICE_TYPE_CLAUSES (c);
- }
- else
- {
- if (device & (1 << code))
+ if (code == GOMP_DEVICE_DEFAULT)
+ seen_default = OMP_CLAUSE_DEVICE_TYPE_CLAUSES (c);
+ else if (code == GOMP_DEVICE_NVIDIA_PTX)
+ seen_nvidia = OMP_CLAUSE_DEVICE_TYPE_CLAUSES (c);
+ else
{
- error_at (OMP_CLAUSE_LOCATION (c),
- "duplicate device_type");
- goto filter_error;
+ /* The OpenACC technical committee advises compilers
+ to silently ignore unknown devices. */
}
+
+ tree *slot = dt_htab->find_slot (t, INSERT);
+ *slot = t;
}
- device |= (1 << code);
}
}
/* Don't do anything if there aren't any device_type clauses. */
- if (device == 0)
+ if (dt_htab->elements () == 0)
return clauses;
+ if (seen_nvidia)
+ dtype = seen_nvidia;
+ else if (seen_default)
+ dtype = seen_default;
+ else
+ goto filter_dtype;
+
dtype = seen_nvidia ? seen_nvidia : seen_default;
/* Now filter out clauses if necessary. */
@@ -1186,7 +1197,7 @@ oacc_filter_device_types (tree clauses)
prev = c;
}
- filter_error:
+ filter_dtype:
/* Remove all device_type clauses. Those clauses are located at the
beginning of the clause list. */
for (c = clauses; c && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEVICE_TYPE;
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index cef56dc..f37a8f7 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -11246,7 +11246,7 @@ c_parser_oacc_clause_device_type (c_parser *parser, omp_clause_mask mask,
{
tree c, clauses;
location_t loc;
- int dev_id = GOMP_DEVICE_NONE;
+ tree dev_id = NULL_TREE;
loc = c_parser_peek_token (parser)->location;
if (!c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>"))
@@ -11255,7 +11255,7 @@ c_parser_oacc_clause_device_type (c_parser *parser, omp_clause_mask mask,
if (c_parser_next_token_is (parser, CPP_MULT))
{
c_parser_consume_token (parser);
- dev_id = GOMP_DEVICE_DEFAULT;
+ dev_id = get_identifier ("*");
if (!c_parser_require (parser, CPP_CLOSE_PAREN, "expected %<)%>"))
return list;
}
@@ -11264,7 +11264,6 @@ c_parser_oacc_clause_device_type (c_parser *parser, omp_clause_mask mask,
do
{
tree keyword = error_mark_node;
- int dev = 0;
if (c_parser_next_token_is (parser, CPP_NAME))
{
@@ -11280,9 +11279,10 @@ c_parser_oacc_clause_device_type (c_parser *parser, omp_clause_mask mask,
return list;
}
- dev = oacc_extract_device_id (IDENTIFIER_POINTER (keyword));
- if (dev)
- dev_id |= 1 << dev;
+ if (dev_id)
+ dev_id = chainon (dev_id, keyword);
+ else
+ dev_id = keyword;
if (c_parser_next_token_is (parser, CPP_COMMA))
c_parser_consume_token (parser);
@@ -11297,8 +11297,7 @@ c_parser_oacc_clause_device_type (c_parser *parser, omp_clause_mask mask,
clauses = c_parser_oacc_all_clauses (parser, mask, "device_type", 0, false,
false);
OMP_CLAUSE_DEVICE_TYPE_CLAUSES (c) = clauses;
- OMP_CLAUSE_DEVICE_TYPE_DEVICES (c) = build_int_cst (integer_type_node,
- dev_id);
+ OMP_CLAUSE_DEVICE_TYPE_DEVICES (c) = dev_id;
OMP_CLAUSE_CHAIN (c) = list;
return c;
}
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index ca4bb68..60bc287 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -28356,7 +28356,7 @@ cp_parser_oacc_clause_device_type (cp_parser *parser, omp_clause_mask mask,
{
tree c, clauses;
location_t loc;
- int dev_id = GOMP_DEVICE_NONE;
+ tree dev_id = NULL_TREE;
loc = cp_lexer_peek_token (parser->lexer)->location;
if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
@@ -28365,7 +28365,7 @@ cp_parser_oacc_clause_device_type (cp_parser *parser, omp_clause_mask mask,
if (cp_lexer_next_token_is (parser->lexer, CPP_MULT))
{
cp_lexer_consume_token (parser->lexer);
- dev_id = GOMP_DEVICE_DEFAULT;
+ dev_id = get_identifier ("*");
if (!cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN))
return list;
}
@@ -28374,7 +28374,6 @@ cp_parser_oacc_clause_device_type (cp_parser *parser, omp_clause_mask mask,
do
{
tree keyword = error_mark_node;
- int dev = 0;
if (cp_lexer_next_token_is (parser->lexer, CPP_NAME))
{
@@ -28390,9 +28389,10 @@ cp_parser_oacc_clause_device_type (cp_parser *parser, omp_clause_mask mask,
return list;
}
- dev = oacc_extract_device_id (IDENTIFIER_POINTER (keyword));
- if (dev)
- dev_id |= 1 << dev;
+ if (dev_id)
+ dev_id = chainon (dev_id, keyword);
+ else
+ dev_id = keyword;
if (cp_lexer_next_token_is (parser->lexer, CPP_COMMA))
cp_lexer_consume_token (parser->lexer);
@@ -28407,8 +28407,7 @@ cp_parser_oacc_clause_device_type (cp_parser *parser, omp_clause_mask mask,
clauses = cp_parser_oacc_all_clauses (parser, mask, "device_type",
pragma_tok, 0, false, false);
OMP_CLAUSE_DEVICE_TYPE_CLAUSES (c) = clauses;
- OMP_CLAUSE_DEVICE_TYPE_DEVICES (c) = build_int_cst (integer_type_node,
- dev_id);
+ OMP_CLAUSE_DEVICE_TYPE_DEVICES (c) = dev_id;
OMP_CLAUSE_CHAIN (c) = list;
return c;
}
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 46bf865..b98a933 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -1171,17 +1171,23 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
{
char n[GFC_MAX_SYMBOL_LEN + 1];
- while (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. */
- }
- gfc_match (" , ");
- }
+ 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);
}
/* Consume the trailing ')'. */
diff --git a/gcc/testsuite/c-c++-common/goacc/dtype-4.c b/gcc/testsuite/c-c++-common/goacc/dtype-4.c
new file mode 100644
index 0000000..f49d522
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/dtype-4.c
@@ -0,0 +1,28 @@
+/* { dg-do compile } */
+
+int
+main (int argc, char **argv)
+{
+ float a, b;
+
+ 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) 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) 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) num_gangs (3) device_type (nvidia) num_gangs (1) device_type (nvidia) num_gangs (2) /* { dg-error "duplicate device_type" } */
+ {
+ }
+
+ return 0;
+}
diff --git a/gcc/testsuite/gfortran.dg/goacc/dtype-2.f95 b/gcc/testsuite/gfortran.dg/goacc/dtype-2.f95
index a4573e9..0d96e37 100644
--- a/gcc/testsuite/gfortran.dg/goacc/dtype-2.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/dtype-2.f95
@@ -28,12 +28,12 @@ program dtype
end program dtype
-! { dg-error "Invalid character" "" { target *-*-* } 8 }
+! { dg-error "Unclassifiable OpenACC directive" "" { target *-*-* } 8 }
! { dg-error "Unexpected" "" { target *-*-* } 10 }
-! { dg-error "Invalid character" "" { target *-*-* } 14 }
+! { dg-error "Unclassifiable OpenACC directive" "" { target *-*-* } 14 }
! { dg-error "Unexpected" "" { target *-*-* } 15 }
-! { dg-error "Invalid character" "" { target *-*-* } 20 }
+! { dg-error "Unclassifiable OpenACC directive" "" { target *-*-* } 20 }
-! { dg-error "Invalid character" "" { target *-*-* } 27 }
+! { dg-error "Unclassifiable OpenACC directive" "" { target *-*-* } 27 }
diff --git a/gcc/testsuite/gfortran.dg/goacc/dtype-3.f b/gcc/testsuite/gfortran.dg/goacc/dtype-3.f
new file mode 100644
index 0000000..2b2d45f
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/dtype-3.f
@@ -0,0 +1,11 @@
+! { dg-do compile }
+
+ IMPLICIT NONE
+
+ INTEGER X
+
+!$ACC PARALLEL DTYPE (NVIDIA) NUM_GANGS(10) COPYOUT(X) ! { dg-error "Unclassifiable OpenACC directive" }
+ X = 0
+!$ACC END PARALLEL ! { dg-error "Unexpected" }
+
+ END
