OpenACC: наличие частного массива для каждого потока графического процессора

Question

Я несу код в графический процессор. Этот код имеет ядро, которое использует приватный массив. Это означает, что массив объявлен внутри цикла ядра.

Когда я портирую код на OpenACC, я получаю ошибочные результаты. Мне кажется, что массив распределяется между векторными потоками графического процессора, и это вызывает несколько условий гонки.

Я организовал следующий пример с внешними вызовами, потому что так выглядит мой исходный код.

header.h:

#define N 100000
#define K 16
#pragma acc routine
void assign_i_to_privj(int * priv, int j, int i);
#pragma acc routinetnumpy
void add_privi_to_sum(int * priv, int i, int *sum);

main.c:

#include "header.h"
int main(void){
int A[N];
#pragma acc data copy(A)
{
#pragma acc parallel loop
     for(int i=0; i<N;i++){
       int priv[K];
       int sum=0;
       int j=0;
       while(1){
       if(j>=K) break;
           assign_i_to_privj(priv, j, i);
           j++;
       }
       j=0;
       while(1){
           if(j>=K) break;
           add_privi_to_sum(priv, j, &sum);
           j++;
       }
       sum/=K; // now sum == i;
       A[i]=sum;
     }
   }
   //now A[i] == i
   for(int i=0; i<123; i++) printf("A[%d]=%d ",i, A[i]);
   printf("\n");
   return 0;
}

Ь:

#include "header.h"
void assign_i_to_privj(int *priv, int j, int i){
       priv[j]=i;
}
void add_privi_to_sum(int *priv, int j, int *sum){
      (*sum)+=priv[j];
}

Я вижу версию компилятора с cc -v это возвращает Export PGI=/opt/pgi/17.5.0,

Скомпилировать:

cc -g -lnvToolsExt -O2  -acc  -ta=tesla:cc60  -c11 -mp -Minfo -Mlarge_arrays -c main.c &&
cc -g -lnvToolsExt -O2 -acc  -ta=tesla:cc60  -c11 -mp -Minfo -Mlarge_arrays -c f.c &&
cc -g -lnvToolsExt -O2 -acc  -ta=tesla:cc60 -c11 -mp -Minfo -Mlarge_arrays f.o main.o -o acc.exe &&
srun  -n 1 acc.exe

Код должен установить все A[i] элементы равные i, Когда я запускаю этот код с поддержкой OpenACC, я получаю совершенно неверные результаты. Я думаю, это состояние гонки. Версия без openaccкомпилируется и работает правильно. В конце пробега A[i]==i

Итак, мой вопрос: как я могу сделать небольшой массив частным для всех потоков графического процессора с OpenACC?

Answer 1

Объявление "priv" выводится из цикла, делая его разделяемым между потоками. Обходной путь - объявить "priv" перед циклом, а затем использовать предложение "private" для его приватизации. Вы также захотите запланировать цикл как "вектор банды", чтобы компилятор не мог автоматически распараллеливать два внутренних цикла.

Например:

% cat main.c
#include "header.h"
int main(void){
int A[N];
int priv[K];
#pragma acc data copy(A)
{
#pragma acc parallel loop gang vector private(priv)
     for(int i=0; i<N;i++){
       int sum=0;
       int j=0;
       while(1){
       if(j>=K) break;
           assign_i_to_privj(priv, j, i);
           j++;
       }
       j=0;
       while(1){
           if(j>=K) break;
           add_privi_to_sum(priv, j, &sum);
           j++;
       }
       sum/=K; // now sum == i;
       A[i]=sum;
     }
   }
   //now A[i] == i
   for(int i=0; i<123; i++) printf("A[%d]=%d ",i, A[i]);
   printf("\n");
   return 0;
}
% pgcc f.c main.c -Minfo=acc -ta=tesla:cc60 -fast -V17.10
f.c:
assign_i_to_privj:
      2, Generating acc routine seq
         Generating Tesla code
add_privi_to_sum:
      5, Generating acc routine seq
         Generating Tesla code
main.c:
main:
      5, Generating copy(A[:])
      7, Accelerator kernel generated
         Generating Tesla code
          8, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
         11, #pragma acc loop seq
         17, #pragma acc loop seq
      7, Local memory used for priv
     11, Loop is parallelizable
     17, Loop is parallelizable
% a.out
A[0]=0 A[1]=1 A[2]=2 A[3]=3 A[4]=4 A[5]=5 A[6]=6 A[7]=7 A[8]=8 A[9]=9 A[10]=10 A[11]=11 A[12]=12 A[13]=13 A[14]=14 A[15]=15 A[16]=16 A[17]=17 A[18]=18 A[19]=19 A[20]=20 A[21]=21 A[22]=22 A[23]=23 A[24]=24 A[25]=25 A[26]=26 A[27]=27 A[28]=28 A[29]=29 A[30]=30 A[31]=31 A[32]=32 A[33]=33 A[34]=34 A[35]=35 A[36]=36 A[37]=37 A[38]=38 A[39]=39 A[40]=40 A[41]=41 A[42]=42 A[43]=43 A[44]=44 A[45]=45 A[46]=46 A[47]=47 A[48]=48 A[49]=49 A[50]=50 A[51]=51 A[52]=52 A[53]=53 A[54]=54 A[55]=55 A[56]=56 A[57]=57 A[58]=58 A[59]=59 A[60]=60 A[61]=61 A[62]=62 A[63]=63 A[64]=64 A[65]=65 A[66]=66 A[67]=67 A[68]=68 A[69]=69 A[70]=70 A[71]=71 A[72]=72 A[73]=73 A[74]=74 A[75]=75 A[76]=76 A[77]=77 A[78]=78 A[79]=79 A[80]=80 A[81]=81 A[82]=82 A[83]=83 A[84]=84 A[85]=85 A[86]=86 A[87]=87 A[88]=88 A[89]=89 A[90]=90 A[91]=91 A[92]=92 A[93]=93 A[94]=94 A[95]=95 A[96]=96 A[97]=97 A[98]=98 A[99]=99 A[100]=100 A[101]=101 A[102]=102 A[103]=103 A[104]=104 A[105]=105 A[106]=106 A[107]=107 A[108]=108 A[109]=109 A[110]=110 A[111]=111 A[112]=112 A[113]=113 A[114]=114 A[115]=115 A[116]=116 A[117]=117 A[118]=118 A[119]=119 A[120]=120 A[121]=121 A[122]=122