This patch adds GCN special casing for most of the OpenACC libgomp tests
that require it. It also disables one testcase that explicitly uses CUDA.
OK to commit?
Andrew
Update OpenACC tests for amdgcn
2019-11-19 Andrew Stubbs <a...@codesourcery.com>
libgomp/
* testsuite/libgomp.oacc-c-c++-common/acc_prof-init-1.c: Handle gcn.
* testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/acc_prof-parallel-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/tile-1.c: Likewise.
* testsuite/libgomp.oacc-fortran/error_stop-1.f: Likewise.
* testsuite/libgomp.oacc-fortran/error_stop-2.f: Likewise.
* testsuite/libgomp.oacc-fortran/error_stop-3.f: Likewise.
* testsuite/libgomp.oacc-c-c++-common/async_queue-1.c: Disable on GCN.
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-init-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-init-1.c
index b356feb8108..e82a03e8f3c 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-init-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-init-1.c
@@ -224,6 +224,8 @@ static void cb_compute_construct_end (acc_prof_info *prof_info, acc_event_info *
if (acc_device_type == acc_device_host)
assert (api_info->device_api == acc_device_api_none);
+ else if (acc_device_type == acc_device_gcn)
+ assert (api_info->device_api == acc_device_api_other);
else
assert (api_info->device_api == acc_device_api_cuda);
assert (api_info->valid_bytes == _ACC_API_INFO_VALID_BYTES);
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c
index 7cfc364e411..ddf647cda9b 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c
@@ -106,6 +106,8 @@ static void cb_enqueue_launch_start (acc_prof_info *prof_info, acc_event_info *e
assert (event_info->launch_event.vector_length >= 1);
else if (acc_device_type == acc_device_nvidia) /* ... is special. */
assert (event_info->launch_event.vector_length == 32);
+ else if (acc_device_type == acc_device_gcn) /* ...and so is this. */
+ assert (event_info->launch_event.vector_length == 64);
else
{
#ifdef __OPTIMIZE__
@@ -118,6 +120,8 @@ static void cb_enqueue_launch_start (acc_prof_info *prof_info, acc_event_info *e
if (acc_device_type == acc_device_host)
assert (api_info->device_api == acc_device_api_none);
+ else if (acc_device_type == acc_device_gcn)
+ assert (api_info->device_api == acc_device_api_other);
else
assert (api_info->device_api == acc_device_api_cuda);
assert (api_info->valid_bytes == _ACC_API_INFO_VALID_BYTES);
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-parallel-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-parallel-1.c
index ac6eb48cbbe..dc7c7582ce2 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-parallel-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-parallel-1.c
@@ -265,6 +265,8 @@ static void cb_enter_data_end (acc_prof_info *prof_info, acc_event_info *event_i
if (acc_device_type == acc_device_host)
assert (api_info->device_api == acc_device_api_none);
+ else if (acc_device_type == acc_device_gcn)
+ assert (api_info->device_api == acc_device_api_other);
else
assert (api_info->device_api == acc_device_api_cuda);
assert (api_info->valid_bytes == _ACC_API_INFO_VALID_BYTES);
@@ -319,6 +321,8 @@ static void cb_exit_data_start (acc_prof_info *prof_info, acc_event_info *event_
if (acc_device_type == acc_device_host)
assert (api_info->device_api == acc_device_api_none);
+ else if (acc_device_type == acc_device_gcn)
+ assert (api_info->device_api == acc_device_api_other);
else
assert (api_info->device_api == acc_device_api_cuda);
assert (api_info->valid_bytes == _ACC_API_INFO_VALID_BYTES);
@@ -371,6 +375,8 @@ static void cb_exit_data_end (acc_prof_info *prof_info, acc_event_info *event_in
if (acc_device_type == acc_device_host)
assert (api_info->device_api == acc_device_api_none);
+ else if (acc_device_type == acc_device_gcn)
+ assert (api_info->device_api == acc_device_api_other);
else
assert (api_info->device_api == acc_device_api_cuda);
assert (api_info->valid_bytes == _ACC_API_INFO_VALID_BYTES);
@@ -510,6 +516,8 @@ static void cb_compute_construct_end (acc_prof_info *prof_info, acc_event_info *
if (acc_device_type == acc_device_host)
assert (api_info->device_api == acc_device_api_none);
+ else if (acc_device_type == acc_device_gcn)
+ assert (api_info->device_api == acc_device_api_other);
else
assert (api_info->device_api == acc_device_api_cuda);
assert (api_info->valid_bytes == _ACC_API_INFO_VALID_BYTES);
@@ -573,6 +581,8 @@ static void cb_enqueue_launch_start (acc_prof_info *prof_info, acc_event_info *e
if (acc_device_type == acc_device_host)
assert (api_info->device_api == acc_device_api_none);
+ else if (acc_device_type == acc_device_gcn)
+ assert (api_info->device_api == acc_device_api_other);
else
assert (api_info->device_api == acc_device_api_cuda);
assert (api_info->valid_bytes == _ACC_API_INFO_VALID_BYTES);
@@ -637,6 +647,8 @@ static void cb_enqueue_launch_end (acc_prof_info *prof_info, acc_event_info *eve
if (acc_device_type == acc_device_host)
assert (api_info->device_api == acc_device_api_none);
+ else if (acc_device_type == acc_device_gcn)
+ assert (api_info->device_api == acc_device_api_other);
else
assert (api_info->device_api == acc_device_api_cuda);
assert (api_info->valid_bytes == _ACC_API_INFO_VALID_BYTES);
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/async_queue-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/async_queue-1.c
index 544b19fe663..4f9e53da85d 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/async_queue-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/async_queue-1.c
@@ -1,3 +1,5 @@
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
+
/* Test mapping of async values to specific underlying queues. */
#undef NDEBUG
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c
index 4ab67363ba6..840052fec12 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c
@@ -26,6 +26,8 @@ main ()
acc_device_t d;
#if defined ACC_DEVICE_TYPE_nvidia
d = acc_device_nvidia;
+#elif defined ACC_DEVICE_TYPE_gcn
+ d = acc_device_gcn;
#elif defined ACC_DEVICE_TYPE_host
d = acc_device_host;
#else
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c
index fdf4eb08f8a..517004a562d 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c
@@ -1,11 +1,11 @@
/* { dg-do link } */
-/* { dg-excess-errors "lto1, mkoffload and lto-wrapper fatal errors" { target openacc_nvidia_accel_selected } } */
+/* { dg-excess-errors "lto1, mkoffload and lto-wrapper fatal errors" { target { openacc_nvidia_accel_selected || openacc_amdgcn_accel_selected } } } */
int var;
#pragma acc declare create (var)
void __attribute__((noinline, noclone))
-foo () /* { dg-error "function 'foo' has been referenced in offloaded code but hasn't been marked to be included in the offloaded code" "" { target openacc_nvidia_accel_selected } } */
+foo () /* { dg-error "function 'foo' has been referenced in offloaded code but hasn't been marked to be included in the offloaded code" "" { target { openacc_nvidia_accel_selected || openacc_amdgcn_accel_selected } } } */
{
var++;
}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/tile-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/tile-1.c
index 5130591dd81..076e3cd75fe 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/tile-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/tile-1.c
@@ -1,3 +1,6 @@
+/* AMD GCN does not use 32-lane vectors.
+ { dg-skip-if "unsuitable dimensions" { openacc_amdgcn_accel_selected } { "*" } { "" } } */
+
/* { dg-additional-options "-fopenacc-dim=32" } */
#include <stdio.h>
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/error_stop-1.f b/libgomp/testsuite/libgomp.oacc-fortran/error_stop-1.f
index 4965e674c27..95810a6ae93 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/error_stop-1.f
+++ b/libgomp/testsuite/libgomp.oacc-fortran/error_stop-1.f
@@ -15,6 +15,6 @@
! { dg-output "ERROR STOP (\n|\r\n|\r)+" }
! PR85463. The "minimal" libgfortran implementation used with nvptx
! offloading is a little bit different.
-! { dg-output "Error termination.*" { target { ! openacc_nvidia_accel_selected } } }
+! { dg-output "Error termination.*" { target { { ! openacc_nvidia_accel_selected } && { ! openacc_amdgcn_accel_selected } } } }
! { dg-output "libgomp: cuStreamSynchronize error.*" { target openacc_nvidia_accel_selected } }
! { dg-shouldfail "" }
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/error_stop-2.f b/libgomp/testsuite/libgomp.oacc-fortran/error_stop-2.f
index 7103fdb5d8e..ce59bbda3c3 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/error_stop-2.f
+++ b/libgomp/testsuite/libgomp.oacc-fortran/error_stop-2.f
@@ -15,6 +15,6 @@
! { dg-output "ERROR STOP 35(\n|\r\n|\r)+" }
! PR85463. The "minimal" libgfortran implementation used with nvptx
! offloading is a little bit different.
-! { dg-output "Error termination.*" { target { ! openacc_nvidia_accel_selected } } }
+! { dg-output "Error termination.*" { target { { ! openacc_nvidia_accel_selected } && { ! openacc_amdgcn_accel_selected } } } }
! { dg-output "libgomp: cuStreamSynchronize error.*" { target openacc_nvidia_accel_selected } }
! { dg-shouldfail "" }
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/error_stop-3.f b/libgomp/testsuite/libgomp.oacc-fortran/error_stop-3.f
index 9c217f14ea1..9b606c83ad9 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/error_stop-3.f
+++ b/libgomp/testsuite/libgomp.oacc-fortran/error_stop-3.f
@@ -15,6 +15,6 @@
! { dg-output "ERROR STOP SiGN(\n|\r\n|\r)+" }
! PR85463. The "minimal" libgfortran implementation used with nvptx
! offloading is a little bit different.
-! { dg-output "Error termination.*" { target { ! openacc_nvidia_accel_selected } } }
+! { dg-output "Error termination.*" { target { { ! openacc_nvidia_accel_selected } && { ! openacc_amdgcn_accel_selected } } } }
! { dg-output "libgomp: cuStreamSynchronize error.*" { target openacc_nvidia_accel_selected } }
! { dg-shouldfail "" }
