[AMD,GCN] Use 'radeon' for the environment variable 'ACC_DEVICE_TYPE' (was: [committed, amdgcn/openacc] Rename acc_device_gcn to acc_device_radeon)

Message ID yxfpd081hvqp.fsf@hertz.schwinge.homeip.net
State New
Headers show
Series
  • [AMD,GCN] Use 'radeon' for the environment variable 'ACC_DEVICE_TYPE' (was: [committed, amdgcn/openacc] Rename acc_device_gcn to acc_device_radeon)
Related show

Commit Message

Thomas Schwinge April 21, 2020, 12:24 p.m.
Hi Andrew!

On 2020-01-17T18:18:41+0000, Andrew Stubbs <ams@codesourcery.com> wrote:
> I've committed this patch, pre-approved by Thomas. It's basically just a

> search and replace.

>

> I had originally invented acc_device_gcn for the device number, but as

> Thomas pointed out the OpenACC documentation already uses

> acc_device_radeon for AMD devices. Presumably this is what any existing

> code will use, if anything, so we ought to be compatible.

>

> There's no official release using the "wrong" name, so I don't believe

> we need to retain that name for any reason.


ACK.

> Rename acc_device_gcn to acc_device_radeon


I wondered whether for symmetry, the GCC-internal 'GOMP_DEVICE_GCN',
'OFFLOAD_TARGET_TYPE_GCN' should also be renamed to '*_RADEON'?  Or,
going by example of '*_NVIDIA_PTX', name them '*_AMD_GCN'.  Or, in fact
then leave them as '*_GCN', given Julian's point in another thread that
"CPU architectures or core implementations [sometimes shift] company
allegiance".  Thank you for listening to me thinking aloud.  ;-)

But more importantly, what about the user-visible 'gcn' in the
'ACC_DEVICE_TYPE' environment variable?  As I'd quoted in
<http://mid.mail-archive.com/87blsqlre6.fsf@euler.schwinge.homeip.net>:

| Per OpenACC 3.0, A.1.2. "AMD
| GPU Targets", for example, there is 'acc_device_radeon'

... which you've addressed...

| (and "the
| case-insensitive name 'radeon' for the environment variable
| 'ACC_DEVICE_TYPE'").

..., but this not yet?

Please see the attached "[AMD GCN] Use 'radeon' for the environment
variable 'ACC_DEVICE_TYPE'", completely untested.  Will you please test
and review that?  If approving this patch, please respond with
"Reviewed-by: NAME <EMAIL>" so that your effort will be recorded in the
commit log, see <https://gcc.gnu.org/wiki/Reviewed-by>.


Grüße
 Thomas


-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter

Comments

Andrew Stubbs April 23, 2020, 4:27 p.m. | #1
On 21/04/2020 13:24, Thomas Schwinge wrote:
> I wondered whether for symmetry, the GCC-internal 'GOMP_DEVICE_GCN',

> 'OFFLOAD_TARGET_TYPE_GCN' should also be renamed to '*_RADEON'?  Or,

> going by example of '*_NVIDIA_PTX', name them '*_AMD_GCN'.  Or, in fact

> then leave them as '*_GCN', given Julian's point in another thread that

> "CPU architectures or core implementations [sometimes shift] company

> allegiance".  Thank you for listening to me thinking aloud.  ;-)


I don't think the GCC internal names need to change. The port name is 
now "gcn", and the target name "amdgcn". I don't think we need a third 
name for it!

> But more importantly, what about the user-visible 'gcn' in the

> 'ACC_DEVICE_TYPE' environment variable?  As I'd quoted in

> <http://mid.mail-archive.com/87blsqlre6.fsf@euler.schwinge.homeip.net>:

> 

> | Per OpenACC 3.0, A.1.2. "AMD

> | GPU Targets", for example, there is 'acc_device_radeon'

> 

> ... which you've addressed...

> 

> | (and "the

> | case-insensitive name 'radeon' for the environment variable

> | 'ACC_DEVICE_TYPE'").

> 

> ..., but this not yet?


Oh, right, the user-visible case ought to accept "radeon", at least. :-(

> Please see the attached "[AMD GCN] Use 'radeon' for the environment

> variable 'ACC_DEVICE_TYPE'", completely untested.  Will you please test

> and review that?  If approving this patch, please respond with

> "Reviewed-by: NAME <EMAIL>" so that your effort will be recorded in the

> commit log, see <https://gcc.gnu.org/wiki/Reviewed-by>.


This is OK, assuming it tests OK.

Andrew
Thomas Schwinge April 29, 2020, 8:08 a.m. | #2
Hi!

On 2020-04-23T17:27:45+0100, Andrew Stubbs <ams@codesourcery.com> wrote:
> On 21/04/2020 13:24, Thomas Schwinge wrote:

>> [...], what about the user-visible 'gcn' in the

>> 'ACC_DEVICE_TYPE' environment variable?  As I'd quoted in

>> <http://mid.mail-archive.com/87blsqlre6.fsf@euler.schwinge.homeip.net>:

>>

>> | Per OpenACC 3.0, A.1.2. "AMD

>> | GPU Targets", for example, there is 'acc_device_radeon'

>>

>> ... which you've addressed...

>>

>> | (and "the

>> | case-insensitive name 'radeon' for the environment variable

>> | 'ACC_DEVICE_TYPE'").

>>

>> ..., but this not yet?

>

> Oh, right, the user-visible case ought to accept "radeon", at least. :-(

>

>> Please see the attached "[AMD GCN] Use 'radeon' for the environment

>> variable 'ACC_DEVICE_TYPE'", [...]

>

> This is OK, assuming it tests OK.


Thanks, tested, and now pushed to master branch in commit
4912a04f8b35fadf65973bffc7037432ff7b7980 "[gcn] Use 'radeon' for the
environment variable 'ACC_DEVICE_TYPE'", see attached.


Grüße
 Thomas


-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
From 4912a04f8b35fadf65973bffc7037432ff7b7980 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>

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 53bb8d23fa15..cfe6e0653c92 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 ef12b4c16d01..5d786a5a2e7c 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 e7ce784314d9..ee5f0e57b190 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 c06c2a097e38..7200ec19c476 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 7b3e670073c9..b046bf2912d6 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 ce59264a60dc..4b1fb5e0e761 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 754e015a280c..7496426c8fa3 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 253f8bf0bd02..2cdd2d1525e7 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 517004a562d2..64f8ab812a63 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 0c9ae9574601..0273c2bddd74 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 30f0539707f5..ca7716466558 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 609f9f6a7da2..9769ee724308 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 c019fe55c7a2..5757917126c1 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 7f13242fd59d..48cbc980731b 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 e7358f4f20df..a3f54d57bc35 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 fca1d960f66d..5d5d20d1bc53 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 2ae0b0d16028..edb063b182bc 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 60f0889a07cf..d60790325053 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.26.2

Patch

From d82df713bbb3401cc346bf07d61ad597d302d5a6 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Tue, 21 Apr 2020 14:16:24 +0200
Subject: [PATCH] [AMD 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                            |  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 +++++++-------
 17 files changed, 49 insertions(+), 47 deletions(-)

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 <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 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 <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 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 <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 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.17.1