Hi Tom!
On 2020-10-30T17:32:56+0100, Tom de Vries <tdevr...@suse.de> wrote:
> On 10/30/20 5:16 PM, Thomas Schwinge wrote:
>> OK to simplify and enhance the
>> testcases as attached, "Simplify and enhance
>> 'libgomp.oacc-c-c++-common/pr85486*.c' [PR85486]"?
>
> Yep, looks good.
As posted, pushed "Simplify and enhance
'libgomp.oacc-c-c++-common/pr85486*.c' [PR85486]" to master branch in
commit 79680c1d5cd3d89c2e7423e20dc8a6e1d6dc8151, and backported to
releases/gcc-10 branch in commit
28aaad48d5aafde3e5f269864ba934c602011328, releases/gcc-9 branch in commit
8860822a91e2e90a5eae726a478cd5ffc0d1fbfa.
Grüße
Thomas
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander
Walter
>From 79680c1d5cd3d89c2e7423e20dc8a6e1d6dc8151 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <tho...@codesourcery.com>
Date: Wed, 28 Oct 2020 10:56:20 +0100
Subject: [PATCH] Simplify and enhance 'libgomp.oacc-c-c++-common/pr85486*.c'
[PR85486]
Avoid code duplication, and better test what we expect to happen.
libgomp/
PR target/85486
* testsuite/libgomp.oacc-c-c++-common/pr85486-2.c: Simplify and enhance.
* testsuite/libgomp.oacc-c-c++-common/pr85486-3.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/pr85486.c: Likewise.
---
.../libgomp.oacc-c-c++-common/pr85486-2.c | 53 ++----------------
.../libgomp.oacc-c-c++-common/pr85486-3.c | 55 ++-----------------
.../libgomp.oacc-c-c++-common/pr85486.c | 9 ++-
3 files changed, 20 insertions(+), 97 deletions(-)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c
index f6ca263166d7..d45326488cd8 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c
@@ -1,52 +1,11 @@
/* { dg-do run { target openacc_nvidia_accel_selected } } */
+/* { dg-additional-options "-DVECTOR_LENGTH=" } */
/* { dg-additional-options "-fopenacc-dim=::128" } */
-/* Minimized from ref-1.C. */
+/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */
+/* { dg-set-target-env-var "GOMP_DEBUG" "1" } */
-#include <stdio.h>
+#include "pr85486.c"
-#pragma acc routine vector
-void __attribute__((noinline, noclone))
-Vector (int *ptr, int n, const int inc)
-{
- #pragma acc loop vector
- for (unsigned ix = 0; ix < n; ix++)
- ptr[ix] += inc;
-}
-
-int
-main (void)
-{
- const int n = 32, m=32;
-
- int ary[m][n];
- unsigned ix, iy;
-
- for (ix = m; ix--;)
- for (iy = n; iy--;)
- ary[ix][iy] = (1 << 16) + (ix << 8) + iy;
-
- int err = 0;
-
-#pragma acc parallel copy (ary)
- {
- Vector (&ary[0][0], m * n, (1 << 24) - (1 << 16));
- }
-
- for (ix = m; ix--;)
- for (iy = n; iy--;)
- if (ary[ix][iy] != ((1 << 24) + (ix << 8) + iy))
- {
- printf ("ary[%u][%u] = %x expected %x\n",
- ix, iy, ary[ix][iy], ((1 << 24) + (ix << 8) + iy));
- err++;
- }
-
- if (err)
- {
- printf ("%d failed\n", err);
- return 1;
- }
-
- return 0;
-}
+/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccdevlow" } } */
+/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=1, vectors=32" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-3.c
index a959b90c29ad..33480a4ae682 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-3.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-3.c
@@ -1,54 +1,11 @@
/* { dg-do run { target openacc_nvidia_accel_selected } } */
+/* { dg-additional-options "-DVECTOR_LENGTH=" } */
/* { dg-set-target-env-var "GOMP_OPENACC_DIM" "::128" } */
-/* Minimized from ref-1.C. */
+/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */
+/* { dg-set-target-env-var "GOMP_DEBUG" "1" } */
-#include <stdio.h>
+#include "pr85486.c"
-#pragma acc routine vector
-void __attribute__((noinline, noclone))
-Vector (int *ptr, int n, const int inc)
-{
- #pragma acc loop vector
- for (unsigned ix = 0; ix < n; ix++)
- ptr[ix] += inc;
-}
-
-int
-main (void)
-{
- const int n = 32, m=32;
-
- int ary[m][n];
- unsigned ix, iy;
-
- for (ix = m; ix--;)
- for (iy = n; iy--;)
- ary[ix][iy] = (1 << 16) + (ix << 8) + iy;
-
- int err = 0;
-
-#pragma acc parallel copy (ary)
- {
- Vector (&ary[0][0], m * n, (1 << 24) - (1 << 16));
- }
-
- for (ix = m; ix--;)
- for (iy = n; iy--;)
- if (ary[ix][iy] != ((1 << 24) + (ix << 8) + iy))
- {
- printf ("ary[%u][%u] = %x expected %x\n",
- ix, iy, ary[ix][iy], ((1 << 24) + (ix << 8) + iy));
- err++;
- }
-
- if (err)
- {
- printf ("%d failed\n", err);
- return 1;
- }
-
- return 0;
-}
-
-/* { dg-prune-output "using vector_length \\(32\\), ignoring runtime setting" } */
+/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccdevlow" } } */
+/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=1, vectors=32" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486.c
index 99c08059d37c..0d98b82f9932 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486.c
@@ -1,4 +1,8 @@
/* { dg-do run { target openacc_nvidia_accel_selected } } */
+/* { dg-additional-options "-DVECTOR_LENGTH=vector_length(128)" } */
+
+/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */
+/* { dg-set-target-env-var "GOMP_DEBUG" "1" } */
/* Minimized from ref-1.C. */
@@ -27,7 +31,7 @@ main (void)
int err = 0;
-#pragma acc parallel copy (ary) vector_length (128) /* { dg-warning "using vector_length \\(32\\) due to call to vector-partitionable routine, ignoring 128" } */
+#pragma acc parallel copy (ary) VECTOR_LENGTH /* { dg-warning "using vector_length \\(32\\) due to call to vector-partitionable routine, ignoring 128" } */
{
Vector (&ary[0][0], m * n, (1 << 24) - (1 << 16));
}
@@ -49,3 +53,6 @@ main (void)
return 0;
}
+
+/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccdevlow" } } */
+/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=1, vectors=32" } */
--
2.17.1
>From 28aaad48d5aafde3e5f269864ba934c602011328 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <tho...@codesourcery.com>
Date: Wed, 28 Oct 2020 10:56:20 +0100
Subject: [PATCH] Simplify and enhance 'libgomp.oacc-c-c++-common/pr85486*.c'
[PR85486]
Avoid code duplication, and better test what we expect to happen.
libgomp/
PR target/85486
* testsuite/libgomp.oacc-c-c++-common/pr85486-2.c: Simplify and enhance.
* testsuite/libgomp.oacc-c-c++-common/pr85486-3.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/pr85486.c: Likewise.
(cherry picked from commit 79680c1d5cd3d89c2e7423e20dc8a6e1d6dc8151)
---
.../libgomp.oacc-c-c++-common/pr85486-2.c | 53 ++----------------
.../libgomp.oacc-c-c++-common/pr85486-3.c | 55 ++-----------------
.../libgomp.oacc-c-c++-common/pr85486.c | 9 ++-
3 files changed, 20 insertions(+), 97 deletions(-)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c
index f6ca263166d7..d45326488cd8 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c
@@ -1,52 +1,11 @@
/* { dg-do run { target openacc_nvidia_accel_selected } } */
+/* { dg-additional-options "-DVECTOR_LENGTH=" } */
/* { dg-additional-options "-fopenacc-dim=::128" } */
-/* Minimized from ref-1.C. */
+/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */
+/* { dg-set-target-env-var "GOMP_DEBUG" "1" } */
-#include <stdio.h>
+#include "pr85486.c"
-#pragma acc routine vector
-void __attribute__((noinline, noclone))
-Vector (int *ptr, int n, const int inc)
-{
- #pragma acc loop vector
- for (unsigned ix = 0; ix < n; ix++)
- ptr[ix] += inc;
-}
-
-int
-main (void)
-{
- const int n = 32, m=32;
-
- int ary[m][n];
- unsigned ix, iy;
-
- for (ix = m; ix--;)
- for (iy = n; iy--;)
- ary[ix][iy] = (1 << 16) + (ix << 8) + iy;
-
- int err = 0;
-
-#pragma acc parallel copy (ary)
- {
- Vector (&ary[0][0], m * n, (1 << 24) - (1 << 16));
- }
-
- for (ix = m; ix--;)
- for (iy = n; iy--;)
- if (ary[ix][iy] != ((1 << 24) + (ix << 8) + iy))
- {
- printf ("ary[%u][%u] = %x expected %x\n",
- ix, iy, ary[ix][iy], ((1 << 24) + (ix << 8) + iy));
- err++;
- }
-
- if (err)
- {
- printf ("%d failed\n", err);
- return 1;
- }
-
- return 0;
-}
+/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccdevlow" } } */
+/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=1, vectors=32" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-3.c
index a959b90c29ad..33480a4ae682 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-3.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-3.c
@@ -1,54 +1,11 @@
/* { dg-do run { target openacc_nvidia_accel_selected } } */
+/* { dg-additional-options "-DVECTOR_LENGTH=" } */
/* { dg-set-target-env-var "GOMP_OPENACC_DIM" "::128" } */
-/* Minimized from ref-1.C. */
+/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */
+/* { dg-set-target-env-var "GOMP_DEBUG" "1" } */
-#include <stdio.h>
+#include "pr85486.c"
-#pragma acc routine vector
-void __attribute__((noinline, noclone))
-Vector (int *ptr, int n, const int inc)
-{
- #pragma acc loop vector
- for (unsigned ix = 0; ix < n; ix++)
- ptr[ix] += inc;
-}
-
-int
-main (void)
-{
- const int n = 32, m=32;
-
- int ary[m][n];
- unsigned ix, iy;
-
- for (ix = m; ix--;)
- for (iy = n; iy--;)
- ary[ix][iy] = (1 << 16) + (ix << 8) + iy;
-
- int err = 0;
-
-#pragma acc parallel copy (ary)
- {
- Vector (&ary[0][0], m * n, (1 << 24) - (1 << 16));
- }
-
- for (ix = m; ix--;)
- for (iy = n; iy--;)
- if (ary[ix][iy] != ((1 << 24) + (ix << 8) + iy))
- {
- printf ("ary[%u][%u] = %x expected %x\n",
- ix, iy, ary[ix][iy], ((1 << 24) + (ix << 8) + iy));
- err++;
- }
-
- if (err)
- {
- printf ("%d failed\n", err);
- return 1;
- }
-
- return 0;
-}
-
-/* { dg-prune-output "using vector_length \\(32\\), ignoring runtime setting" } */
+/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccdevlow" } } */
+/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=1, vectors=32" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486.c
index 99c08059d37c..0d98b82f9932 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486.c
@@ -1,4 +1,8 @@
/* { dg-do run { target openacc_nvidia_accel_selected } } */
+/* { dg-additional-options "-DVECTOR_LENGTH=vector_length(128)" } */
+
+/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */
+/* { dg-set-target-env-var "GOMP_DEBUG" "1" } */
/* Minimized from ref-1.C. */
@@ -27,7 +31,7 @@ main (void)
int err = 0;
-#pragma acc parallel copy (ary) vector_length (128) /* { dg-warning "using vector_length \\(32\\) due to call to vector-partitionable routine, ignoring 128" } */
+#pragma acc parallel copy (ary) VECTOR_LENGTH /* { dg-warning "using vector_length \\(32\\) due to call to vector-partitionable routine, ignoring 128" } */
{
Vector (&ary[0][0], m * n, (1 << 24) - (1 << 16));
}
@@ -49,3 +53,6 @@ main (void)
return 0;
}
+
+/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccdevlow" } } */
+/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=1, vectors=32" } */
--
2.17.1
>From 8860822a91e2e90a5eae726a478cd5ffc0d1fbfa Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <tho...@codesourcery.com>
Date: Wed, 28 Oct 2020 10:56:20 +0100
Subject: [PATCH] Simplify and enhance 'libgomp.oacc-c-c++-common/pr85486*.c'
[PR85486]
Avoid code duplication, and better test what we expect to happen.
libgomp/
PR target/85486
* testsuite/libgomp.oacc-c-c++-common/pr85486-2.c: Simplify and enhance.
* testsuite/libgomp.oacc-c-c++-common/pr85486-3.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/pr85486.c: Likewise.
(cherry picked from commit 79680c1d5cd3d89c2e7423e20dc8a6e1d6dc8151)
---
.../libgomp.oacc-c-c++-common/pr85486-2.c | 53 ++----------------
.../libgomp.oacc-c-c++-common/pr85486-3.c | 55 ++-----------------
.../libgomp.oacc-c-c++-common/pr85486.c | 9 ++-
3 files changed, 20 insertions(+), 97 deletions(-)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c
index f6ca263166d7..d45326488cd8 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c
@@ -1,52 +1,11 @@
/* { dg-do run { target openacc_nvidia_accel_selected } } */
+/* { dg-additional-options "-DVECTOR_LENGTH=" } */
/* { dg-additional-options "-fopenacc-dim=::128" } */
-/* Minimized from ref-1.C. */
+/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */
+/* { dg-set-target-env-var "GOMP_DEBUG" "1" } */
-#include <stdio.h>
+#include "pr85486.c"
-#pragma acc routine vector
-void __attribute__((noinline, noclone))
-Vector (int *ptr, int n, const int inc)
-{
- #pragma acc loop vector
- for (unsigned ix = 0; ix < n; ix++)
- ptr[ix] += inc;
-}
-
-int
-main (void)
-{
- const int n = 32, m=32;
-
- int ary[m][n];
- unsigned ix, iy;
-
- for (ix = m; ix--;)
- for (iy = n; iy--;)
- ary[ix][iy] = (1 << 16) + (ix << 8) + iy;
-
- int err = 0;
-
-#pragma acc parallel copy (ary)
- {
- Vector (&ary[0][0], m * n, (1 << 24) - (1 << 16));
- }
-
- for (ix = m; ix--;)
- for (iy = n; iy--;)
- if (ary[ix][iy] != ((1 << 24) + (ix << 8) + iy))
- {
- printf ("ary[%u][%u] = %x expected %x\n",
- ix, iy, ary[ix][iy], ((1 << 24) + (ix << 8) + iy));
- err++;
- }
-
- if (err)
- {
- printf ("%d failed\n", err);
- return 1;
- }
-
- return 0;
-}
+/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccdevlow" } } */
+/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=1, vectors=32" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-3.c
index a959b90c29ad..33480a4ae682 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-3.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-3.c
@@ -1,54 +1,11 @@
/* { dg-do run { target openacc_nvidia_accel_selected } } */
+/* { dg-additional-options "-DVECTOR_LENGTH=" } */
/* { dg-set-target-env-var "GOMP_OPENACC_DIM" "::128" } */
-/* Minimized from ref-1.C. */
+/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */
+/* { dg-set-target-env-var "GOMP_DEBUG" "1" } */
-#include <stdio.h>
+#include "pr85486.c"
-#pragma acc routine vector
-void __attribute__((noinline, noclone))
-Vector (int *ptr, int n, const int inc)
-{
- #pragma acc loop vector
- for (unsigned ix = 0; ix < n; ix++)
- ptr[ix] += inc;
-}
-
-int
-main (void)
-{
- const int n = 32, m=32;
-
- int ary[m][n];
- unsigned ix, iy;
-
- for (ix = m; ix--;)
- for (iy = n; iy--;)
- ary[ix][iy] = (1 << 16) + (ix << 8) + iy;
-
- int err = 0;
-
-#pragma acc parallel copy (ary)
- {
- Vector (&ary[0][0], m * n, (1 << 24) - (1 << 16));
- }
-
- for (ix = m; ix--;)
- for (iy = n; iy--;)
- if (ary[ix][iy] != ((1 << 24) + (ix << 8) + iy))
- {
- printf ("ary[%u][%u] = %x expected %x\n",
- ix, iy, ary[ix][iy], ((1 << 24) + (ix << 8) + iy));
- err++;
- }
-
- if (err)
- {
- printf ("%d failed\n", err);
- return 1;
- }
-
- return 0;
-}
-
-/* { dg-prune-output "using vector_length \\(32\\), ignoring runtime setting" } */
+/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccdevlow" } } */
+/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=1, vectors=32" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486.c
index 99c08059d37c..0d98b82f9932 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486.c
@@ -1,4 +1,8 @@
/* { dg-do run { target openacc_nvidia_accel_selected } } */
+/* { dg-additional-options "-DVECTOR_LENGTH=vector_length(128)" } */
+
+/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */
+/* { dg-set-target-env-var "GOMP_DEBUG" "1" } */
/* Minimized from ref-1.C. */
@@ -27,7 +31,7 @@ main (void)
int err = 0;
-#pragma acc parallel copy (ary) vector_length (128) /* { dg-warning "using vector_length \\(32\\) due to call to vector-partitionable routine, ignoring 128" } */
+#pragma acc parallel copy (ary) VECTOR_LENGTH /* { dg-warning "using vector_length \\(32\\) due to call to vector-partitionable routine, ignoring 128" } */
{
Vector (&ary[0][0], m * n, (1 << 24) - (1 << 16));
}
@@ -49,3 +53,6 @@ main (void)
return 0;
}
+
+/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccdevlow" } } */
+/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=1, vectors=32" } */
--
2.17.1
