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(-)
@@ -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;
@@ -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
@@ -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)"
}
@@ -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
@@ -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>
@@ -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;
@@ -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
@@ -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++;
}
@@ -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" } */
@@ -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. */
@@ -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
@@ -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" } */
@@ -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)"
}
@@ -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 } }
@@ -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 } }
@@ -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 } }
@@ -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