diff options
18 files changed, 61 insertions, 47 deletions
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 53bb8d2..cfe6e06 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,5 +1,17 @@ 2020-04-29 Thomas Schwinge <thomas@codesourcery.com> + * oacc-init.c (get_openacc_name): Handle 'gcn'. + * testsuite/lib/libgomp.exp + (offload_target_to_openacc_device_type) [amdgcn*]: Return + 'radeon'. Adjust all users. + (check_effective_target_openacc_amdgcn_accel_present): Rename + to... + (check_effective_target_openacc_radeon_accel_present): ... this. + Adjust all users. + (check_effective_target_openacc_amdgcn_accel_selected): Rename to... + (check_effective_target_openacc_radeon_accel_selected): ... this. + Adjust all users. + * testsuite/libgomp.fortran/use_device_ptr-optional-2.f90: Add 'dg-do run'. diff --git a/libgomp/oacc-init.c b/libgomp/oacc-init.c index ef12b4c..5d786a5 100644 --- a/libgomp/oacc-init.c +++ b/libgomp/oacc-init.c @@ -99,7 +99,9 @@ unknown_device_type_error (acc_device_t invalid_type) static const char * get_openacc_name (const char *name) { - if (strcmp (name, "nvptx") == 0) + if (strcmp (name, "gcn") == 0) + return "radeon"; + else if (strcmp (name, "nvptx") == 0) return "nvidia"; else return name; diff --git a/libgomp/testsuite/lib/libgomp.exp b/libgomp/testsuite/lib/libgomp.exp index e7ce784..ee5f0e5 100644 --- a/libgomp/testsuite/lib/libgomp.exp +++ b/libgomp/testsuite/lib/libgomp.exp @@ -319,7 +319,7 @@ proc libgomp_option_proc { option } { proc offload_target_to_openacc_device_type { offload_target } { switch -glob $offload_target { amdgcn* { - return "gcn" + return "radeon" } disable { return "host" @@ -483,10 +483,10 @@ proc check_effective_target_hsa_offloading_selected {} { }] } -# Return 1 if at least one AMD GCN board is present. +# Return 1 if at least one AMD GPU is accessible. -proc check_effective_target_openacc_amdgcn_accel_present { } { - return [check_runtime openacc_amdgcn_accel_present { +proc check_effective_target_openacc_radeon_accel_present { } { + return [check_runtime openacc_radeon_accel_present { #include <openacc.h> int main () { return !(acc_get_num_devices (acc_device_radeon) > 0); @@ -494,11 +494,11 @@ proc check_effective_target_openacc_amdgcn_accel_present { } { } "" ] } -# Return 1 if at least one AMD GCN board is present, and the AMD GCN device -# type is selected by default. +# Return 1 if at least one AMD GPU is accessible, and the OpenACC 'radeon' +# device type is selected. -proc check_effective_target_openacc_amdgcn_accel_selected { } { - if { ![check_effective_target_openacc_amdgcn_accel_present] } { +proc check_effective_target_openacc_radeon_accel_selected { } { + if { ![check_effective_target_openacc_radeon_accel_present] } { return 0; } global offload_target diff --git a/libgomp/testsuite/libgomp.oacc-c++/c++.exp b/libgomp/testsuite/libgomp.oacc-c++/c++.exp index c06c2a0..7200ec1 100644 --- a/libgomp/testsuite/libgomp.oacc-c++/c++.exp +++ b/libgomp/testsuite/libgomp.oacc-c++/c++.exp @@ -88,15 +88,6 @@ if { $lang_test_file_found } { unsupported "$subdir $offload_target offloading" continue } - gcn { - if { ![check_effective_target_openacc_amdgcn_accel_present] } { - # Don't bother; execution testing is going to FAIL. - untested "$subdir $offload_target offloading: supported, but hardware not accessible" - continue - } - - set acc_mem_shared 0 - } host { set acc_mem_shared 1 } @@ -115,6 +106,15 @@ if { $lang_test_file_found } { set acc_mem_shared 0 } + radeon { + if { ![check_effective_target_openacc_radeon_accel_present] } { + # Don't bother; execution testing is going to FAIL. + untested "$subdir $offload_target offloading: supported, but hardware not accessible" + continue + } + + set acc_mem_shared 0 + } default { error "Unknown OpenACC device type: $openacc_device_type (offload target: $offload_target)" } diff --git a/libgomp/testsuite/libgomp.oacc-c++/firstprivate-mappings-1.C b/libgomp/testsuite/libgomp.oacc-c++/firstprivate-mappings-1.C index 7b3e670..b046bf2 100644 --- a/libgomp/testsuite/libgomp.oacc-c++/firstprivate-mappings-1.C +++ b/libgomp/testsuite/libgomp.oacc-c++/firstprivate-mappings-1.C @@ -3,7 +3,7 @@ /* PR middle-end/48591 */ /* PR other/71064 */ /* Set to 0 for offloading targets not supporting long double. */ -#if defined(ACC_DEVICE_TYPE_nvidia) || defined(ACC_DEVICE_TYPE_gcn) +#if defined(ACC_DEVICE_TYPE_nvidia) || defined(ACC_DEVICE_TYPE_radeon) # define DO_LONG_DOUBLE 0 #else # define DO_LONG_DOUBLE 1 diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_get_property-gcn.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_get_property-gcn.c index ce59264..4b1fb5e 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_get_property-gcn.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_get_property-gcn.c @@ -3,7 +3,7 @@ those obtained through the HSA API. */ /* { dg-additional-sources acc_get_property-aux.c } */ /* { dg-additional-options "-ldl" } */ -/* { dg-do run { target openacc_amdgcn_accel_selected } } */ +/* { dg-do run { target openacc_radeon_accel_selected } } */ #include <dlfcn.h> #include <stdint.h> 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 754e015..7496426 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,7 +26,7 @@ main () acc_device_t d; #if defined ACC_DEVICE_TYPE_nvidia d = acc_device_nvidia; -#elif defined ACC_DEVICE_TYPE_gcn +#elif defined ACC_DEVICE_TYPE_radeon d = acc_device_radeon; #elif defined ACC_DEVICE_TYPE_host d = acc_device_host; diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-mappings-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-mappings-1.c index 253f8bf..2cdd2d1 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-mappings-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-mappings-1.c @@ -6,7 +6,7 @@ /* PR middle-end/48591 */ /* PR other/71064 */ /* Set to 0 for offloading targets not supporting long double. */ -#if defined(ACC_DEVICE_TYPE_nvidia) || defined(ACC_DEVICE_TYPE_gcn) +#if defined(ACC_DEVICE_TYPE_nvidia) || defined(ACC_DEVICE_TYPE_radeon) # define DO_LONG_DOUBLE 0 #else # define DO_LONG_DOUBLE 1 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 517004a..64f8ab8 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 || openacc_amdgcn_accel_selected } } } */ +/* { dg-excess-errors "lto1, mkoffload and lto-wrapper fatal errors" { target { openacc_nvidia_accel_selected || openacc_radeon_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 || openacc_amdgcn_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_radeon_accel_selected } } } */ { var++; } diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c index 0c9ae95..0273c2b 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c @@ -1,5 +1,5 @@ /* AMD GCN does not use 32-lane vectors. - { dg-skip-if "unsuitable dimensions" { openacc_amdgcn_accel_selected } { "*" } { "" } } */ + { dg-skip-if "unsuitable dimensions" { openacc_radeon_accel_selected } { "*" } { "" } } */ /* { dg-additional-options "-fopenacc-dim=32" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-dim-default.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-dim-default.c index 30f0539..ca77164 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-dim-default.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-dim-default.c @@ -128,7 +128,7 @@ int test_1 (int gp, int wp, int vp) int main () { -#ifdef ACC_DEVICE_TYPE_gcn +#ifdef ACC_DEVICE_TYPE_radeon /* AMD GCN uses the autovectorizer for the vector dimension: the use of a function call in vector-partitioned code in this test is not currently supported. */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-2.c index 609f9f6..9769ee7 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-2.c @@ -2,7 +2,7 @@ #include <openacc.h> #include <gomp-constants.h> -#ifdef ACC_DEVICE_TYPE_gcn +#ifdef ACC_DEVICE_TYPE_radeon #define NUM_WORKERS 16 #define NUM_VECTORS 1 #else 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 c019fe5..57579171 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/tile-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/tile-1.c @@ -1,5 +1,5 @@ /* AMD GCN does not use 32-lane vectors, so the expected use counts mismatch. - { dg-skip-if "unsuitable dimensions" { openacc_amdgcn_accel_selected } { "*" } { "" } } */ + { dg-skip-if "unsuitable dimensions" { openacc_radeon_accel_selected } { "*" } { "" } } */ /* { dg-additional-options "-fopenacc-dim=32" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c/c.exp b/libgomp/testsuite/libgomp.oacc-c/c.exp index 7f13242..48cbc98 100644 --- a/libgomp/testsuite/libgomp.oacc-c/c.exp +++ b/libgomp/testsuite/libgomp.oacc-c/c.exp @@ -51,15 +51,6 @@ foreach offload_target [concat [split $offload_targets ","] "disable"] { unsupported "$subdir $offload_target offloading" continue } - gcn { - if { ![check_effective_target_openacc_amdgcn_accel_present] } { - # Don't bother; execution testing is going to FAIL. - untested "$subdir $offload_target offloading: supported, but hardware not accessible" - continue - } - - set acc_mem_shared 0 - } host { set acc_mem_shared 1 } @@ -78,6 +69,15 @@ foreach offload_target [concat [split $offload_targets ","] "disable"] { set acc_mem_shared 0 } + radeon { + if { ![check_effective_target_openacc_radeon_accel_present] } { + # Don't bother; execution testing is going to FAIL. + untested "$subdir $offload_target offloading: supported, but hardware not accessible" + continue + } + + set acc_mem_shared 0 + } default { error "Unknown OpenACC device type: $openacc_device_type (offload target: $offload_target)" } diff --git a/libgomp/testsuite/libgomp.oacc-fortran/error_stop-1.f b/libgomp/testsuite/libgomp.oacc-fortran/error_stop-1.f index e7358f4..a3f54d5 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/error_stop-1.f +++ b/libgomp/testsuite/libgomp.oacc-fortran/error_stop-1.f @@ -17,7 +17,7 @@ ! In gfortran's main program, libfortran's set_options is called - which sets ! compiler_options.backtrace = 1 by default. For an offload libgfortran, this ! is never called and, hence, "Error termination." is never printed. Thus: -! { dg-output "Error termination.*" { target { ! { openacc_nvidia_accel_selected || openacc_amdgcn_accel_selected } } } } +! { dg-output "Error termination.*" { target { ! { openacc_nvidia_accel_selected || openacc_radeon_accel_selected } } } } ! ! PR85463: ! { dg-output "libgomp: cuStreamSynchronize error.*" { target openacc_nvidia_accel_selected } } diff --git a/libgomp/testsuite/libgomp.oacc-fortran/error_stop-2.f b/libgomp/testsuite/libgomp.oacc-fortran/error_stop-2.f index fca1d96..5d5d20d 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/error_stop-2.f +++ b/libgomp/testsuite/libgomp.oacc-fortran/error_stop-2.f @@ -17,7 +17,7 @@ ! In gfortran's main program, libfortran's set_options is called - which sets ! compiler_options.backtrace = 1 by default. For an offload libgfortran, this ! is never called and, hence, "Error termination." is never printed. Thus: -! { dg-output "Error termination.*" { target { ! { openacc_nvidia_accel_selected || openacc_amdgcn_accel_selected } } } } +! { dg-output "Error termination.*" { target { ! { openacc_nvidia_accel_selected || openacc_radeon_accel_selected } } } } ! ! PR85463: ! { dg-output "libgomp: cuStreamSynchronize error.*" { target openacc_nvidia_accel_selected } } diff --git a/libgomp/testsuite/libgomp.oacc-fortran/error_stop-3.f b/libgomp/testsuite/libgomp.oacc-fortran/error_stop-3.f index 2ae0b0d..edb063b 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/error_stop-3.f +++ b/libgomp/testsuite/libgomp.oacc-fortran/error_stop-3.f @@ -17,7 +17,7 @@ ! In gfortran's main program, libfortran's set_options is called - which sets ! compiler_options.backtrace = 1 by default. For an offload libgfortran, this ! is never called and, hence, "Error termination." is never printed. Thus: -! { dg-output "Error termination.*" { target { ! { openacc_nvidia_accel_selected || openacc_amdgcn_accel_selected } } } } +! { dg-output "Error termination.*" { target { ! { openacc_nvidia_accel_selected || openacc_radeon_accel_selected } } } } ! ! PR85463: ! { dg-output "libgomp: cuStreamSynchronize error.*" { target openacc_nvidia_accel_selected } } diff --git a/libgomp/testsuite/libgomp.oacc-fortran/fortran.exp b/libgomp/testsuite/libgomp.oacc-fortran/fortran.exp index 60f0889..d607903 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/fortran.exp +++ b/libgomp/testsuite/libgomp.oacc-fortran/fortran.exp @@ -82,8 +82,11 @@ if { $lang_test_file_found } { unsupported "$subdir $offload_target offloading" continue } - gcn { - if { ![check_effective_target_openacc_amdgcn_accel_present] } { + host { + set acc_mem_shared 1 + } + nvidia { + if { ![check_effective_target_openacc_nvidia_accel_present] } { # Don't bother; execution testing is going to FAIL. untested "$subdir $offload_target offloading: supported, but hardware not accessible" continue @@ -91,11 +94,8 @@ if { $lang_test_file_found } { set acc_mem_shared 0 } - host { - set acc_mem_shared 1 - } - nvidia { - if { ![check_effective_target_openacc_nvidia_accel_present] } { + radeon { + if { ![check_effective_target_openacc_radeon_accel_present] } { # Don't bother; execution testing is going to FAIL. untested "$subdir $offload_target offloading: supported, but hardware not accessible" continue |