diff options
author | Tobias Burnus <tobias@codesourcery.com> | 2020-11-06 11:13:47 +0100 |
---|---|---|
committer | Tobias Burnus <tobias@codesourcery.com> | 2020-11-06 11:13:47 +0100 |
commit | a2c11935b010ee55f7ccd14d27f62c6fbed3745e (patch) | |
tree | f1183d41725d18ca1ee0e5994fa96c5adce21d9e /gcc | |
parent | 8ebedfcd86aa5e3fc902fb442ce12c9d440c23c8 (diff) | |
download | gcc-a2c11935b010ee55f7ccd14d27f62c6fbed3745e.zip gcc-a2c11935b010ee55f7ccd14d27f62c6fbed3745e.tar.gz gcc-a2c11935b010ee55f7ccd14d27f62c6fbed3745e.tar.bz2 |
OpenACC (C/C++): Fix 'acc atomic' parsing
gcc/c/ChangeLog:
* c-parser.c (c_parser_omp_atomic): Add openacc parameter and update
OpenACC matching.
(c_parser_omp_construct): Update call.
gcc/cp/ChangeLog:
* parser.c (cp_parser_omp_atomic): Add openacc parameter and update
OpenACC matching.
(cp_parser_omp_construct): Update call.
gcc/testsuite/ChangeLog:
* c-c++-common/goacc-gomp/atomic.c: New test.
* c-c++-common/goacc/atomic.c: New test.
Diffstat (limited to 'gcc')
-rw-r--r-- | gcc/c/c-parser.c | 24 | ||||
-rw-r--r-- | gcc/cp/parser.c | 23 | ||||
-rw-r--r-- | gcc/testsuite/c-c++-common/goacc-gomp/atomic.c | 43 | ||||
-rw-r--r-- | gcc/testsuite/c-c++-common/goacc/atomic.c | 30 |
4 files changed, 110 insertions, 10 deletions
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index fc97aa3..dedfb84 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -17304,7 +17304,7 @@ c_parser_oacc_wait (location_t loc, c_parser *parser, char *p_name) LOC is the location of the #pragma token. */ static void -c_parser_omp_atomic (location_t loc, c_parser *parser) +c_parser_omp_atomic (location_t loc, c_parser *parser, bool openacc) { tree lhs = NULL_TREE, rhs = NULL_TREE, v = NULL_TREE; tree lhs1 = NULL_TREE, rhs1 = NULL_TREE; @@ -17343,6 +17343,12 @@ c_parser_omp_atomic (location_t loc, c_parser *parser) new_code = OMP_ATOMIC; else if (!strcmp (p, "capture")) new_code = OMP_ATOMIC_CAPTURE_NEW; + else if (openacc) + { + p = NULL; + error_at (cloc, "expected %<read%>, %<write%>, %<update%>, " + "or %<capture%> clause"); + } else if (!strcmp (p, "seq_cst")) new_memory_order = OMP_MEMORY_ORDER_SEQ_CST; else if (!strcmp (p, "acq_rel")) @@ -17370,7 +17376,12 @@ c_parser_omp_atomic (location_t loc, c_parser *parser) { if (new_code != ERROR_MARK) { - if (code != ERROR_MARK) + /* OpenACC permits 'update capture'. */ + if (openacc + && code == OMP_ATOMIC + && new_code == OMP_ATOMIC_CAPTURE_NEW) + code = new_code; + else if (code != ERROR_MARK) error_at (cloc, "too many atomic clauses"); else code = new_code; @@ -17392,7 +17403,9 @@ c_parser_omp_atomic (location_t loc, c_parser *parser) if (code == ERROR_MARK) code = OMP_ATOMIC; - if (memory_order == OMP_MEMORY_ORDER_UNSPECIFIED) + if (openacc) + memory_order = OMP_MEMORY_ORDER_RELAXED; + else if (memory_order == OMP_MEMORY_ORDER_UNSPECIFIED) { omp_requires_mask = (enum omp_requires) (omp_requires_mask @@ -17448,6 +17461,7 @@ c_parser_omp_atomic (location_t loc, c_parser *parser) } break; case OMP_ATOMIC: + /* case OMP_ATOMIC_CAPTURE_NEW: - or update to OpenMP 5.1 */ if (memory_order == OMP_MEMORY_ORDER_ACQ_REL || memory_order == OMP_MEMORY_ORDER_ACQUIRE) { @@ -21489,7 +21503,7 @@ c_parser_omp_construct (c_parser *parser, bool *if_p) switch (p_kind) { case PRAGMA_OACC_ATOMIC: - c_parser_omp_atomic (loc, parser); + c_parser_omp_atomic (loc, parser, true); return; case PRAGMA_OACC_CACHE: strcpy (p_name, "#pragma acc"); @@ -21516,7 +21530,7 @@ c_parser_omp_construct (c_parser *parser, bool *if_p) stmt = c_parser_oacc_wait (loc, parser, p_name); break; case PRAGMA_OMP_ATOMIC: - c_parser_omp_atomic (loc, parser); + c_parser_omp_atomic (loc, parser, false); return; case PRAGMA_OMP_CRITICAL: stmt = c_parser_omp_critical (loc, parser, if_p); diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index e7bfbf6..f030cad 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -37994,7 +37994,7 @@ cp_parser_omp_structured_block (cp_parser *parser, bool *if_p) where x and v are lvalue expressions with scalar type. */ static void -cp_parser_omp_atomic (cp_parser *parser, cp_token *pragma_tok) +cp_parser_omp_atomic (cp_parser *parser, cp_token *pragma_tok, bool openacc) { tree lhs = NULL_TREE, rhs = NULL_TREE, v = NULL_TREE, lhs1 = NULL_TREE; tree rhs1 = NULL_TREE, orig_lhs; @@ -38029,6 +38029,12 @@ cp_parser_omp_atomic (cp_parser *parser, cp_token *pragma_tok) new_code = OMP_ATOMIC; else if (!strcmp (p, "capture")) new_code = OMP_ATOMIC_CAPTURE_NEW; + else if (openacc) + { + p = NULL; + error_at (cloc, "expected %<read%>, %<write%>, %<update%>, " + "or %<capture%> clause"); + } else if (!strcmp (p, "seq_cst")) new_memory_order = OMP_MEMORY_ORDER_SEQ_CST; else if (!strcmp (p, "acq_rel")) @@ -38056,7 +38062,12 @@ cp_parser_omp_atomic (cp_parser *parser, cp_token *pragma_tok) { if (new_code != ERROR_MARK) { - if (code != ERROR_MARK) + /* OpenACC permits 'update capture'. */ + if (openacc + && code == OMP_ATOMIC + && new_code == OMP_ATOMIC_CAPTURE_NEW) + code = new_code; + else if (code != ERROR_MARK) error_at (cloc, "too many atomic clauses"); else code = new_code; @@ -38078,7 +38089,9 @@ cp_parser_omp_atomic (cp_parser *parser, cp_token *pragma_tok) if (code == ERROR_MARK) code = OMP_ATOMIC; - if (memory_order == OMP_MEMORY_ORDER_UNSPECIFIED) + if (openacc) + memory_order = OMP_MEMORY_ORDER_RELAXED; + else if (memory_order == OMP_MEMORY_ORDER_UNSPECIFIED) { omp_requires_mask = (enum omp_requires) (omp_requires_mask @@ -43517,7 +43530,7 @@ cp_parser_omp_construct (cp_parser *parser, cp_token *pragma_tok, bool *if_p) switch (cp_parser_pragma_kind (pragma_tok)) { case PRAGMA_OACC_ATOMIC: - cp_parser_omp_atomic (parser, pragma_tok); + cp_parser_omp_atomic (parser, pragma_tok, true); return; case PRAGMA_OACC_CACHE: stmt = cp_parser_oacc_cache (parser, pragma_tok); @@ -43552,7 +43565,7 @@ cp_parser_omp_construct (cp_parser *parser, cp_token *pragma_tok, bool *if_p) stmt = cp_parser_oacc_wait (parser, pragma_tok); break; case PRAGMA_OMP_ATOMIC: - cp_parser_omp_atomic (parser, pragma_tok); + cp_parser_omp_atomic (parser, pragma_tok, false); return; case PRAGMA_OMP_CRITICAL: stmt = cp_parser_omp_critical (parser, pragma_tok, if_p); diff --git a/gcc/testsuite/c-c++-common/goacc-gomp/atomic.c b/gcc/testsuite/c-c++-common/goacc-gomp/atomic.c new file mode 100644 index 0000000..4d18f23 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc-gomp/atomic.c @@ -0,0 +1,43 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-original" } */ + +#pragma omp requires atomic_default_mem_order(acq_rel) + +void +foo () +{ + int i, v; + +#pragma omp atomic read + i = v; + +#pragma acc atomic read + i = v; + +#pragma omp atomic write + i = v; + +#pragma acc atomic write + i = v; + +#pragma omp atomic update + i += 1; + +#pragma acc atomic update + i += 1; + +#pragma omp atomic capture + v = i += 1; + +#pragma acc atomic capture + v = i += 1; +#pragma acc atomic update capture + v = i += 1; +} + +/* { dg-final { scan-tree-dump-times "i = #pragma omp atomic read acquire" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "i = #pragma omp atomic read relaxed" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "#pragma omp atomic release" 2 "original" } } */ +/* { dg-final { scan-tree-dump-times "#pragma omp atomic relaxed" 2 "original" } } */ +/* { dg-final { scan-tree-dump-times "v = #pragma omp atomic capture acq_rel" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "v = #pragma omp atomic capture relaxed" 2 "original" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/atomic.c b/gcc/testsuite/c-c++-common/goacc/atomic.c new file mode 100644 index 0000000..ff3b25e --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/atomic.c @@ -0,0 +1,30 @@ +/* { dg-do compile } */ + +void +foo () +{ + int i, v; +#pragma acc atomic read bar /* { dg-error "expected 'read', 'write', 'update', or 'capture' clause" } */ + i = v; /* { dg-error "expected end of line before 'bar'" "" { target *-*-* } .-1 } */ + +#pragma acc atomic read write /* { dg-error "too many atomic clauses" } */ + i = v; + +#pragma acc atomic read seq_cst /* { dg-error "expected 'read', 'write', 'update', or 'capture' clause" } */ + i = v; /* { dg-error "expected end of line before 'seq_cst'" "" { target *-*-* } .-1 } */ + +#pragma acc atomic read relaxed /* { dg-error "expected 'read', 'write', 'update', or 'capture' clause" } */ + i = v; /* { dg-error "expected end of line before 'relaxed'" "" { target *-*-* } .-1 } */ + +#pragma acc atomic update hint(1) /* { dg-error "expected 'read', 'write', 'update', or 'capture' clause" } */ + i += 1; /* { dg-error "expected end of line before 'hint'" "" { target *-*-* } .-1 } */ + +#pragma acc atomic update update capture /* { dg-error "too many atomic clauses" } */ + v = i += 1; + +#pragma acc atomic update capture capture /* { dg-error "too many atomic clauses" } */ + v = i += 1; + +#pragma acc atomic write capture /* { dg-error "too many atomic clauses" } */ + i = 1; +} |