mirror of
git://gcc.gnu.org/git/gcc.git
synced 2025-03-21 23:41:13 +08:00
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.
This commit is contained in:
Tobias Burnus
parent
8ebedfcd86
commit
a2c11935b0
gcc
@ -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);
|
||||
|
||||
@ -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);
|
||||
|
||||
43
gcc/testsuite/c-c++-common/goacc-gomp/atomic.c
Normal file
43
gcc/testsuite/c-c++-common/goacc-gomp/atomic.c
Normal file
@ -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" } } */
|
||||
30
gcc/testsuite/c-c++-common/goacc/atomic.c
Normal file
30
gcc/testsuite/c-c++-common/goacc/atomic.c
Normal file
@ -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;
|
||||
}
|
||||
Loading…
x
Reference in New Issue
Block a user