From b0f9199a17911966ee24ec27b23bfb7ed7846700 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge 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 +#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 +#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