[openacc] Add __builtin_goacc_parlevel_{id,size}

[thirdparty/gcc.git] / libgomp / testsuite / libgomp.oacc-c-c++-common / routine-g-1.c

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c
index b6ab7134c0bc640620970e8aa56b76637ffdb2fd..a164f576bc3faa81b419693e1dc1508988e182b9 100644 (file)
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c
@@ -1,8 +1,6 @@
-/* This code uses nvptx inline assembly guarded with acc_on_device, which is
-   not optimized away at -O0, and then confuses the target assembler.
-   { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */
-
 #include <stdio.h>
+#include <openacc.h>
+#include <gomp-constants.h>
 
 #define N (32*32*32+17)
 
@@ -12,13 +10,13 @@ void __attribute__ ((noinline)) gang (int ary[N])
 #pragma acc loop gang
     for (unsigned ix = 0; ix < N; ix++)
       {
-       if (__builtin_acc_on_device (5))
+       if (acc_on_device (acc_device_not_host))
          {
-           int g = 0, w = 0, v = 0;
+           int g, w, v;
 
-           __asm__ volatile ("mov.u32 %0,%%ctaid.x;" : "=r" (g));
-           __asm__ volatile ("mov.u32 %0,%%tid.y;" : "=r" (w));
-           __asm__ volatile ("mov.u32 %0,%%tid.x;" : "=r" (v));
+           g = __builtin_goacc_parlevel_id (GOMP_DIM_GANG);
+           w = __builtin_goacc_parlevel_id (GOMP_DIM_WORKER);
+           v = __builtin_goacc_parlevel_id (GOMP_DIM_VECTOR);
            ary[ix] = (g << 16) | (w << 8) | v;
          }
        else
@@ -38,7 +36,7 @@ int main ()
   
 #pragma acc parallel num_gangs(32) copy(ary) copy(ondev)
   {
-    ondev = __builtin_acc_on_device (5);
+    ondev = acc_on_device (acc_device_not_host);
     gang (ary);
   }