Skip to content

R
riscv-gcc-1
Commit 26b74ed0 by Andrew Stubbs Committed by Andrew Stubbs
Options

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
parent faab8a70
Showing with 38 additions and 2 deletions
libgomp/ChangeLog
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
......
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);
......
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);
......
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);
......
libgomp/testsuite/libgomp.oacc-c-c++-common/async_queue-1.c
/* { dg-do run { target openacc_nvidia_accel_selected } } */
/* Test mapping of async values to specific underlying queues. */
#undef NDEBUG
......
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
......
libgomp/testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c
/* { 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++;
}
......
libgomp/testsuite/libgomp.oacc-c-c++-common/tile-1.c
/* 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>
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment