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.

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 (file)
--- 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 (file)
--- 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);
 }