diff options
| author | Tobias Burnus <tobias@codesourcery.com> | 2022-05-17 22:09:16 +0200 |
|---|---|---|
| committer | Tobias Burnus <tobias@codesourcery.com> | 2022-05-17 22:09:16 +0200 |
| commit | 47554478a13f64bff1ee4b9bb0319ae63d42ca52 (patch) | |
| tree | 42acab11fa43a8ff23a2b29c37ed310be3a15f04 /gcc | |
| parent | ddb1427defe95341ac2eb672e7bea7303f8c7db9 (diff) | |
| download | gcc-47554478a13f64bff1ee4b9bb0319ae63d42ca52.zip gcc-47554478a13f64bff1ee4b9bb0319ae63d42ca52.tar.gz gcc-47554478a13f64bff1ee4b9bb0319ae63d42ca52.tar.bz2 | |
OpenMP: Skip target-nesting warning for reverse offload
gcc/ChangeLog:
* omp-low.cc (check_omp_nesting_restrictions): Skip warning for
target inside target if inner is reverse offload.
gcc/testsuite/ChangeLog:
* c-c++-common/gomp/target-device-ancestor-5.c: New test.
Diffstat (limited to 'gcc')
| -rw-r--r-- | gcc/omp-low.cc | 10 | ||||
| -rw-r--r-- | gcc/testsuite/c-c++-common/gomp/target-device-ancestor-5.c | 28 |
2 files changed, 38 insertions, 0 deletions
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index b97e0e9..c83af6c 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -3883,6 +3883,16 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx) } else { + if ((gimple_omp_target_kind (ctx->stmt) + == GF_OMP_TARGET_KIND_REGION) + && (gimple_omp_target_kind (stmt) + == GF_OMP_TARGET_KIND_REGION)) + { + c = omp_find_clause (gimple_omp_target_clauses (stmt), + OMP_CLAUSE_DEVICE); + if (c && OMP_CLAUSE_DEVICE_ANCESTOR (c)) + break; + } warning_at (gimple_location (stmt), 0, "%qs construct inside of %qs region", stmt_name, ctx_stmt_name); diff --git a/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-5.c b/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-5.c new file mode 100644 index 0000000..b6ff84b --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-5.c @@ -0,0 +1,28 @@ +#pragma omp requires reverse_offload /* { dg-message "sorry, unimplemented: 'reverse_offload' clause on 'requires' directive not supported yet" } */ + +void +foo () +{ + /* Good nesting - as reverse offload */ + #pragma omp target + #pragma omp target device(ancestor:1) /* valid -> no warning */ /* { dg-bogus "'target' construct inside of 'target' region" } */ + { } + + /* Bad nesting */ + #pragma omp target + #pragma omp target /* { dg-warning "'target' construct inside of 'target' region" } */ + #pragma omp target /* { dg-warning "'target' construct inside of 'target' region" } */ + { } + + /* Good nesting - as reverse offload */ + #pragma omp target + #pragma omp target /* { dg-warning "'target' construct inside of 'target' region" } */ + #pragma omp target device(ancestor:1) /* valid -> no warning */ /* { dg-bogus "'target' construct inside of 'target' region" } */ + { } + + #pragma omp target + #pragma omp target device(ancestor:1) /* valid -> no warning */ /* { dg-bogus "'target' construct inside of 'target' region" } */ + #pragma omp target device(ancestor:1) /* { dg-error "OpenMP constructs are not allowed in target region with 'ancestor'" } */ + { } + +} |