From: Chung-Lin Tang <cltang@codesourcery.com>
To: Thomas Schwinge <thomas@codesourcery.com>,
Jakub Jelinek <jakub@redhat.com>,
Cesar Philippidis <cesar_philippidis@mentor.com>
Cc: gcc-patches <gcc-patches@gcc.gnu.org>
Subject: Re: [PATCH, OpenACC] Make reduction arguments addressable
Date: Tue, 31 May 2016 11:10:00 -0000 [thread overview]
Message-ID: <3c5f38f7-ac02-6757-13bf-de48a1192048@codesourcery.com> (raw)
In-Reply-To: <878tyqsn32.fsf@kepler.schwinge.homeip.net>
On 2016/5/31 3:28 PM, Thomas Schwinge wrote:
> Hi!
>
> On Mon, 30 May 2016 18:53:41 +0200, Jakub Jelinek <jakub@redhat.com> wrote:
>> On Mon, May 30, 2016 at 10:38:59PM +0800, Chung-Lin Tang wrote:
>>> Hi, a previous patch of Cesar's has made the middle-end omp-lowering
>>> automatically create and insert a tofrom (i.e. present_or_copy) map for
>>> parallel reductions. This allowed the user to not need explicit
>>> clauses to copy out the reduction result, but because reduction arguments
>>> are not marked addressable, async does not work as expected,
>>> i.e. the asynchronous copy-out results are not used in the compiler generated code.
>>
>> If you need it only for async parallel/kernels? regions, can't you do that
>> only for those and not for others?
That is achievable, but not in line with how we currently treat all other
data clause OMP_CLAUSE_MAPs, which are all marked addressable. Is this special
case handling really better here?
> Also, please add comments to the source code to document the need for
> such special handling.
>
>>> This patch fixes this in the front-ends, I've tested this patch without
>>> new regressions, and fixes some C++ OpenACC tests that regressed after
>>> my last OpenACC async patch. Is this okay for trunk?
>>
>> Testcases in the testsuite or others? If the latter, we should add them.
>
> The r236772 commit "[PATCH, libgomp] Rewire OpenACC async",
> <http://news.gmane.org/find-root.php?message_id=%3C56FA4F69.1060101%40codesourcery.com%3E>
> regressed (or, triggered/exposed the existing wrong behavior?)
> libgomp.oacc-c++/template-reduction.C execution testing for nvptx
> offloading. (Please always send email about such known regressions, and
> XFAIL them with your commit -- that would have saved me an hour
> yesterday, when I bisected recent changes to figure out why that test
> suddenly fails.)
Sorry, Thomas. I was going to quickly send this follow-up patch, so glossed over XFAILing.
> For reference, here is a test case, a reduced C version of
> libgomp/testsuite/libgomp.oacc-c++/template-reduction.C. This test case
> works (without Chung-Lin's "[PATCH, OpenACC] Make reduction arguments
> addressable" patch) if I enable "POCs", which surprises me a bit, because
> I thought after Cesar's recent changes, the gimplifier is doing the same
> thing of adding a data clause next to the reduction clause. Probably
> it's not doing the exactly same thing, though. Should it? Cesar, do you
> have any comments on this? For example (just guessing), should
> TREE_ADDRESSABLE be set where the gimplifier does its work, instead of in
> the three front ends?
There's really nothing wrong about Cesar's patch. The marking addressable needs
to be done earlier, or it may be too late during gimplification. I already
tried to fix this in gimplify.c before, but didn't completely work.
I'll add more testcases for this before I commit any final patches.
Thanks,
Chung-Lin
>
> // Reduced C version of libgomp/testsuite/libgomp.oacc-c++/template-reduction.C.
>
> const int n = 100;
>
> // Check present and async and an explicit firstprivate
>
> int
> async_sum (int c)
> {
> int s = 0;
>
> #define POCs //present_or_copy(s)
> #pragma acc parallel loop num_gangs (10) gang reduction (+:s) POCs firstprivate (c) async
> for (int i = 0; i < n; i++)
> s += i+c;
>
> #pragma acc wait
>
> return s;
> }
>
> int
> main()
> {
> int result = 0;
>
> for (int i = 0; i < n; i++)
> {
> result += i+1;
> }
>
> if (async_sum (1) != result)
> __builtin_abort ();
>
> return 0;
> }
>
>
> GrüÃe
> Thomas
>
next prev parent reply other threads:[~2016-05-31 9:51 UTC|newest]
Thread overview: 11+ messages / expand[flat|nested] mbox.gz Atom feed top
2016-05-30 18:42 Chung-Lin Tang
2016-05-30 21:43 ` Jakub Jelinek
2016-05-31 9:59 ` Thomas Schwinge
2016-05-31 11:10 ` Chung-Lin Tang [this message]
2016-06-01 13:32 ` Chung-Lin Tang
2016-06-01 13:39 ` Jakub Jelinek
2016-06-02 13:55 ` Chung-Lin Tang
2016-06-02 14:00 ` Jakub Jelinek
2016-06-03 7:13 ` Chung-Lin Tang
2016-06-03 7:15 ` Jakub Jelinek
2016-06-10 10:22 ` Thomas Schwinge
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=3c5f38f7-ac02-6757-13bf-de48a1192048@codesourcery.com \
--to=cltang@codesourcery.com \
--cc=cesar_philippidis@mentor.com \
--cc=gcc-patches@gcc.gnu.org \
--cc=jakub@redhat.com \
--cc=thomas@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).