OpenACC Scheduling

Question

Say that I have a construct like this:

for(int i=0;i<5000;i++){
  const int upper_bound = f(i);
  #pragma acc parallel loop
  for(int j=0;j<upper_bound;j++){
    //Do work...
  }
}

Where f is a monotonically-decreasing function of i.

Since num_gangs, num_workers, and vector_length are not set, OpenACC chooses what it thinks is an appropriate scheduling.

But does it choose such a scheduling afresh each time it encounters the pragma, or only once the first time the pragma is encountered?

Looking at the output of PGI_ACC_TIME suggests that scheduling is only performed once.

jefflarkin jefflarkin · Accepted Answer · 2018-08-09T20:09:23

The PGI compiler will choose how to decompose the work at compile-time, but will generally determine the number of gangs at runtime. Gangs are inherently scalable parallelism, so the decision on how many can be deferred until runtime. The vector length and number of workers affects how the underlying kernel gets generated, so they're generally selected at compile-time to maximize optimization opportunities. With loops like these, where the bounds aren't really known at compile-time, the compiler has to generate some extra code in the kernel to ensure exactly the correct number of iterations are performed.

OpenACC Scheduling

2 Answers