From: Chung-Lin Tang <cltang@codesourcery.com>
To: Tom de Vries <Tom_deVries@mentor.com>,
gcc-patches <gcc-patches@gcc.gnu.org>
Cc: Tom de Vries <vries@codesourcery.com>,
Thomas Schwinge <thomas@codesourcery.com>
Subject: Re: [gomp4] Generate sequential loop for OpenACC loop directive inside kernels
Date: Tue, 23 Jun 2015 12:56:00 -0000 [thread overview]
Message-ID: <55895746.50506@codesourcery.com> (raw)
In-Reply-To: <557FE6C8.6000706@mentor.com>
[-- Attachment #1: Type: text/plain, Size: 1340 bytes --]
On 2015/6/16 05:05 PM, Tom de Vries wrote:
> 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
Just committed the attached testcase patch to gomp-4_0-branch.
Chung-Lin
2015-06-23 Chung-Lin Tang <cltang@codesourcery.com>
gcc/testsuite/
* c-c++-common/goacc/kernels-loop.c (ACC_LOOP): Add #ifndef/#define.
(main): Tag loops inside kernels construct with '#pragma ACC_LOOP'.
* c-c++-common/goacc/kernels-loop-2.c: Likewise.
* c-c++-common/goacc/kernels-loop-3.c: Likewise.
* c-c++-common/goacc/kernels-loop-n.c: Likewise.
* c-c++-common/goacc/kernels-loop-acc-loop.c: New test.
* c-c++-common/goacc/kernels-loop-2-acc-loop.c: New test.
* c-c++-common/goacc/kernels-loop-3-acc-loop.c: New test.
* c-c++-common/goacc/kernels-loop-n-acc-loop.c: New test.
[-- Attachment #2: kernels-acc-loop.patch --]
[-- Type: text/x-patch, Size: 8286 bytes --]
Index: gcc/testsuite/c-c++-common/goacc/kernels-loop-3-acc-loop.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/kernels-loop-3-acc-loop.c (revision 0)
+++ gcc/testsuite/c-c++-common/goacc/kernels-loop-3-acc-loop.c (revision 0)
@@ -0,0 +1,20 @@
+/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "-ftree-parallelize-loops=32" } */
+/* { dg-additional-options "-fdump-tree-parloops_oacc_kernels-all" } */
+/* { dg-additional-options "-fdump-tree-optimized" } */
+
+/* Check that loops with '#pragma acc loop' tagged gets properly parallelized. */
+#define ACC_LOOP acc loop
+#include "kernels-loop-3.c"
+
+/* Check that only one loop is analyzed, and that it can be parallelized. */
+/* { dg-final { scan-tree-dump-times "SUCCESS: may be parallelized" 1 "parloops_oacc_kernels" } } */
+/* { dg-final { scan-tree-dump-not "FAILED:" "parloops_oacc_kernels" } } */
+
+/* Check that the loop has been split off into a function. */
+/* { dg-final { scan-tree-dump-times "(?n);; Function .*main._omp_fn.0" 1 "optimized" } } */
+
+/* { dg-final { scan-tree-dump-times "(?n)pragma omp target oacc_parallel.*num_gangs\\(32\\)" 1 "parloops_oacc_kernels" } } */
+
+/* { dg-final { cleanup-tree-dump "parloops_oacc_kernels" } } */
+/* { dg-final { cleanup-tree-dump "optimized" } } */
Index: gcc/testsuite/c-c++-common/goacc/kernels-loop-2.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/kernels-loop-2.c (revision 224836)
+++ gcc/testsuite/c-c++-common/goacc/kernels-loop-2.c (working copy)
@@ -8,6 +8,10 @@
#define N (1024 * 512)
#define COUNTERTYPE unsigned int
+#ifndef ACC_LOOP
+#define ACC_LOOP
+#endif
+
int
main (void)
{
@@ -21,18 +25,21 @@ main (void)
#pragma acc kernels copyout (a[0:N])
{
+ #pragma ACC_LOOP
for (COUNTERTYPE i = 0; i < N; i++)
a[i] = i * 2;
}
#pragma acc kernels copyout (b[0:N])
{
+ #pragma ACC_LOOP
for (COUNTERTYPE i = 0; i < N; i++)
b[i] = i * 4;
}
#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N])
{
+ #pragma ACC_LOOP
for (COUNTERTYPE ii = 0; ii < N; ii++)
c[ii] = a[ii] + b[ii];
}
Index: gcc/testsuite/c-c++-common/goacc/kernels-loop.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/kernels-loop.c (revision 224836)
+++ gcc/testsuite/c-c++-common/goacc/kernels-loop.c (working copy)
@@ -8,6 +8,10 @@
#define N (1024 * 512)
#define COUNTERTYPE unsigned int
+#ifndef ACC_LOOP
+#define ACC_LOOP
+#endif
+
int
main (void)
{
@@ -27,6 +31,7 @@ main (void)
#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N])
{
+ #pragma ACC_LOOP
for (COUNTERTYPE ii = 0; ii < N; ii++)
c[ii] = a[ii] + b[ii];
}
Index: gcc/testsuite/c-c++-common/goacc/kernels-loop-2-acc-loop.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/kernels-loop-2-acc-loop.c (revision 0)
+++ gcc/testsuite/c-c++-common/goacc/kernels-loop-2-acc-loop.c (revision 0)
@@ -0,0 +1,23 @@
+/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "-ftree-parallelize-loops=32" } */
+/* { dg-additional-options "-fdump-tree-parloops_oacc_kernels-all" } */
+/* { dg-additional-options "-fdump-tree-optimized" } */
+
+/* Check that loops with '#pragma acc loop' tagged gets properly parallelized. */
+#define ACC_LOOP acc loop
+#include "kernels-loop-2.c"
+
+/* Check that only three loops are analyzed, and that all can be
+ parallelized. */
+/* { dg-final { scan-tree-dump-times "SUCCESS: may be parallelized" 3 "parloops_oacc_kernels" } } */
+/* { dg-final { scan-tree-dump-not "FAILED:" "parloops_oacc_kernels" } } */
+
+/* Check that the loop has been split off into a function. */
+/* { dg-final { scan-tree-dump-times "(?n);; Function .*main._omp_fn.0" 1 "optimized" } } */
+/* { dg-final { scan-tree-dump-times "(?n);; Function .*main._omp_fn.1" 1 "optimized" } } */
+/* { dg-final { scan-tree-dump-times "(?n);; Function .*main._omp_fn.2" 1 "optimized" } } */
+
+/* { dg-final { scan-tree-dump-times "(?n)pragma omp target oacc_parallel.*num_gangs\\(32\\)" 3 "parloops_oacc_kernels" } } */
+
+/* { dg-final { cleanup-tree-dump "parloops_oacc_kernels" } } */
+/* { dg-final { cleanup-tree-dump "optimized" } } */
Index: gcc/testsuite/c-c++-common/goacc/kernels-loop-n.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/kernels-loop-n.c (revision 224836)
+++ gcc/testsuite/c-c++-common/goacc/kernels-loop-n.c (working copy)
@@ -8,6 +8,10 @@
#define N ((1024 * 512) + 1)
#define COUNTERTYPE unsigned int
+#ifndef ACC_LOOP
+#define ACC_LOOP
+#endif
+
int
foo (COUNTERTYPE n)
{
@@ -27,6 +31,7 @@ foo (COUNTERTYPE n)
#pragma acc kernels copyin (a[0:n], b[0:n]) copyout (c[0:n])
{
+ #pragma ACC_LOOP
for (COUNTERTYPE ii = 0; ii < n; ii++)
c[ii] = a[ii] + b[ii];
}
Index: gcc/testsuite/c-c++-common/goacc/kernels-loop-acc-loop.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/kernels-loop-acc-loop.c (revision 0)
+++ gcc/testsuite/c-c++-common/goacc/kernels-loop-acc-loop.c (revision 0)
@@ -0,0 +1,20 @@
+/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "-ftree-parallelize-loops=32" } */
+/* { dg-additional-options "-fdump-tree-parloops_oacc_kernels-all" } */
+/* { dg-additional-options "-fdump-tree-optimized" } */
+
+/* Check that loops with '#pragma acc loop' tagged gets properly parallelized. */
+#define ACC_LOOP acc loop
+#include "kernels-loop.c"
+
+/* Check that only one loop is analyzed, and that it can be parallelized. */
+/* { dg-final { scan-tree-dump-times "SUCCESS: may be parallelized" 1 "parloops_oacc_kernels" } } */
+/* { dg-final { scan-tree-dump-not "FAILED:" "parloops_oacc_kernels" } } */
+
+/* Check that the loop has been split off into a function. */
+/* { dg-final { scan-tree-dump-times "(?n);; Function .*main._omp_fn.0" 1 "optimized" } } */
+
+/* { dg-final { scan-tree-dump-times "(?n)pragma omp target oacc_parallel.*num_gangs\\(32\\)" 1 "parloops_oacc_kernels" } } */
+
+/* { dg-final { cleanup-tree-dump "parloops_oacc_kernels" } } */
+/* { dg-final { cleanup-tree-dump "optimized" } } */
Index: gcc/testsuite/c-c++-common/goacc/kernels-loop-3.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/kernels-loop-3.c (revision 224836)
+++ gcc/testsuite/c-c++-common/goacc/kernels-loop-3.c (working copy)
@@ -8,6 +8,10 @@
#define N (1024 * 512)
#define COUNTERTYPE unsigned int
+#ifndef ACC_LOOP
+#define ACC_LOOP
+#endif
+
int
main (void)
{
@@ -22,6 +26,7 @@ main (void)
#pragma acc kernels copy (c[0:N])
{
+ #pragma ACC_LOOP
for (COUNTERTYPE ii = 0; ii < N; ii++)
c[ii] = c[ii] + ii + 1;
}
Index: gcc/testsuite/c-c++-common/goacc/kernels-loop-n-acc-loop.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/kernels-loop-n-acc-loop.c (revision 0)
+++ gcc/testsuite/c-c++-common/goacc/kernels-loop-n-acc-loop.c (revision 0)
@@ -0,0 +1,20 @@
+/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "-ftree-parallelize-loops=32" } */
+/* { dg-additional-options "-fdump-tree-parloops_oacc_kernels-all" } */
+/* { dg-additional-options "-fdump-tree-optimized" } */
+
+/* Check that loops with '#pragma acc loop' tagged gets properly parallelized. */
+#define ACC_LOOP acc loop
+#include "kernels-loop-n.c"
+
+/* Check that only one loop is analyzed, and that it can be parallelized. */
+/* { dg-final { scan-tree-dump-times "SUCCESS: may be parallelized" 1 "parloops_oacc_kernels" } } */
+/* { dg-final { scan-tree-dump-not "FAILED:" "parloops_oacc_kernels" } } */
+
+/* Check that the loop has been split off into a function. */
+/* { dg-final { scan-tree-dump-times "(?n);; Function .*foo.*._omp_fn.0" 1 "optimized" } } */
+
+/* { dg-final { scan-tree-dump-times "(?n)pragma omp target oacc_parallel.*num_gangs\\(32\\)" 1 "parloops_oacc_kernels" } } */
+
+/* { dg-final { cleanup-tree-dump "parloops_oacc_kernels" } } */
+/* { dg-final { cleanup-tree-dump "optimized" } } */
prev parent reply other threads:[~2015-06-23 12:55 UTC|newest]
Thread overview: 3+ messages / expand[flat|nested] mbox.gz Atom feed top
2015-06-16 9:01 Chung-Lin Tang
2015-06-16 9:05 ` Tom de Vries
2015-06-23 12:56 ` Chung-Lin Tang [this message]
Reply instructions:
You may reply publicly to this message via plain-text email
using any one of the following methods:
* Save the following mbox file, import it into your mail client,
and reply-to-all from there: mbox
Avoid top-posting and favor interleaved quoting:
https://en.wikipedia.org/wiki/Posting_style#Interleaved_style
* Reply using the --to, --cc, and --in-reply-to
switches of git-send-email(1):
git send-email \
--in-reply-to=55895746.50506@codesourcery.com \
--to=cltang@codesourcery.com \
--cc=Tom_deVries@mentor.com \
--cc=gcc-patches@gcc.gnu.org \
--cc=thomas@codesourcery.com \
--cc=vries@codesourcery.com \
/path/to/YOUR_REPLY
https://kernel.org/pub/software/scm/git/docs/git-send-email.html
* If your mail client supports setting the In-Reply-To header
via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line
before the message body.
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).