Hi Jim!
On Tue, 4 Nov 2014 14:18:43 -0600, James Norris <james_nor...@mentor.com> wrote:
> void
> f (int flag,, int *a)
> {
> if (flag)
> #pragma acc update host (a[0:10])
>
> return;
> }
>
> I read bullet one under "Restrictions" on page 43, to mean the
> placement of the update directive statement in the above example
> is wrong. Given the poor grammar, this could be read a different way.
That's also my interpretation, and I think it should be that way for all
executable directives (as listed 2.12 Executable Directives, plus 2.8
Cache Directive -- am I missing additional ones?), for the reason we just
talked about: if -fopenacc is not active, this code will turn into:
if (flag)
return;
..., which probably was not intended. I think this is spelled out
clearly in the OpenMP specification (but didn't verify right now), and
see also the discussion in
<https://gcc.gnu.org/bugzilla/show_bug.cgi?id=63326>.
As to fixing this, see how I'm handling PRAGMA_OACC_UPDATE in
gcc/c/c-parser:c_parser_pragma; modelled after PRAGMA_OMP_* executable
directives. Plus, I started testing this in
gcc/testsuite/c-c++-common/goacc/pragma_context.c. We should also be
adding appropritate documentation to the C/C++ parser source code (or
generally GCC user documentation?) -- it's easy to understand once you
got it, but it also took me a too long time to... As you're saying, this
is a different kind of C/C++ pragma usage that what they're "normally"
used for.
This is not a priority right now, but if you or somebody wants to pick it
up, here is my WIP patch for the C PRAGMA_OACC_CACHE:
commit a294b3299d4214f2cbe611cc71ef190b247716cf
Author: Thomas Schwinge <tho...@codesourcery.com>
Date: Wed Nov 5 18:57:41 2014 +0100
OpenACC cache: pragma context
TODO: gcc/cp/
---
gcc/c/c-parser.c | 33 +++++++++++++----------
gcc/testsuite/c-c++-common/goacc/pragma_context.c | 13 +++++++--
2 files changed, 30 insertions(+), 16 deletions(-)
diff --git gcc/c/c-parser.c gcc/c/c-parser.c
index 9f4b013..57d29e5 100644
--- gcc/c/c-parser.c
+++ gcc/c/c-parser.c
@@ -1247,6 +1247,7 @@ static vec<tree, va_gc> *c_parser_expr_list (c_parser *,
bool, bool,
static tree c_parser_oacc_loop (location_t, c_parser *, char *);
static void c_parser_omp_construct (c_parser *);
static void c_parser_omp_threadprivate (c_parser *);
+static void c_parser_oacc_cache (c_parser *);
static void c_parser_oacc_enter_exit_data (c_parser *, bool);
static void c_parser_oacc_update (c_parser *);
static void c_parser_omp_barrier (c_parser *);
@@ -9575,6 +9576,17 @@ c_parser_pragma (c_parser *parser, enum pragma_context
context)
switch (id)
{
+ case PRAGMA_OACC_CACHE:
+ if (context != pragma_compound)
+ {
+ if (context == pragma_stmt)
+ c_parser_error (parser, "%<#pragma acc cache%> may only be "
+ "used in compound statements");
+ goto bad_stmt;
+ }
+ c_parser_oacc_cache (parser);
+ return false;
+
case PRAGMA_OACC_ENTER_DATA:
c_parser_oacc_enter_exit_data (parser, true);
return false;
@@ -11926,27 +11938,24 @@ c_parser_omp_structured_block (c_parser *parser)
/* OpenACC 2.0:
# pragma acc cache (variable-list) new-line
-
- LOC is the location of the #pragma token.
*/
-static tree
-c_parser_oacc_cache (location_t loc, c_parser *parser)
+static void
+c_parser_oacc_cache (c_parser *parser)
{
- tree stmt, clauses;
+ location_t loc = c_parser_peek_token (parser)->location;
- clauses = c_parser_omp_var_list_parens (parser, OMP_CLAUSE__CACHE_, NULL);
+ c_parser_consume_pragma (parser);
+
+ tree clauses = c_parser_omp_var_list_parens (parser, OMP_CLAUSE__CACHE_,
NULL);
clauses = c_finish_omp_clauses (clauses);
-
c_parser_skip_to_pragma_eol (parser);
- stmt = make_node (OACC_CACHE);
+ tree stmt = make_node (OACC_CACHE);
TREE_TYPE (stmt) = void_type_node;
OACC_CACHE_CLAUSES (stmt) = clauses;
SET_EXPR_LOCATION (stmt, loc);
add_stmt (stmt);
-
- return stmt;
}
/* OpenACC 2.0:
@@ -14749,10 +14758,6 @@ c_parser_omp_construct (c_parser *parser)
switch (p_kind)
{
- case PRAGMA_OACC_CACHE:
- strcpy (p_name, "#pragma acc");
- stmt = c_parser_oacc_cache (loc, parser);
- break;
case PRAGMA_OACC_DATA:
stmt = c_parser_oacc_data (loc, parser);
break;
diff --git gcc/testsuite/c-c++-common/goacc/pragma_context.c
gcc/testsuite/c-c++-common/goacc/pragma_context.c
index ad33d92..e76d2d6 100644
--- gcc/testsuite/c-c++-common/goacc/pragma_context.c
+++ gcc/testsuite/c-c++-common/goacc/pragma_context.c
@@ -1,15 +1,18 @@
// pragma_external
+#pragma acc cache /* { dg-error "expected declaration specifiers before
'#pragma'" } */
#pragma acc update /* { dg-error "expected declaration specifiers before
'#pragma'" } */
// pragma_struct
struct s_pragma_struct
{
+#pragma acc cache /* { dg-error "expected declaration specifiers before
'#pragma'" } */
#pragma acc update /* { dg-error "expected declaration specifiers before
'#pragma'" } */
};
// pragma_param
void
f_pragma_param (
+#pragma acc cache /* { dg-error "expected declaration specifiers before
'#pragma'" } */
#pragma acc update /* { dg-error "expected declaration specifiers before
'#pragma'" } */
void)
{
@@ -20,13 +23,19 @@ void
f2 (void)
{
if (0)
+#pragma acc cache /* { dg-error "'#pragma acc cache' may only be used in
compound statements before '#pragma'" } */
+ ;
+ if (0)
#pragma acc update /* { dg-error "'#pragma acc update' may only be used in
compound statements before '#pragma'" } */
+ ;
}
// pragma_compound
void
f3 (void)
{
- int i = 0;
-#pragma acc update device(i)
+ int i[10] = { 0 };
+
+#pragma acc cache (i[1:3])
+#pragma acc update device(i[3:2])
}
Grüße,
Thomas
pgpbOPFYsfN7u.pgp
Description: PGP signature
