diff mbox

[gomp4] OpenACC Executable Directives

Message ID 87egtgs6gx.fsf@schwinge.name
State New
Headers show

Commit Message

Thomas Schwinge Nov. 6, 2014, 3:18 p.m. UTC 
Hi Jim!

On Tue, 4 Nov 2014 14:18:43 -0600, James Norris <James_Norris@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 <thomas@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(-)



Grüße,
 Thomas
diff mbox

Patch

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])
 }