From 4912a04f8b35fadf65973bffc7037432ff7b7980 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Tue, 21 Apr 2020 14:16:24 +0200 Subject: [PATCH] [gcn] Use 'radeon' for the environment variable 'ACC_DEVICE_TYPE' ..., per OpenACC 3.0, A.1.2. "AMD GPU Targets". This complements commit 6687d13a87c42dddc7d1c7adade38d31ba0d1401 "Rename acc_device_gcn to acc_device_radeon". libgomp/ * 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. --- libgomp/ChangeLog | 12 ++++++++++++ libgomp/oacc-init.c | 4 +++- libgomp/testsuite/lib/libgomp.exp | 16 ++++++++-------- libgomp/testsuite/libgomp.oacc-c++/c++.exp | 18 +++++++++--------- .../libgomp.oacc-c++/firstprivate-mappings-1.C | 2 +- .../acc_get_property-gcn.c | 2 +- .../asyncwait-nop-1.c | 2 +- .../firstprivate-mappings-1.c | 2 +- .../function-not-offloaded.c | 4 ++-- .../libgomp.oacc-c-c++-common/loop-auto-1.c | 2 +- .../loop-dim-default.c | 2 +- .../libgomp.oacc-c-c++-common/routine-wv-2.c | 2 +- .../libgomp.oacc-c-c++-common/tile-1.c | 2 +- libgomp/testsuite/libgomp.oacc-c/c.exp | 18 +++++++++--------- .../libgomp.oacc-fortran/error_stop-1.f | 2 +- .../libgomp.oacc-fortran/error_stop-2.f | 2 +- .../libgomp.oacc-fortran/error_stop-3.f | 2 +- .../testsuite/libgomp.oacc-fortran/fortran.exp | 14 +++++++------- 18 files changed, 61 insertions(+), 47 deletions(-) diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 53bb8d23fa1..cfe6e0653c9 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,5 +1,17 @@ 2020-04-29 Thomas Schwinge + * 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 ef12b4c16d0..5d786a5a2e7 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 e7ce784314d..ee5f0e57b19 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 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 c06c2a097e3..7200ec19c47 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 7b3e670073c..b046bf2912d 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 ce59264a60d..4b1fb5e0e76 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 #include 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 754e015a280..7496426c8fa 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 253f8bf0bd0..2cdd2d1525e 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 517004a562d..64f8ab812a6 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 0c9ae957460..0273c2bddd7 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 30f0539707f..ca771646655 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 609f9f6a7da..9769ee72430 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 #include -#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 c019fe55c7a..5757917126c 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 7f13242fd59..48cbc980731 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 e7358f4f20d..a3f54d57bc3 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 fca1d960f66..5d5d20d1bc5 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 2ae0b0d1602..edb063b182b 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 60f0889a07c..d6079032505 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 -- 2.30.2