OpenACC : 모든 GPU 스레드에 대한 개인 배열이 있습니다.

Question

코드를 GPU로 가져오고 있습니다. 이 코드에는 개인 배열을 사용하는 커널이 있습니다. 즉, 배열은 커널 루프 내에서 선언됩니다.

OpenACC에 코드를 포팅하면 버그가 발생합니다. 나에게있어서 어레이가 GPU 벡터 쓰레드들 사이에서 공유되는 것처럼 보이고 이로 인해 여러 경쟁 조건이 발생한다.

외부 호출을 사용하여 다음 예제를 구성 했으므로 이것이 원래 코드처럼 보입니다.

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; 
} 

f.c :

#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]; 
} 

내가 Export PGI=/opt/pgi/17.5.0을 반환 cc -v와 컴파일러 버전을 볼 수 있습니다.

컴파일하려면

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 

코드는 i와 같은 모든 A[i] 요소를 설정해야합니다. OpenACC 지원으로이 코드를 실행하면 완전히 잘못된 결과를 얻습니다. 제 추측은 경쟁 조건입니다. openacc이없는 버전은 올바르게 컴파일되고 실행됩니다. A[i]==i

그래서 실행의 끝에서, 내 질문은 :은 어떻게 작은 배열이 OpenACC 모든 GPU 스레드 에 개인으로 만들 수 있습니까?

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