diff options
author | Andrew Stubbs <ams@codesourcery.com> | 2019-12-13 17:40:06 +0000 |
---|---|---|
committer | Andrew Stubbs <ams@gcc.gnu.org> | 2019-12-13 17:40:06 +0000 |
commit | 26b74ed0223d108d7d7818c3c860f20cfe81a4af (patch) | |
tree | 2284c2179d38d561b52796023312bb416d85b32d /libgomp | |
parent | faab8a70f2c40758c8bb15303098f3b824bafb60 (diff) | |
download | gcc-26b74ed0223d108d7d7818c3c860f20cfe81a4af.zip gcc-26b74ed0223d108d7d7818c3c860f20cfe81a4af.tar.gz gcc-26b74ed0223d108d7d7818c3c860f20cfe81a4af.tar.bz2 |
Update OpenACC tests for amdgcn
2019-12-13 Andrew Stubbs <ams@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/async_queue-1.c: Disable on GCN.
* testsuite/libgomp.oacc-c-c++-common/tile-1.c: Likewise.
From-SVN: r279378
Diffstat (limited to 'libgomp')
8 files changed, 38 insertions, 2 deletions
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 136ba6c..ce44039 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,14 @@ +2019-12-13 Andrew Stubbs <ams@codesourcery.com> + + * 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/async_queue-1.c: Disable on GCN. + * testsuite/libgomp.oacc-c-c++-common/tile-1.c: Likewise. + 2019-12-13 Tobias Burnus <tobias@codesourcery.com> * openacc.f90 (module openacc_kinds): Use 'PUBLIC' to mark all symbols 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 b356feb..e82a03e 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 7cfc364..ddf647c 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 ac6eb48..dc7c758 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 544b19f..4f9e53d 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 4ab6736..840052f 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 fdf4eb0..517004a 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 5130591..c019fe5 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, so the expected use counts mismatch. + { dg-skip-if "unsuitable dimensions" { openacc_amdgcn_accel_selected } { "*" } { "" } } */ + /* { dg-additional-options "-fopenacc-dim=32" } */ #include <stdio.h> |