From: Tom de Vries <Tom_deVries@mentor.com>
To: Cesar Philippidis <cesar@codesourcery.com>,
"gcc-patches@gcc.gnu.org" <gcc-patches@gcc.gnu.org>,
"Schwinge, Thomas" <Thomas_Schwinge@mentor.com>
Subject: Re: [og7] vector_length extension part 5: libgomp and tests
Date: Thu, 05 Apr 2018 16:36:00 -0000 [thread overview]
Message-ID: <b252156e-ac9f-cb4c-dfd6-1aeea4e4a85d@mentor.com> (raw)
In-Reply-To: <c024e89c-6bd8-2cc2-28a2-f5ca17b30aa8@codesourcery.com>
[-- Attachment #1: Type: text/plain, Size: 307 bytes --]
On 03/02/2018 09:47 PM, Cesar Philippidis wrote:
> libgomp/
> * plugin/plugin-nvptx.c (nvptx_exec): Adjust calculations of
> workers and vectors.
I wrote a test case that triggers this code, and added it to this code.
Build x86_64 with nvptx accelerator and tested libgomp.
Committed.
Thanks,
- Tom
[-- Attachment #2: 0004-nvptx-Handle-large-vectors-in-libgomp.patch --]
[-- Type: text/x-patch, Size: 3027 bytes --]
[nvptx] Handle large vectors in libgomp
2018-04-05 Cesar Philippidis <cesar@codesourcery.com>
Tom de Vries <tom@codesourcery.com>
* plugin/plugin-nvptx.c (nvptx_exec): Adjust calculations of
workers and vectors.
* testsuite/libgomp.oacc-c-c++-common/vector-length-128-7.c: New test.
---
libgomp/plugin/plugin-nvptx.c | 10 +++---
.../vector-length-128-7.c | 41 ++++++++++++++++++++++
2 files changed, 47 insertions(+), 4 deletions(-)
diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index bdc0c30..9b4768f 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -734,8 +734,6 @@ nvptx_exec (void (*fn), size_t mapnum, void **hostaddrs, void **devaddrs,
int threads_per_block = threads_per_sm > block_size
? block_size : threads_per_sm;
- threads_per_block /= warp_size;
-
if (threads_per_sm > cpu_size)
threads_per_sm = cpu_size;
@@ -802,6 +800,10 @@ nvptx_exec (void (*fn), size_t mapnum, void **hostaddrs, void **devaddrs,
if (seen_zero)
{
+ int vectors = dims[GOMP_DIM_VECTOR] > 0
+ ? dims[GOMP_DIM_VECTOR] : warp_size;
+ int workers = threads_per_block / vectors;
+
for (i = 0; i != GOMP_DIM_MAX; i++)
if (!dims[i])
{
@@ -819,10 +821,10 @@ nvptx_exec (void (*fn), size_t mapnum, void **hostaddrs, void **devaddrs,
: 2 * dev_size;
break;
case GOMP_DIM_WORKER:
- dims[i] = threads_per_block;
+ dims[i] = workers;
break;
case GOMP_DIM_VECTOR:
- dims[i] = warp_size;
+ dims[i] = vectors;
break;
default:
abort ();
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-7.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-7.c
new file mode 100644
index 0000000..60c264c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-7.c
@@ -0,0 +1,41 @@
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
+/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */
+/* { dg-additional-options "-foffload=-mlong-vector-in-workers" } */
+/* { dg-set-target-env-var "GOMP_DEBUG" "1" } */
+
+#include <stdlib.h>
+
+#define N 1024
+
+unsigned int a[N];
+unsigned int b[N];
+unsigned int c[N];
+unsigned int n = N;
+
+int
+main (void)
+{
+ for (unsigned int i = 0; i < n; ++i)
+ {
+ a[i] = i % 3;
+ b[i] = i % 5;
+ }
+
+#pragma acc parallel vector_length (128) copyin (a,b) copyout (c)
+ {
+#pragma acc loop worker
+ for (unsigned int i = 0; i < 4; i++)
+#pragma acc loop vector
+ for (unsigned int j = 0; j < n / 4; j++)
+ c[(i * N / 4) + j] = a[(i * N / 4) + j] + b[(i * N / 4) + j];
+ }
+
+ for (unsigned int i = 0; i < n; ++i)
+ if (c[i] != (i % 3) + (i % 5))
+ abort ();
+
+ return 0;
+}
+
+/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 0, 128\\)" "oaccdevlow" } } */
+/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=8, vectors=128" } */
next prev parent reply other threads:[~2018-04-05 16:36 UTC|newest]
Thread overview: 50+ messages / expand[flat|nested] mbox.gz Atom feed top
2018-03-01 21:17 [og7] vector_length extension part 1: generalize function and variable names Cesar Philippidis
2018-03-02 16:55 ` [og7] vector_length extension part 2: Generalize state propagation and synchronization Cesar Philippidis
2018-03-21 17:16 ` Tom de Vries
2018-03-22 8:05 ` Cesar Philippidis
2018-03-22 14:16 ` Tom de Vries
2018-03-22 14:35 ` Cesar Philippidis
2018-03-22 14:24 ` Tom de Vries
2018-03-22 15:18 ` Cesar Philippidis
2018-03-22 16:20 ` Tom de Vries
2018-03-22 17:26 ` Cesar Philippidis
2018-03-22 17:58 ` Tom de Vries
2018-03-22 19:32 ` Cesar Philippidis
2018-03-23 8:56 ` Tom de Vries
2018-03-23 14:35 ` Tom de Vries
2018-03-22 15:04 ` Tom de Vries
2018-03-22 17:14 ` Cesar Philippidis
2018-03-22 17:47 ` Tom de Vries
2018-03-22 17:48 ` Cesar Philippidis
2018-03-22 18:00 ` Tom de Vries
2018-03-23 13:14 ` Tom de Vries
2018-03-23 13:16 ` Tom de Vries
2018-03-23 14:18 ` Tom de Vries
2018-03-23 16:30 ` Tom de Vries
2018-03-30 1:50 ` Tom de Vries
2018-03-30 14:48 ` Tom de Vries
2018-03-30 15:06 ` Cesar Philippidis
2018-03-30 15:35 ` Tom de Vries
2018-04-05 16:33 ` Tom de Vries
2018-04-03 14:52 ` [nvptx] Use MAX, MIN, ROUND_UP macros Tom de Vries
2018-04-03 15:00 ` [og7] vector_length extension part 2: Generalize state propagation and synchronization Tom de Vries
2018-04-05 14:06 ` Tom de Vries
2018-04-05 14:14 ` Tom de Vries
2018-03-02 17:51 ` [og7] vector_length extension part 3: reductions Cesar Philippidis
2018-04-05 14:07 ` Tom de Vries
2018-04-05 16:26 ` Tom de Vries
2018-03-02 19:18 ` [og7] vector_length extension part 4: target hooks and automatic parallelism Cesar Philippidis
2018-03-21 15:55 ` Tom de Vries
2018-03-21 20:28 ` Cesar Philippidis
2018-03-26 14:25 ` Tom de Vries
2018-03-26 14:37 ` Cesar Philippidis
2018-03-26 16:52 ` Tom de Vries
2018-03-27 12:16 ` Tom de Vries
2018-03-26 17:13 ` Tom de Vries
2018-04-05 16:32 ` Tom de Vries
2018-03-02 20:47 ` [og7] vector_length extension part 5: libgomp and tests Cesar Philippidis
2018-03-16 13:50 ` Thomas Schwinge
2018-03-27 13:00 ` Tom de Vries
2018-04-05 16:36 ` Tom de Vries [this message]
2018-03-09 15:29 ` [og7] vector_length extension part 1: generalize function and variable names Thomas Schwinge
2018-03-09 15:31 ` Cesar Philippidis
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=b252156e-ac9f-cb4c-dfd6-1aeea4e4a85d@mentor.com \
--to=tom_devries@mentor.com \
--cc=Thomas_Schwinge@mentor.com \
--cc=cesar@codesourcery.com \
--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).