Hi!
On 2020-04-23T17:27:45+0100, Andrew Stubbs <a...@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 <tho...@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 <tho...@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
