065e25f633
When running test-case libgomp.oacc-c-c++-common/vector-length-128-7.c on an
RTX A2000 (sm_86) with driver 510.60.02 I run into:
...
FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-7.c \
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 \
output pattern test
...
The failing check verifies the launch dimensions:
...
/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: \
launch gangs=1, workers=8, vectors=128" } */
...
which fails because (as we can see with GOMP_DEBUG=1) the actual num_workers
is 6:
...
nvptx_exec: kernel main$_omp_fn$0: launch gangs=1, workers=6, vectors=128
...
This is due to the result of cuOccupancyMaxPotentialBlockSize (which suggests
'a launch configuration with reasonable occupancy') printed just before:
...
cuOccupancyMaxPotentialBlockSize: grid = 52, block = 768
...
[ Note: 6 * 128 == 768. ]
Fix this by updating the check to allow num_workers in the range 1 to 8.
Tested on x86_64 with nvptx accelerator.
libgomp/ChangeLog:
2022-04-01 Tom de Vries <tdevries@suse.de>
* testsuite/libgomp.oacc-c-c++-common/vector-length-128-7.c: Fix
num_workers check.
40 lines
999 B
C
40 lines
999 B
C
/* { dg-do run { target openacc_nvidia_accel_selected } } */
|
|
/* { dg-additional-options "-foffload=-fdump-tree-oaccloops" } */
|
|
/* { 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\\)" "oaccloops" } } */
|
|
/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=\[1-8\], vectors=128" } */
|