From 32099c0d24adb93a031e0301ffd77b065b6f5472 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Fri, 4 Jun 2021 15:31:53 +0200
Subject: [PATCH] Fix 'libgomp.oacc-fortran/parallel-dims.f90' for
 'acc_device_radeon'

..., by simplifying 'libgomp.oacc-c-c++-common/parallel-dims.c', and updating
the former correspondingly.  '__builtin_goacc_parlevel_id' does the right thing
for all 'acc_device_*'.

Follow-up to commit 09e0ad6253f4330977e1b2f116b5e289dc2c2a02 "Update OpenACC
tests for amdgcn".

	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Simplify.
	* testsuite/libgomp.oacc-fortran/parallel-dims-aux.c: Update.
---
 .../libgomp.oacc-c-c++-common/parallel-dims.c      | 32 ++++------------------
 .../libgomp.oacc-fortran/parallel-dims-aux.c       | 31 ++++-----------------
 2 files changed, 12 insertions(+), 51 deletions(-)

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c
index c7412c2..974e150 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c
@@ -10,42 +10,22 @@
 #include <openacc.h>
 #include <gomp-constants.h>
 
-/* TODO: "(int) acc_device_*" casts because of the C++ acc_on_device wrapper
-   not behaving as expected for -O0.  */
 #pragma acc routine seq
-static unsigned int __attribute__ ((optimize ("O2"))) acc_gang ()
+static int acc_gang ()
 {
-  if (acc_on_device ((int) acc_device_host))
-    return 0;
-  else if (acc_on_device ((int) acc_device_nvidia)
-	   || acc_on_device ((int) acc_device_radeon))
-    return __builtin_goacc_parlevel_id (GOMP_DIM_GANG);
-  else
-    __builtin_abort ();
+  return __builtin_goacc_parlevel_id (GOMP_DIM_GANG);
 }
 
 #pragma acc routine seq
-static unsigned int __attribute__ ((optimize ("O2"))) acc_worker ()
+static int acc_worker ()
 {
-  if (acc_on_device ((int) acc_device_host))
-    return 0;
-  else if (acc_on_device ((int) acc_device_nvidia)
-	   || acc_on_device ((int) acc_device_radeon))
-    return __builtin_goacc_parlevel_id (GOMP_DIM_WORKER);
-  else
-    __builtin_abort ();
+  return __builtin_goacc_parlevel_id (GOMP_DIM_WORKER);
 }
 
 #pragma acc routine seq
-static unsigned int __attribute__ ((optimize ("O2"))) acc_vector ()
+static int acc_vector ()
 {
-  if (acc_on_device ((int) acc_device_host))
-    return 0;
-  else if (acc_on_device ((int) acc_device_nvidia)
-	   || acc_on_device ((int) acc_device_radeon))
-    return __builtin_goacc_parlevel_id (GOMP_DIM_VECTOR);
-  else
-    __builtin_abort ();
+  return __builtin_goacc_parlevel_id (GOMP_DIM_VECTOR);
 }
 
 
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims-aux.c b/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims-aux.c
index b5986f4..cdece32 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims-aux.c
+++ b/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims-aux.c
@@ -5,41 +5,22 @@
 
 /* Used by 'parallel-dims.f90'.  */
 
-#include <limits.h>
-#include <openacc.h>
 #include <gomp-constants.h>
 
-/* TODO: "(int) acc_device_*" casts because of the C++ acc_on_device wrapper
-   not behaving as expected for -O0.  */
 #pragma acc routine seq
-/* static */ unsigned int __attribute__ ((optimize ("O2"))) acc_gang ()
+/* static */ int acc_gang ()
 {
-  if (acc_on_device ((int) acc_device_host))
-    return 0;
-  else if (acc_on_device ((int) acc_device_nvidia))
-    return __builtin_goacc_parlevel_id (GOMP_DIM_GANG);
-  else
-    __builtin_abort ();
+  return __builtin_goacc_parlevel_id (GOMP_DIM_GANG);
 }
 
 #pragma acc routine seq
-/* static */ unsigned int __attribute__ ((optimize ("O2"))) acc_worker ()
+/* static */ int acc_worker ()
 {
-  if (acc_on_device ((int) acc_device_host))
-    return 0;
-  else if (acc_on_device ((int) acc_device_nvidia))
-    return __builtin_goacc_parlevel_id (GOMP_DIM_WORKER);
-  else
-    __builtin_abort ();
+  return __builtin_goacc_parlevel_id (GOMP_DIM_WORKER);
 }
 
 #pragma acc routine seq
-/* static */ unsigned int __attribute__ ((optimize ("O2"))) acc_vector ()
+/* static */ int acc_vector ()
 {
-  if (acc_on_device ((int) acc_device_host))
-    return 0;
-  else if (acc_on_device ((int) acc_device_nvidia))
-    return __builtin_goacc_parlevel_id (GOMP_DIM_VECTOR);
-  else
-    __builtin_abort ();
+  return __builtin_goacc_parlevel_id (GOMP_DIM_VECTOR);
 }
-- 
2.7.4