From: Jakub Jelinek <jakub@redhat.com>
To: Alexander Monakov <amonakov@ispras.ru>
Cc: gcc-patches@gcc.gnu.org
Subject: Re: [PATCH 2/5] omp-low: implement SIMT privatization, part 1
Date: Thu, 20 Apr 2017 15:32:00 -0000 [thread overview]
Message-ID: <20170420151622.GV1809@tucnak> (raw)
In-Reply-To: <alpine.LNX.2.20.13.1703311800330.7566@monopod.intra.ispras.ru>
On Fri, Mar 31, 2017 at 06:38:09PM +0300, Alexander Monakov wrote:
> I've noticed while re-reading that this patch incorrectly handled SIMT case
> in lower_lastprivate_clauses. The code was changed to look for variables
> with "omp simt private" attribute, and was left under
> 'simduid && DECL_HAS_VALUE_EXPR_P (new_var)' condition. This effectively
> constrained processing to privatized (i.e. addressable) variables, as
> non-addressable variables now have neither the value-expr nor the attribute.
Sorry for the review delay.
> This wasn't caught in testing, as apparently all testcases that have target
> simd loops with linear/lastprivate clauses also have the corresponding variables
> mentioned in target map clause, which makes them addressable (is that necessary?),
Yes, in order to map something you need to map its address (+ size) on the
host to its address on the device, so it needs to be addressable.
Compared to that, firstprivate on target should not make it addressable.
> and I didn't think to check something like that manually.
>
> The following patch fixes it by adjusting the if's in lower_lastprivate_clauses;
> alternatively it may be possible to keep that code as-is, and instead set up
> value-expr and "omp simt private" attributes for all formally private variables,
> not just addressable ones, but to me that sounds less preferable. OK for trunk?
>
> I think it's possible to improve test coverage for NVPTX by running all OpenMP
> testcases via nvptx-none-run (it'll need some changes, but shouldn't be hard).
>
> gcc/
> * omp-low.c (lower_lastprivate_clauses): Correct handling of linear and
> lastprivate clauses in SIMT case.
>
> libgomp/
> * testsuite/libgomp.c/target-36.c: New testcase.
Ok for trunk/gcc-7-branch, thanks.
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.c/target-36.c
> @@ -0,0 +1,18 @@
> +int
> +main ()
> +{
> + int ah, bh, n = 1024;
> +#pragma omp target map(from: ah, bh)
> + {
> + int a, b;
> +#pragma omp simd lastprivate(b)
> + for (a = 0; a < n; a++)
> + {
> + b = a + n + 1;
> + asm volatile ("" : "+r"(b));
> + }
> + ah = a, bh = b;
> + }
> + if (ah != n || bh != 2 * n)
> + __builtin_abort ();
> +}
Would be nice to also test explicit linear, perhaps in the same testcase,
just add ch and c and use say linear(c:2).
Jakub
next prev parent reply other threads:[~2017-04-20 15:16 UTC|newest]
Thread overview: 27+ messages / expand[flat|nested] mbox.gz Atom feed top
2017-03-22 15:46 [PATCHv2 0/5] OpenMP/PTX: improve correctness in SIMD regions Alexander Monakov
2017-03-22 15:46 ` [PATCH 1/5] nvptx: implement SIMT enter/exit insns Alexander Monakov
2017-03-27 11:12 ` Alexander Monakov
2017-03-27 15:03 ` Bernd Schmidt
2017-03-22 15:46 ` [PATCH 3/5] omp-offload: implement SIMT privatization, part 2 Alexander Monakov
2017-03-23 10:37 ` Jakub Jelinek
2017-03-23 10:53 ` Alexander Monakov
2017-03-23 11:19 ` Jakub Jelinek
2017-03-22 15:46 ` [PATCH 4/5] tree-inline: implement SIMT privatization, part 3 Alexander Monakov
2017-03-23 10:47 ` Jakub Jelinek
2017-03-23 11:13 ` Alexander Monakov
2017-03-23 11:25 ` Jakub Jelinek
2017-03-23 16:15 ` Alexander Monakov
2017-03-23 16:23 ` Jakub Jelinek
2017-03-23 17:02 ` Alexander Monakov
2017-03-23 17:09 ` Jakub Jelinek
2017-03-22 15:46 ` [PATCH 2/5] omp-low: implement SIMT privatization, part 1 Alexander Monakov
2017-03-23 10:32 ` Jakub Jelinek
2017-03-31 16:05 ` Alexander Monakov
2017-04-07 7:58 ` [PATCH] omp-low: fix lastprivate/linear lowering for SIMT Alexander Monakov
2017-04-20 12:55 ` Alexander Monakov
2017-04-20 15:32 ` Jakub Jelinek [this message]
2017-04-20 16:48 ` [PATCH 2/5] omp-low: implement SIMT privatization, part 1 Alexander Monakov
2017-04-20 17:29 ` Jakub Jelinek
2017-03-22 15:46 ` [PATCH 5/5] address-taken: optimize SIMT privatized variables Alexander Monakov
2017-03-23 10:48 ` Jakub Jelinek
2017-03-31 10:22 ` [PATCHv2 0/5] OpenMP/PTX: improve correctness in SIMD regions 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=20170420151622.GV1809@tucnak \
--to=jakub@redhat.com \
--cc=amonakov@ispras.ru \
--cc=gcc-patches@gcc.gnu.org \
/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).