@@ -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
@@ -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 },
@@ -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.
@@ -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;
@@ -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>
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}