diff options
| author | Sandra Loosemore <sandra@codesourcery.com> | 2020-08-19 19:18:57 -0700 |
|---|---|---|
| committer | Kwok Cheung Yeung <kcy@codesourcery.com> | 2023-05-12 19:13:46 +0100 |
| commit | 82abd6233bf09b7c2e0b3526b7ea74aea865b978 (patch) | |
| tree | e3043201cb29e46d5b445351fddda8a10b18cbc7 /gcc/c-family | |
| parent | 020f4009ee8291c842cd64a0eacbb244546261d9 (diff) | |
Annotate inner loops in "acc kernels loop" directives (C/C++).
Normally explicit loop directives in a kernels region inhibit
automatic annotation of other loops in the same nest, on the theory
that users have indicated they want manual control over that section
of code. However there seems to be an expectation in user code that
the combined "kernels loop" directive should still allow annotation of
inner loops. This patch implements this behavior for C and C++.
2020-08-19 Sandra Loosemore <sandra@codesourcery.com>
gcc/c-family/
* c-omp.cc (annotate_loops_in_kernels_regions): Process inner
loops in combined "acc kernels loop" directives.
gcc/testsuite/
* c-c++-common/goacc/kernels-loop-annotation-18.c: New.
* c-c++-common/goacc/kernels-loop-annotation-19.c: New.
* c-c++-common/goacc/combined-directives.c: Adjust expected
patterns.
Diffstat (limited to 'gcc/c-family')
| -rw-r--r-- | gcc/c-family/ChangeLog.omp | 7 | ||||
| -rw-r--r-- | gcc/c-family/c-omp.cc | 36 |
2 files changed, 31 insertions, 12 deletions
diff --git a/gcc/c-family/ChangeLog.omp b/gcc/c-family/ChangeLog.omp index 5da3b329dc8..3961227eb2a 100644 --- a/gcc/c-family/ChangeLog.omp +++ b/gcc/c-family/ChangeLog.omp @@ -1,3 +1,10 @@ +2020-08-19 Sandra Loosemore <sandra@codesourcery.com> + + Annotate inner loops in "acc kernels loop" directives (C/C++). + + * c-omp.cc (annotate_loops_in_kernels_regions): Process inner + loops in combined "acc kernels loop" directives. + 2020-03-27 Sandra Loosemore <sandra@codesourcery.com> * c-common.h (c_oacc_annotate_loops_in_kernels_regions): Declare. diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc index 2c25dff3f88..d9961fbf982 100644 --- a/gcc/c-family/c-omp.cc +++ b/gcc/c-family/c-omp.cc @@ -3484,18 +3484,30 @@ annotate_loops_in_kernels_regions (tree *nodeptr, int *walk_subtrees, /* Do not try to add automatic OpenACC annotations inside manually annotated loops. Presumably, the user avoided doing it on purpose; for example, all available levels of parallelism may - have been used up. */ - { - struct annotation_info nested_info - = { NULL_TREE, NULL_TREE, false, as_explicit_annotation, - node, info }; - if (info->state >= as_in_kernels_region) - do_not_annotate_loop_nest (info, as_explicit_annotation, - node); - walk_tree (&OMP_BODY (node), annotate_loops_in_kernels_regions, - (void *) &nested_info, NULL); - *walk_subtrees = 0; - } + have been used up. However, assume that the combined construct + "#pragma acc kernels loop" means to try to process the whole + loop nest. + Note that a single OACC_LOOP construct represents an entire set + of collapsed loops so we do not have to deal explicitly with the + collapse clause here, as the Fortran front end does. */ + if (info->state == as_in_kernels_region && OACC_LOOP_COMBINED (node)) + { + walk_tree (&OMP_BODY (node), annotate_loops_in_kernels_regions, + (void *) info, NULL); + *walk_subtrees = 0; + } + else + { + struct annotation_info nested_info + = { NULL_TREE, NULL_TREE, false, as_explicit_annotation, + node, info }; + if (info->state >= as_in_kernels_region) + do_not_annotate_loop_nest (info, as_explicit_annotation, + node); + walk_tree (&OMP_BODY (node), annotate_loops_in_kernels_regions, + (void *) &nested_info, NULL); + *walk_subtrees = 0; + } break; case FOR_STMT: |