Don't mark OpenACC auto loops as independent inside acc parallel regions

2018-09-20  Cesar Philippidis  <cesar@codesourcery.com>

	gcc/
	* omp-low.cc (lower_oacc_head_mark): Don't mark OpenACC auto
	loops as independent inside acc parallel regions.

	gcc/testsuite/
	* c-c++-common/goacc/loop-auto-1.c: Adjust test case to conform to
	the new behavior of the auto clause in OpenACC 2.5.
	* c-c++-common/goacc/loop-auto-2.c: Likewise.
	* gcc.dg/goacc/loop-processing-1.c: Likewise.
	* c-c++-common/goacc/loop-auto-3.c: New test.
	* gfortran.dg/goacc/loop-auto-1.f90: New test.

	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c: Adjust test case
	to conform to the new behavior of the auto clause in OpenACC 2.5.

diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index ac03493f34c8730bce21c8c94e8e538ce597953c..5fd096c27bc76722ec752c7bc83ea3b144fa1163 100644 (file)
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,8 @@
+2018-09-20  Cesar Philippidis  <cesar@codesourcery.com>
+
+       * omp-low.cc (lower_oacc_head_mark): Don't mark OpenACC auto
+       loops as independent inside acc parallel regions.
+
 2018-12-22  Cesar Philippidis  <cesar@codesourcery.com>
             Julian Brown  <julian@codesourcery.com>
             Tobias Burnus  <tobias@codesourcery.com>

diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 779cc26e2049c1cad43cd7b94babb26d3aca5073..5575ae47ef740a4d3aab27ebc8515c792003aa03 100644 (file)
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -8658,8 +8658,10 @@ lower_oacc_head_mark (location_t loc, tree ddvar, tree clauses,
   else
     gcc_unreachable ();
 
-  /* In a parallel region, loops are implicitly INDEPENDENT.  */
-  if (!tgt || is_oacc_parallel_or_serial (tgt))
+  /* In a parallel region, loops without auto and seq clauses are
+     implicitly INDEPENDENT.  */
+  if ((!tgt || is_oacc_parallel_or_serial (tgt))
+      && !(tag & (OLF_SEQ | OLF_AUTO)))
     tag |= OLF_INDEPENDENT;
 
   /* Loops inside OpenACC 'kernels' decomposed parts' regions are expected to

diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index 9be386620c1c394dda815461905854323a90b9fc..a5477c45a763edc60bd735471f4436250d9ee70b 100644 (file)
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,12 @@
+2018-09-20  Cesar Philippidis  <cesar@codesourcery.com>
+
+       * c-c++-common/goacc/loop-auto-1.c: Adjust test case to conform to
+       the new behavior of the auto clause in OpenACC 2.5.
+       * c-c++-common/goacc/loop-auto-2.c: Likewise.
+       * gcc.dg/goacc/loop-processing-1.c: Likewise.
+       * c-c++-common/goacc/loop-auto-3.c: New test.
+       * gfortran.dg/goacc/loop-auto-1.f90: New test.
+
 2018-08-28  Julian Brown  <julian@codesourcery.com>
            Cesar Philippidis  <cesar@codesourcery.com>
 

diff --git a/gcc/testsuite/c-c++-common/goacc/loop-auto-1.c b/gcc/testsuite/c-c++-common/goacc/loop-auto-1.c
index 124befc400222e1c9ceca1ca669abca418842c13..dcad07f11c887015bbd864a9c8d6e69a68fd4cc8 100644 (file)
--- a/gcc/testsuite/c-c++-common/goacc/loop-auto-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/loop-auto-1.c
@@ -10,7 +10,7 @@ void Foo ()
 #pragma acc loop seq
        for (int jx = 0; jx < 10; jx++) {}
 
-#pragma acc loop auto /* { dg-warning "insufficient partitioning" } */
+#pragma acc loop auto independent /* { dg-warning "insufficient partitioning" } */
        for (int jx = 0; jx < 10; jx++) {}
       }
 
@@ -20,7 +20,7 @@ void Foo ()
 #pragma acc loop auto
        for (int jx = 0; jx < 10; jx++) {}
 
-#pragma acc loop auto /* { dg-warning "insufficient partitioning" } */
+#pragma acc loop auto independent /* { dg-warning "insufficient partitioning" } */
        for (int jx = 0; jx < 10; jx++)
          {
 #pragma acc loop vector
@@ -51,7 +51,7 @@ void Foo ()
 #pragma acc loop vector
        for (int jx = 0; jx < 10; jx++)
          {
-#pragma acc loop auto /* { dg-warning "insufficient partitioning" } */
+#pragma acc loop auto independent /* { dg-warning "insufficient partitioning" } */
            for (int kx = 0; kx < 10; kx++) {}
          }
 
@@ -64,27 +64,27 @@ void Foo ()
 
       }
     
-#pragma acc loop auto
+#pragma acc loop auto independent
     for (int ix = 0; ix < 10; ix++)
       {
-#pragma acc loop auto
+#pragma acc loop auto independent
        for (int jx = 0; jx < 10; jx++)
          {
-#pragma acc loop auto
+#pragma acc loop auto independent
            for (int kx = 0; kx < 10; kx++) {}
          }
       }
 
-#pragma acc loop auto
+#pragma acc loop auto independent
     for (int ix = 0; ix < 10; ix++)
       {
-#pragma acc loop auto
+#pragma acc loop auto independent
        for (int jx = 0; jx < 10; jx++)
          {
-#pragma acc loop auto /* { dg-warning "insufficient partitioning" } */
+#pragma acc loop auto independent /* { dg-warning "insufficient partitioning" } */
            for (int kx = 0; kx < 10; kx++)
              {
-#pragma acc loop auto
+#pragma acc loop auto independent
                for (int lx = 0; lx < 10; lx++) {}
              }
          }
@@ -101,7 +101,7 @@ void Gang (void)
 #pragma acc loop seq
        for (int jx = 0; jx < 10; jx++) {}
 
-#pragma acc loop auto /* { dg-warning "insufficient partitioning" } */
+#pragma acc loop auto independent /* { dg-warning "insufficient partitioning" } */
        for (int jx = 0; jx < 10; jx++) {}
       }
 
@@ -111,7 +111,7 @@ void Gang (void)
 #pragma acc loop auto
        for (int jx = 0; jx < 10; jx++) {}
 
-#pragma acc loop auto /* { dg-warning "insufficient partitioning" } */
+#pragma acc loop auto independent /* { dg-warning "insufficient partitioning" } */
        for (int jx = 0; jx < 10; jx++)
          {
 #pragma acc loop vector
@@ -142,7 +142,7 @@ void Gang (void)
 #pragma acc loop vector
        for (int jx = 0; jx < 10; jx++)
          {
-#pragma acc loop auto /* { dg-warning "insufficient partitioning" } */
+#pragma acc loop auto independent /* { dg-warning "insufficient partitioning" } */
            for (int kx = 0; kx < 10; kx++) {}
          }
 
@@ -176,7 +176,7 @@ void Worker (void)
 #pragma acc loop seq
        for (int jx = 0; jx < 10; jx++) {}
 
-#pragma acc loop auto /* { dg-warning "insufficient partitioning" } */
+#pragma acc loop auto independent /* { dg-warning "insufficient partitioning" } */
        for (int jx = 0; jx < 10; jx++) {}
       }
 
@@ -186,7 +186,7 @@ void Worker (void)
 #pragma acc loop auto
        for (int jx = 0; jx < 10; jx++) {}
 
-#pragma acc loop auto /* { dg-warning "insufficient partitioning" } */
+#pragma acc loop auto independent /* { dg-warning "insufficient partitioning" } */
        for (int jx = 0; jx < 10; jx++)
          {
 #pragma acc loop vector
@@ -194,20 +194,20 @@ void Worker (void)
          }
       }
 
-#pragma acc loop auto
+#pragma acc loop
     for (int ix = 0; ix < 10; ix++)
       {
-#pragma acc loop auto
+#pragma acc loop
        for (int jx = 0; jx < 10; jx++) {}
       }
 
-#pragma acc loop auto
+#pragma acc loop
     for (int ix = 0; ix < 10; ix++)
       {
-#pragma acc loop auto /* { dg-warning "insufficient partitioning" } */
+#pragma acc loop /* { dg-warning "insufficient partitioning" } */
        for (int jx = 0; jx < 10; jx++)
          {
-#pragma acc loop auto
+#pragma acc loop
            for (int kx = 0; kx < 10; kx++) {}
          }
       }
@@ -222,17 +222,17 @@ void Vector (void)
 #pragma acc loop seq
        for (int jx = 0; jx < 10; jx++) {}
 
-#pragma acc loop auto /* { dg-warning "insufficient partitioning" } */
+#pragma acc loop auto independent /* { dg-warning "insufficient partitioning" } */
        for (int jx = 0; jx < 10; jx++) {}
       }
 
-#pragma acc loop auto
+#pragma acc loop auto independent
     for (int ix = 0; ix < 10; ix++) {}
 
-#pragma acc loop auto /* { dg-warning "insufficient partitioning" } */
+#pragma acc loop auto independent /* { dg-warning "insufficient partitioning" } */
     for (int ix = 0; ix < 10; ix++)
       {
-#pragma acc loop auto
+#pragma acc loop auto independent
        for (int jx = 0; jx < 10; jx++) {}
       }
 }
@@ -240,6 +240,6 @@ void Vector (void)
 #pragma acc routine seq
 void Seq (void)
 {
-#pragma acc loop auto /* { dg-warning "insufficient partitioning" } */
+#pragma acc loop auto independent /* { dg-warning "insufficient partitioning" } */
     for (int ix = 0; ix < 10; ix++) {}
 }

diff --git a/gcc/testsuite/c-c++-common/goacc/loop-auto-2.c b/gcc/testsuite/c-c++-common/goacc/loop-auto-2.c
index af3f0bddf2cbb8e4f189bd23e127b8008d84ba09..5aa36e93ab8f5848606193edd2e7ee700a85ec6d 100644 (file)
--- a/gcc/testsuite/c-c++-common/goacc/loop-auto-2.c
+++ b/gcc/testsuite/c-c++-common/goacc/loop-auto-2.c
@@ -72,12 +72,12 @@ void Bad ()
 #pragma acc loop tile(*) gang vector
     for (int ix = 0; ix < 10; ix++)
       {
-       #pragma acc loop auto /* { dg-warning "insufficient partitioning" } */
+       #pragma acc loop auto independent /* { dg-warning "insufficient partitioning" } */
        for (int jx = 0; jx < 10; jx++)
          ;
       }
 
-#pragma acc loop tile(*) auto /* { dg-warning "insufficient partitioning" } */
+#pragma acc loop tile(*) auto independent /* { dg-warning "insufficient partitioning" } */
     for (int ix = 0; ix < 10; ix++)
       {
        #pragma acc loop worker

diff --git a/gcc/testsuite/c-c++-common/goacc/loop-auto-3.c b/gcc/testsuite/c-c++-common/goacc/loop-auto-3.c
new file mode 100644 (file)
index 0000000..8f79ead
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/loop-auto-3.c
@@ -0,0 +1,78 @@
+/* Ensure that the auto clause falls back to seq parallelism when the
+   OpenACC loop is not explicitly independent.  */
+
+/* { dg-additional-options "-fopt-info-optimized-omp" } */
+
+void
+test ()
+{
+  int i, j, k, l, n = 100;
+  
+#pragma acc parallel loop auto /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */
+  for (i = 0; i < n; i++)
+#pragma acc loop auto independent /* { dg-message "optimized: assigned OpenACC gang loop parallelism" } */
+    for (j = 0; j < n; j++)
+#pragma acc loop worker vector /* { dg-message "optimized: assigned OpenACC worker vector loop parallelism" } */
+      for (k = 0; k < n; k++)
+       ;
+
+#pragma acc parallel loop auto independent /* { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } */
+  for (i = 0; i < n; i++)
+#pragma acc loop auto /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */
+    for (j = 0; j < n; j++)
+#pragma acc loop auto /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */
+      for (k = 0; k < n; k++)
+#pragma acc loop auto independent /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */
+       for (l = 0; l < n; l++)
+         ;
+
+#pragma acc parallel loop gang /* { dg-message "optimized: assigned OpenACC gang loop parallelism" } */
+  for (i = 0; i < n; i++)
+#pragma acc loop worker /* { dg-message "optimized: assigned OpenACC worker loop parallelism" } */
+    for (j = 0; j < n; j++)
+#pragma acc loop vector /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */
+      for (k = 0; k < n; k++)
+       {
+#pragma acc loop auto independent /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */
+         /* { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } */
+         for (l = 0; l < n; l++)
+           ;
+#pragma acc loop auto /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */
+         for (l = 0; l < n; l++)
+           ;
+       }
+
+#pragma acc parallel loop /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */
+  /* { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } */
+  for (i = 0; i < n; i++)
+    {
+#pragma acc loop gang worker /* { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } */
+      for (j = 0; j < n; j++)
+#pragma acc loop auto /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */
+       for (k = 0; k < n; k++)
+         {
+#pragma acc loop vector /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */
+           for (l = 0; l < n; l++)
+             ;
+#pragma acc loop auto independent /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */
+           for (l = 0; l < n; l++)
+             ;
+         }
+#pragma acc loop worker /* { dg-message "optimized: assigned OpenACC worker loop parallelism" } */
+      for (j = 0; j < n; j++)
+#pragma acc loop vector /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */
+       for (k = 0; k < n; k++)
+         ;
+    }
+
+#pragma acc parallel loop /* { dg-message "optimized: assigned OpenACC gang loop parallelism" } */
+  for (i = 0; i < n; i++)
+#pragma acc loop /* { dg-message "optimized: assigned OpenACC worker loop parallelism" } */
+    for (j = 0; j < n; j++)
+#pragma acc loop /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */
+      /* { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } */
+      for (k = 0; k < n; k++)
+#pragma acc loop /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */
+         for (l = 0; l < n; l++)
+           ;
+}

diff --git a/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c b/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c
index 698f008580e32d17e50ee20ea0f38b46cffaa721..4230ddc33f70de6f5a1938d8c7c8fa40e972c0ca 100644 (file)
--- a/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c
+++ b/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c
@@ -9,7 +9,7 @@ void vector_1 (int *ary, int size)
   {
 #pragma acc loop gang
     for (int jx = 0; jx < 1; jx++)
-#pragma acc loop auto
+#pragma acc loop auto independent
       for (int ix = 0; ix < size; ix++)
        ary[ix] = place ();
   }

diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-auto-1.f90 b/gcc/testsuite/gfortran.dg/goacc/loop-auto-1.f90
new file mode 100644 (file)
index 0000000..8d600f4
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/loop-auto-1.f90
@@ -0,0 +1,88 @@
+! Ensure that the auto clause falls back to seq parallelism when the
+! OpenACC loop is not explicitly independent.
+
+! { dg-additional-options "-fopt-info-optimized-omp" }
+
+program test
+  implicit none
+  integer, parameter :: n = 100
+  integer i, j, k, l
+  
+  !$acc parallel loop auto ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+  do i = 1, n
+     !$acc loop auto independent ! { dg-message "optimized: assigned OpenACC gang loop parallelism" }
+     do j = 1, n
+        !$acc loop worker vector ! { dg-message "optimized: assigned OpenACC worker vector loop parallelism" }
+        do k = 1, n
+        end do
+     end do
+  end do
+ 
+  !$acc parallel loop auto independent ! { dg-message "optimized: assigned OpenACC gang worker loop parallelism" }
+  do i = 1, n
+     !$acc loop auto ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+     do j = 1, n
+        !$acc loop auto ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+        do k = 1, n
+           !$acc loop auto independent ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+           do l = 1, n
+           end do
+        end do
+     end do
+  end do
+
+  !$acc parallel loop gang ! { dg-message "optimized: assigned OpenACC gang loop parallelism" }
+  do i = 1, n
+     !$acc loop worker ! { dg-message "optimized: assigned OpenACC worker loop parallelism" }
+     do j = 1, n
+        !$acc loop vector ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+        do k = 1, n
+           !$acc loop auto independent ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+           ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
+           do l = 1, n
+           end do
+           !$acc loop auto ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+           do l = 1, n
+           end do
+        end do
+     end do
+  end do
+  
+
+  !$acc parallel loop ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+  ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
+  do i = 1, n
+     !$acc loop gang worker ! { dg-message "optimized: assigned OpenACC gang worker loop parallelism" }
+     do j = 1, n
+        !$acc loop auto ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+       do k = 1, n
+          !$acc loop vector ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+          do l = 1, n
+          end do
+       end do
+       !$acc loop auto independent ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+       do l = 1, n
+       end do
+    end do
+    !$acc loop worker ! { dg-message "optimized: assigned OpenACC worker loop parallelism" }
+    do j = 1, n
+       !$acc loop vector ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+       do k = 1, n
+       end do
+    end do
+  end do
+
+  !$acc parallel loop ! { dg-message "optimized: assigned OpenACC gang loop parallelism" }
+  do i = 1, n
+     !$acc loop ! { dg-message "optimized: assigned OpenACC worker loop parallelism" }
+     do j = 1, n
+        !$acc loop ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+        ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
+        do k = 1, n
+           !$acc loop ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+           do l = 1, n
+           end do
+        end do
+     end do
+  end do
+end program test

diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index 9db651f6983f5b695dc8ba8821f5f076fbcc29ea..d06a78f321b3846e33cfe279043bad2472418e09 100644 (file)
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,8 @@
+2018-09-20  Cesar Philippidis  <cesar@codesourcery.com>
+
+       * testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c: Adjust test case
+       to conform to the new behavior of the auto clause in OpenACC 2.5.
+
 2018-12-22  Cesar Philippidis  <cesar@codesourcery.com>
             Julian Brown  <julian@codesourcery.com>
 

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c
index c13cab780cb502a144467fd9239c45ab36bfa633..418275592972aa8cce8040e13ad6ba8b853ac88f 100644 (file)
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c
@@ -107,7 +107,7 @@ int vector_1 (int *ary, int size)
   {
 #pragma acc loop gang
     for (int jx = 0; jx < 1; jx++)
-#pragma acc loop auto
+#pragma acc loop auto independent
       for (int ix = 0; ix < size; ix++)
        ary[ix] = place ();
   }
@@ -123,7 +123,7 @@ int vector_2 (int *ary, int size)
   {
 #pragma acc loop worker
     for (int jx = 0; jx < size  / 64; jx++)
-#pragma acc loop auto
+#pragma acc loop auto independent
       for (int ix = 0; ix < 64; ix++)
        ary[ix + jx * 64] = place ();
   }
@@ -139,7 +139,7 @@ int worker_1 (int *ary, int size)
   {
 #pragma acc loop gang
     for (int kx = 0; kx < 1; kx++)
-#pragma acc loop auto
+#pragma acc loop auto independent
       for (int jx = 0; jx <  size  / 64; jx++)
 #pragma acc loop vector
        for (int ix = 0; ix < 64; ix++)
@@ -156,7 +156,7 @@ int gang_1 (int *ary, int size)
 #pragma acc parallel num_gangs (32) num_workers (32) vector_length(32) copy(ary[0:size]) firstprivate (size)
   /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-1 } */
   {
-#pragma acc loop auto
+#pragma acc loop auto independent
     for (int jx = 0; jx <  size  / 64; jx++)
 #pragma acc loop worker
       for (int ix = 0; ix < 64; ix++)
@@ -172,11 +172,11 @@ int gang_2 (int *ary, int size)
   
 #pragma acc parallel num_gangs (32) num_workers (32) vector_length(32) copy(ary[0:size]) firstprivate (size)
   {
-#pragma acc loop auto
+#pragma acc loop auto independent
     for (int kx = 0; kx < size / (32 * 32); kx++)
-#pragma acc loop auto
+#pragma acc loop auto independent
       for (int jx = 0; jx <  32; jx++)
-#pragma acc loop auto
+#pragma acc loop auto independent
        for (int ix = 0; ix < 32; ix++)
          ary[ix + jx * 32 + kx * 32 * 32] = place ();
   }
@@ -190,9 +190,9 @@ int gang_3 (int *ary, int size)
   
 #pragma acc parallel num_workers (32) vector_length(32) copy(ary[0:size]) firstprivate (size)
   {
-#pragma acc loop auto
+#pragma acc loop auto independent
     for (int jx = 0; jx <  size  / 64; jx++)
-#pragma acc loop auto
+#pragma acc loop auto independent
       for (int ix = 0; ix < 64; ix++)
        ary[ix + jx * 64] = place ();
   }
@@ -206,7 +206,7 @@ int gang_4 (int *ary, int size)
   
 #pragma acc parallel vector_length(32) copy(ary[0:size]) firstprivate (size)
   {
-#pragma acc loop auto
+#pragma acc loop auto independent
     for (int jx = 0; jx <  size; jx++)
       ary[jx] = place ();
   }