Skip to content

R
riscv-gcc-1
Commit 4912a04f by Thomas Schwinge
Options

[gcn] Use 'radeon' for the environment variable 'ACC_DEVICE_TYPE' 

..., per OpenACC 3.0, A.1.2. "AMD GPU Targets".

This complements commit 6687d13a "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.
parent b9dc11b6
Showing with 61 additions and 47 deletions
libgomp/ChangeLog
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'.
......
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;
......
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
......
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)"
}
......
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
......
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>
......
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;
......
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
......
libgomp/testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c
/* { 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++;
}
......
libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c
/* 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" } */
......
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. */
......
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
......
libgomp/testsuite/libgomp.oacc-c-c++-common/tile-1.c
/* 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" } */
......
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)"
}
......
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 } }
......
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 } }
......
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 } }
......
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
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment