Parallel programming dependency openacc

Question

I am trying to parallelize this loops, but get some error in PGI compiler, I don't understand what's wrong

#pragma acc kernels
{
#pragma acc loop independent
    for (i = 0;i < k; i++)
    {
        for(;dt*j <= Ms[i+1].t;j++)
        {
            w = (j*dt - Ms[i].t)/(Ms[i+1].t-Ms[i].t);
            X[j] = Ms[i].x*(1-w)+Ms[i+1].x*w;
            Y[j] = Ms[i].y*(1-w)+Ms[i+1].y*w;
        }
    }
}

Error 85, Generating Multicore code 87, #pragma acc loop gang 89, Accelerator restriction: size of the GPU copy of Y,X is unknown Complex loop carried dependence of Ms->t,Ms->x,X->,Ms->y,Y-> prevents parallelization Loop carried reuse of Y->,X-> prevents parallelization So what i can do to solve this dependence problem?

Mat Colgrove Mat Colgrove · Accepted Answer · 2018-04-06T15:13:29

I see a few issues here. Also given the output, I'm assuming that you're compiling with "-ta=multicore,tesla" (i.e. targeting both a multicore CPU and a GPU)

First, since "j" is not initialized in the "i" loop, the starting value of "j" will depended on the ending value of "j" from the previous iteration of "i". Hence, the loops are not parallelizable. By using "loop independent", you have forced parallelization on the outer loop, but you will get differing answers from running the code sequentially. You will need to rethink your algorithm.

I would suggest making X and Y two dimensional. With the first dimension of size "k". The second dimension can be a jagged array (i.e. each having a differing size) with the size corresponding to the "Ms[i+1].t" value.

I wrote an example of using jagged arrays as part of my Chapter (#5) of the Parallel Programming with OpenACC book. See: https://github.com/rmfarber/ParallelProgrammingWithOpenACC/blob/master/Chapter05/jagged_array.c

Alternatively, you might be able to set "j=Ms[i].t" assuming "Ms[0].t" is set.

for(j=Ms[i].t;dt*j <= Ms[i+1].t;j++)

"Accelerator restriction: size of the GPU copy of Y,X is unknown"

This is telling you that the compiler can not implicitly copy the X and Y arrays on the device. In C/C++, unbounded pointers don't have sizes so the compiler can't tell how big these arrays are. Often it can derive this information from the loop trip counts, but since the loop trip count is unknown (see above), it can't in this case. To fix, you need to include a data directive on the "kernels" directive or add a data region to your code. For example:

#pragma acc kernels copyout(X[0:size], Y[0:size])

or

 #pragma acc data copyout(X[0:size], Y[0:size])
 {
    ...
    #pragma acc kernels 
    ...
 }

Another thing to keep in mind is pointer aliasing. In C/C++, pointers of the same type are allowed to point at the same object. Hence, without additional information such as the "restrict" attribute, the "independent" clause, or the PGI compiler flag "-Msafeptr", the compiler must assume your pointers do point to the same object making the loop not parallelizable.

Parallel programming dependency openacc

2 Answers