Hi! The following patch adds additional testing of the declare directive and fixes for issues that arose from the testing.
Committed to gomp-4_0-branch. Jim
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index e7df751..bcbd163 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -1767,12 +1767,15 @@ finish_oacc_declare (tree fnbody, tree decls) break; } - stmt = make_node (OACC_DECLARE); - TREE_TYPE (stmt) = void_type_node; - OACC_DECLARE_CLAUSES (stmt) = ret_clauses; - SET_EXPR_LOCATION (stmt, loc); + if (ret_clauses) + { + stmt = make_node (OACC_DECLARE); + TREE_TYPE (stmt) = void_type_node; + OACC_DECLARE_CLAUSES (stmt) = ret_clauses; + SET_EXPR_LOCATION (stmt, loc); - tsi_link_before (&i, stmt, TSI_CONTINUE_LINKING); + tsi_link_before (&i, stmt, TSI_CONTINUE_LINKING); + } DECL_ATTRIBUTES (fndecl) = remove_attribute ("oacc declare", DECL_ATTRIBUTES (fndecl)); @@ -12812,6 +12815,14 @@ c_parser_oacc_declare (c_parser *parser) error = true; continue; } + else if (TREE_PUBLIC (decl)) + { + error_at (loc, + "invalid use of %<global%> variable %qD " + "in %<#pragma acc declare%>", decl); + error = true; + continue; + } break; } diff --git a/gcc/cp/decl.c b/gcc/cp/decl.c index 15da51e..a35f599 100644 --- a/gcc/cp/decl.c +++ b/gcc/cp/decl.c @@ -14343,7 +14343,17 @@ finish_oacc_declare (tree fndecl, tree decls) { t = tsi_stmt (i); if (TREE_CODE (t) == BIND_EXPR) - list = BIND_EXPR_BODY (t); + { + list = BIND_EXPR_BODY (t); + if (TREE_CODE (list) != STATEMENT_LIST) + { + stmt = list; + list = alloc_stmt_list (); + BIND_EXPR_BODY (t) = list; + i = tsi_start (list); + tsi_link_after (&i, stmt, TSI_CONTINUE_LINKING); + } + } } if (clauses) @@ -14371,11 +14381,11 @@ finish_oacc_declare (tree fndecl, tree decls) } } - if (!found) - { - i = tsi_start (list); - tsi_link_before (&i, stmt, TSI_CONTINUE_LINKING); - } + if (!found) + { + i = tsi_start (list); + tsi_link_before (&i, stmt, TSI_CONTINUE_LINKING); + } } while (oacc_returns) @@ -14405,18 +14415,21 @@ finish_oacc_declare (tree fndecl, tree decls) free (r); } - for (i = tsi_start (list); !tsi_end_p (i); tsi_next (&i)) + if (ret_clauses) { - if (tsi_end_p (i)) - break; - } + for (i = tsi_start (list); !tsi_end_p (i); tsi_next (&i)) + { + if (tsi_end_p (i)) + break; + } - stmt = make_node (OACC_DECLARE); - TREE_TYPE (stmt) = void_type_node; - OMP_STANDALONE_CLAUSES (stmt) = ret_clauses; - SET_EXPR_LOCATION (stmt, loc); + stmt = make_node (OACC_DECLARE); + TREE_TYPE (stmt) = void_type_node; + OMP_STANDALONE_CLAUSES (stmt) = ret_clauses; + SET_EXPR_LOCATION (stmt, loc); - tsi_link_before (&i, stmt, TSI_CONTINUE_LINKING); + tsi_link_before (&i, stmt, TSI_CONTINUE_LINKING); + } DECL_ATTRIBUTES (fndecl) = remove_attribute ("oacc declare", DECL_ATTRIBUTES (fndecl)); diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index 78bcb0a1..41fb35e 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -32123,6 +32123,14 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok) error = true; continue; } + else if (TREE_PUBLIC (decl)) + { + error_at (loc, + "invalid use of %<global%> variable %qD " + "in %<#pragma acc declare%>", decl); + error = true; + continue; + } break; } diff --git a/gcc/testsuite/c-c++-common/goacc/declare-2.c b/gcc/testsuite/c-c++-common/goacc/declare-2.c index ce12463..7979f0c 100644 --- a/gcc/testsuite/c-c++-common/goacc/declare-2.c +++ b/gcc/testsuite/c-c++-common/goacc/declare-2.c @@ -63,4 +63,6 @@ f (void) extern int ve6; #pragma acc declare present_or_create(ve6) /* { dg-error "invalid use of" } */ + +#pragma acc declare present (v9) /* { dg-error "invalid use of" } */ } diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-1.c index 59cfe51..584b921 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-1.c @@ -4,6 +4,26 @@ #include <stdlib.h> #include <stdio.h> +#define N 8 + +void +subr1 (int *a) +{ + int f[N]; +#pragma acc declare copy (f) + +#pragma acc parallel copy (a[0:N]) + { + int i; + + for (i = 0; i < N; i++) + { + f[i] = a[i]; + a[i] = f[i] + f[i]; + } + } +} + int b[8]; #pragma acc declare create (b) @@ -13,7 +33,6 @@ int d[8] = { 1, 2, 3, 4, 5, 6, 7, 8 }; int main (int argc, char **argv) { - const int N = 8; int a[N]; int e[N]; #pragma acc declare create (e) @@ -61,5 +80,18 @@ main (int argc, char **argv) abort (); } + for (i = 0; i < N; i++) + { + a[i] = 1234; + } + + subr1 (&a[0]); + + for (i = 0; i < N; i++) + { + if (a[i] != 1234 * 2) + abort (); + } + return 0; } diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-5.c new file mode 100644 index 0000000..8bb7dec --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-5.c @@ -0,0 +1,16 @@ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ + +#include <openacc.h> +#include <stdlib.h> +#include <stdio.h> + +#define N 8 + +int +main (int argc, char **argv) +{ + int a[N] __attribute__((unused)); +#pragma acc declare present (a) +} + +/* { dg-shouldfail "" } */