[gomp4] OpenACC cache directive for C.
Thomas Schwinge
thomas@codesourcery.com
Wed Nov 5 16:29:00 GMT 2014
Hi!
In r217145, I applied Jim's patch to gomp-4_0-branch:
commit 4361f9b6b2c74c2961c3a5290a4945abe2d7a444
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Wed Nov 5 16:26:47 2014 +0000
OpenACC cache directive for C.
gcc/c-family/
* c-pragma.c (oacc_pragmas): Add "cache".
gcc/c/
* c-parser.c (c_parser_omp_variable_list): Handle
OMP_NO_CLAUSE_CACHE.
(c_parser_oacc_cache): New function.
(c_parser_omp_construct): Use it for PRAGMA_OACC_CACHE.
libgomp/
* testsuite/libgomp.oacc-c/cache-1.c: New file.
* testsuite/libgomp.oacc-c++/cache-1.C: Likewise.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@217145 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/c-family/ChangeLog.gomp | 4 ++
gcc/c-family/c-pragma.c | 1 +
gcc/c/ChangeLog.gomp | 7 +++
gcc/c/c-parser.c | 50 ++++++++++++++++++++
libgomp/ChangeLog.gomp | 5 ++
libgomp/testsuite/libgomp.oacc-c++/cache-1.C | 50 ++++++++++++++++++++
libgomp/testsuite/libgomp.oacc-c/cache-1.c | 69 ++++++++++++++++++++++++++++
7 files changed, 186 insertions(+)
diff --git gcc/c-family/ChangeLog.gomp gcc/c-family/ChangeLog.gomp
index 2f8c8a6..5f3b641 100644
--- gcc/c-family/ChangeLog.gomp
+++ gcc/c-family/ChangeLog.gomp
@@ -1,3 +1,7 @@
+2014-11-05 James Norris <jnorris@codesourcery.com>
+
+ * c-pragma.c (oacc_pragmas): Add "cache".
+
2014-11-03 Cesar Philippidis <cesar@codesourcery.com>
* c-pragma.c (oacc_pragmas): Add entries for PRAGMA_OACC_ENTER_DATA
diff --git gcc/c-family/c-pragma.c gcc/c-family/c-pragma.c
index a28727e..e98b555 100644
--- gcc/c-family/c-pragma.c
+++ gcc/c-family/c-pragma.c
@@ -1181,6 +1181,7 @@ static vec<pragma_ns_name> registered_pp_pragmas;
struct omp_pragma_def { const char *name; unsigned int id; };
static const struct omp_pragma_def oacc_pragmas[] = {
+ { "cache", PRAGMA_OACC_CACHE },
{ "data", PRAGMA_OACC_DATA },
{ "enter", PRAGMA_OACC_ENTER_DATA },
{ "exit", PRAGMA_OACC_EXIT_DATA },
diff --git gcc/c/ChangeLog.gomp gcc/c/ChangeLog.gomp
index f4e2010..7acd7b3 100644
--- gcc/c/ChangeLog.gomp
+++ gcc/c/ChangeLog.gomp
@@ -1,3 +1,10 @@
+2014-11-05 James Norris <jnorris@codesourcery.com>
+
+ * c-parser.c (c_parser_omp_variable_list): Handle
+ OMP_NO_CLAUSE_CACHE.
+ (c_parser_oacc_cache): New function.
+ (c_parser_omp_construct): Use it for PRAGMA_OACC_CACHE.
+
2014-11-03 Cesar Philippidis <cesar@codesourcery.com>
* c-parser.c (c_parser_oacc_enter_exit_data): New function.
diff --git gcc/c/c-parser.c gcc/c/c-parser.c
index 6ac1ace..410b19f 100644
--- gcc/c/c-parser.c
+++ gcc/c/c-parser.c
@@ -10053,6 +10053,14 @@ c_parser_omp_variable_list (c_parser *parser,
{
switch (kind)
{
+ case OMP_NO_CLAUSE_CACHE:
+ if (c_parser_peek_token (parser)->type != CPP_OPEN_SQUARE)
+ {
+ c_parser_error (parser, "expected %<[%>");
+ t = error_mark_node;
+ break;
+ }
+ /* FALL THROUGH. */
case OMP_CLAUSE_MAP:
case OMP_CLAUSE_FROM:
case OMP_CLAUSE_TO:
@@ -10091,6 +10099,29 @@ c_parser_omp_variable_list (c_parser *parser,
t = error_mark_node;
break;
}
+
+ if (kind == OMP_NO_CLAUSE_CACHE)
+ {
+ mark_exp_read (low_bound);
+ mark_exp_read (length);
+
+ if (TREE_CODE (low_bound) != INTEGER_CST
+ && !TREE_READONLY (low_bound))
+ {
+ error_at (clause_loc,
+ "%qD is not a constant", low_bound);
+ t = error_mark_node;
+ }
+
+ if (TREE_CODE (length) != INTEGER_CST
+ && !TREE_READONLY (length))
+ {
+ error_at (clause_loc,
+ "%qD is not a constant", length);
+ t = error_mark_node;
+ }
+ }
+
t = tree_cons (low_bound, length, t);
}
break;
@@ -11864,6 +11895,21 @@ 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 __attribute__((unused)), c_parser *parser)
+{
+ c_parser_omp_var_list_parens (parser, OMP_NO_CLAUSE_CACHE, NULL);
+ c_parser_skip_to_pragma_eol (parser);
+
+ return NULL_TREE;
+}
+
+/* OpenACC 2.0:
# pragma acc data oacc-data-clause[optseq] new-line
structured-block
@@ -14506,6 +14552,10 @@ 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 libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp
index 8dc947d..4ac348d 100644
--- libgomp/ChangeLog.gomp
+++ libgomp/ChangeLog.gomp
@@ -1,3 +1,8 @@
+2014-11-05 James Norris <jnorris@codesourcery.com>
+
+ * testsuite/libgomp.oacc-c/cache-1.c: New file.
+ * testsuite/libgomp.oacc-c++/cache-1.C: Likewise.
+
2014-11-05 Thomas Schwinge <thomas@codesourcery.com>
James Norris <jnorris@codesourcery.com>
diff --git libgomp/testsuite/libgomp.oacc-c++/cache-1.C libgomp/testsuite/libgomp.oacc-c++/cache-1.C
new file mode 100644
index 0000000..eb94a68
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c++/cache-1.C
@@ -0,0 +1,50 @@
+/* { dg-do compile } */
+
+#include <stdlib.h>
+
+int
+main (int argc, char **argv)
+{
+#define N 2
+ int a[N], b[N];
+ int i;
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 3;
+ b[i] = 0;
+ }
+
+#pragma acc parallel copyin (a[0:N]) copyout (b[0:N])
+{
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ {
+ const int idx = ii;
+ int n = 1;
+ const int len = n;
+
+#pragma acc cache /* { dg-error "error: expected '\\(' before end of line" } */
+
+#pragma acc cache (a) /* { dg-error "error: expected '\\\['" } */
+
+#pragma acc cache (a[0:N]) copyin (a[0:N]) /* { dg-error "error: expected end of line before 'copyin'" } */
+
+#pragma acc cache () /* { dg-error "error: expected unqualified-id before '\\)' token" } */
+
+#pragma acc cache (a[0:N] b[0:N}) /* { dg-error "error: expected end of line before '\\\}' token|error: expected '\\)' before 'b'" } */
+
+ b[ii] = a[ii];
+ }
+}
+
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != b[i])
+ abort ();
+ }
+
+ return 0;
+}
diff --git libgomp/testsuite/libgomp.oacc-c/cache-1.c libgomp/testsuite/libgomp.oacc-c/cache-1.c
new file mode 100644
index 0000000..5473475
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c/cache-1.c
@@ -0,0 +1,69 @@
+/* { dg-do compile } */
+
+#include <openacc.h>
+#include <stdlib.h>
+
+int
+main (int argc, char **argv)
+{
+#define N 2
+ int a[N], b[N];
+ int i;
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 3;
+ b[i] = 0;
+ }
+
+#pragma acc parallel copyin (a[0:N]) copyout (b[0:N])
+{
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ {
+ const int idx = ii;
+ int n = 1;
+ const int len = n;
+
+#pragma acc cache /* { dg-error "error: expected '\\(' before end of line" } */
+
+#pragma acc cache (a) /* { dg-error "error: expected '\\\[' before '\\)' token" } */
+
+#pragma acc cache (a[0:N]) copyin (a[0:N]) /* { dg-error "error: expected end of line before 'copyin'" } */
+
+#pragma acc cache () /* { dg-error "error: expected identifier before '\\)' token" } */
+
+#pragma acc cache (a[0:N] b[0:N}) /* { dg-error "error: expected '\\)' before 'b'" } */
+
+#pragma acc cache (a[0:N] /* { dg-error "error: expected '\\)' before end of line" } */
+
+#pragma acc cache (a[ii]) /* { dg-error "error: 'ii' is not a constant" } */
+
+#pragma acc cache (a[idx:n]) /* { dg-error "error: 'n' is not a constant" } */
+
+#pragma acc cache (a[0:N])
+
+#pragma acc cache (a[0:N], b[0:N])
+
+#pragma acc cache (a[0])
+
+#pragma acc cache (a[0], a[1], b[0:N])
+
+#pragma acc cache (a[idx])
+
+#pragma acc cache (a[idx:len])
+
+ b[ii] = a[ii];
+ }
+}
+
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != b[i])
+ abort ();
+ }
+
+ return 0;
+}
Grüße,
Thomas
-------------- next part --------------
A non-text attachment was scrubbed...
Name: not available
Type: application/pgp-signature
Size: 472 bytes
Desc: not available
URL: <http://gcc.gnu.org/pipermail/gcc-patches/attachments/20141105/93bc9f01/attachment.sig>
More information about the Gcc-patches
mailing list