Private Array that goes through recursion for each gang in openacc

Question

How to properly make an array private in openacc? for my use, i have declared an array that goes through a recursion function and updates itself. In C++ I did it this way,

int recursion(int array[],int i)
{
   ....
   return array[i] = recursion(array,i);
}
int main()
{
  ....
   for(int run=0;run<N;++run)
   { 
    .... 
      for(int i=0;i<some N;++i)
      int f = recursion(array,i);
   }
 .....
}

Now the main problem begins when i try to do this in parallel using openacc directives. I want to copy this array into the parallel for loop region in such a way that each gang is going to have a copy of this array and will be able to change their own version of array[] and not the others using the recursion function. The way i tried doing this is

#pragma acc routine seq
int recursion(int array[],int i)
{
 ....
 return array[i] = recursion(array,i);
}
int main()
{
 ....
 #pragma acc data copyin(array[0:N])
 #pragma acc parallel loop gang private(array[0:N])

 for(int run=0;run<N;++run)
 {  
  ....
    for(int i=0;i<N;++i)
    { 
     ....
     int f = recursion(array,i);
    }
  }
 }

But it seems that the array is not getting passed to the recursion function as i have checked that it does not change. What is the perfect way to do this?

p.s. I also have tried #pragma acc data pcreate(array[0:N]) with #pragma acc parallel loop independent private(array[0:N]) but the result is the same

you can find the whole code here. It works perfectly without the directives. only thing you need to change is comment the line with curand and un-comment the line with rand at 251 268 line. Please Help!

Mat Colgrove Mat Colgrove · Accepted Answer · 2017-10-31T16:23:50

A "private" array is uninitialized but your code is expecting that they have an initial value. You'll either want to initialize the arrays within the loop our use the "firstprivate" clause which will initialize each private array with the initial value from the host.

Also, you have "random" and "ptr" in both a private clause and a data clause. You should remove them from the data directive since this will make them global.

Recursion on the device is problematic given there is a very small stack on the GPU (8MB). You can increase this by setting the PGI environment variable "PGI_CUDA_STACKSIZE" to a larger value, but if you recurse too deep, the program will crash.

I tried running your program but it kept erroring, presumably due to a stack overflow. I wasn't sure what input values to use so it could also be pilot error on my part.

If using "firstprivate" doesn't help, please let me know what input values to use, and I'll look into it further.

Private Array that goes through recursion for each gang in openacc

1 Answers