[llvm-dev] Automatic Insertion of OpenACC/OpenMP directives

[Fernando Magno Quintao Pereira via llvm-dev]

Dear Jonathan and Mehdi,

    you are right. To extend the ability of the OpenACC compliant
compiler, we use a technique called restrictification to disambiguate
pointers at runtime. So, if you use this program below as input in
http://cuda.dcc.ufmg.br/dawn:

float saxpy(float a, float *x, float *y, int n) {
 int j = 0;
 for (int i = 0; i < n; ++i) {
   y[j] = a*x[i] + y[j];
   ++j;
 }
}

    then the final program that you obtain is this one:

float saxpy(float a, float *x, float *y, int n) {
  int j = 0;
  long long int AI1[6];
  AI1[0] = n + -1;
  AI1[1] = 4 * AI1[0];
  AI1[2] = AI1[1] + 4;
  AI1[3] = AI1[2] / 4;
  AI1[4] = (AI1[3] > 0);
  AI1[5] = (AI1[4] ? AI1[3] : 0);
  char RST_AI1 = 0;
  RST_AI1 |= !((x + 0 > y + AI1[5]) || (y + 0 > x + AI1[5])); //
Restrictification test!
  #pragma acc data pcopyin(x[0:AI1[5]]) pcopy(y[0:AI1[5]]) if(!RST_AI1)
  #pragma acc kernels if(!RST_AI1) // See the test being used here!
  for (int i = 0; i < n; ++i) {
    y[j] = a * x[i] + y[j];
    ++j;
  }
}

    The input kernel, saxpy is, in the absence of aliasing, a map.
Thus, it is parallelizable (assuming no aliasing between arrays x and
y). So, how does dawn-cc handle aliasing? Our tool does it via a
restrictification check. In the code above, this restrictification
check is encoded in the variable RST_AI1. This variable is the result
of a test that ensures that the limits of arrays x and y do not
overlap. LLVM itself uses some form of restrictification to carry out
vectorization. Our checks are produced after the paper:

Pericles Rafael Alves, Fabian Gruber, Johannes Doerfert, Alexandros
Labrineas, Tobias Grosser, Fabrice Rastello and Fernando Magno Quintão
Pereira. Runtime Pointer Disambiguation. OOPSLA'2015.

    In this example, aliasing does not occur as long as &x + n <= &y
(or vice versa). Dawn-cc can create restrictification checks as long
as all the involved variables encode affine expressions of induction
variables.

Regards,

Fernando