我正在为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];
}
我可以看到带有返回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支持运行此代码时，我得到了完全错误的结果。我的猜测是竞争条件。没有openacccompile的版本并且运行正常。在运行结束时A[i]==i
所以，我的问题是：如何使用OpenACC将一个小型数组专用于所有GPU线程？
“priv”的声明从循环中被提升，从而使它在线程之间共享。解决方法是在循环之前声明“priv”，然后使用“private”子句对其进行私有化。您还需要将循环计划为“gang vector”以防止编译器自动并行化两个内部循环。
例如：
% 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