This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
Re: [gomp4] Generate sequential loop for OpenACC loop directive inside kernels
- From: Tom de Vries <Tom_deVries at mentor dot com>
- To: Chung-Lin Tang <cltang at codesourcery dot com>, gcc-patches <gcc-patches at gcc dot gnu dot org>
- Cc: Tom de Vries <vries at codesourcery dot com>, Thomas Schwinge <thomas at codesourcery dot com>
- Date: Tue, 16 Jun 2015 11:05:12 +0200
- Subject: Re: [gomp4] Generate sequential loop for OpenACC loop directive inside kernels
- Authentication-results: sourceware.org; auth=none
- References: <557FE579 dot 30304 at codesourcery dot com>
On 16/06/15 10:59, Chung-Lin Tang wrote:
This patch adjusts omp-low.c:expand_omp_for_generic() to expand to a "sequential"
loop form (without the OMP runtime calls), used for loop directives inside
OpenACC kernels constructs. Tom mentions that this allows the kernels parallelization
to work when '#pragma acc loop' makes the front-ends create OMP_FOR, which the
loop analysis phases don't understand.
Tested and committed to gomp-4_0-branch.
Hi Chung-Lin,
can you commit a test-case to exercise the code?
Thanks,
- Tom
Chung-Lin
2015-06-16 Chung-Lin Tang <cltang@codesourcery.com>
* omp-low.c (struct omp_region): Add inside_kernels_p field.
(expand_omp_for_generic): Adjust to generate a 'sequential' loop
when GOMP builtin arguments are BUILT_IN_NONE.
(expand_omp_for): Use expand_omp_for_generic() to generate a
non-parallelized loop for OMP_FORs inside OpenACC kernels regions.
(expand_omp): Mark inside_kernels_p field true for regions
nested inside OpenACC kernels constructs.