A動的に割り当てられた配列が extern として定義され、OpenACC を使用して GPU にアップロードされる簡単なテスト ケースを実行しようとしています。すべて PGI コンパイラを使用します。
私の header.hファイル:
extern int *A;
#pragma acc declare create(A)
次に、私のheader.c実装：
int *A;
#pragma acc declare copyin(A)
それから、main.c私は持っています
#include "header.h"
int main(int argc, char* argv[]){
printf("main() start\n");
int sum=0;
int N=0;
if(argc==1){
printf("usage: ./main.exe N");
}else{
N=atoi(argv[1]);
}
printf("N =%d\n", N);
A=(int*)malloc(N*sizeof(int));
for(int i=0;i<N;i++){A[i]=i;}
printf("almost data region\n");
#pragma acc data copy(sum)
{
printf("inside data region\n");
#pragma acc update device(A[0:N])
#pragma acc parallel loop reduction(+:sum)
for(int i=0;i<N;i++){
sum+=A[i];
}
}
printf("sum = %d\n",sum);
}
次のコマンドでコードをコンパイルします。
$ cc -g -lnvToolsExt -O2 -acc -ta=tesla:cc60 -c11 -mp -Minfo -Mlarge_arrays -c -o header.o header.c
$ cc -g -lnvToolsExt -O2 -acc -ta=tesla:cc60 -c11 -mp -Minfo -Mlarge_arrays -c -o main.o main.c
PGC-W-0155-Pointer value created from a nonlong integral type (main.c: 12)
main:
13, Generated 2 alternate versions of the loop
Generated vector simd code for the loop
17, Generating copy(sum)
21, Generating update device(A[:N])
Accelerator kernel generated
Generating Tesla code
21, Generating reduction(+:sum)
22, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
PGC/x86-64 Linux 17.5-0: compilation completed with warnings
$ cc -g -lnvToolsExt -O2 -acc -ta=tesla:cc60 -c11 -mp -Minfo -Mlarge_arrays header.o main.o -o main.exe
私のPGIコンパイラのバージョンは次のとおりです。
$ cc -v
Export PGI=/opt/pgi/17.5.0
コードを実行するには:
$ ACC_NOTIFY=3 srun cuda-memcheck --show-backtrace yes main.exe 10000
upload CUDA data file=/scratch/snx3000/ragagnin/2017/prova/main.c function=main line=17 device=0 threadid=1 variable=A bytes=8
upload CUDA data file=/scratch/snx3000/ragagnin/2017/prova/main.c function=main line=17 device=0 threadid=1 variable=sum bytes=4
Present table dump for device[1]: NVIDIA Tesla GPU 0, compute capability 6.0, threadid=1
host:0x606780 device:0x10216200000 size:8 presentcount:0+1 line:-1 name:A
host:0x7fffffff67ac device:0x1021a400000 size:4 presentcount:1+0 line:17 name:sum
allocated block device:0x1021a400000 size:512 thread:1
FATAL ERROR: data in update device clause was not found on device 1: name=A
file:/scratch/snx3000/ragagnin/2017/prova/main.c main line:21
main() start
N =10000
almost data region
inside data region
========= CUDA-MEMCHECK
========= Program hit CUDA_ERROR_INVALID_DEVICE (error 101) due to "invalid device ordinal" on CUDA API call to cuDevicePrimaryCtxRetain.
========= Saved host backtrace up to driver entry point at error
========= Host Frame:/opt/cray/nvidia/default/lib64/libcuda.so (cuDevicePrimaryCtxRetain + 0x15c) [0x1e497c]
========= Host Frame:/opt/pgi/17.5.0/linux86-64/17.5/lib/libaccnmp.so (__pgi_uacc_cuda_initdev + 0x962) [0x140e1]
========= Host Frame:/opt/pgi/17.5.0/linux86-64/17.5/lib/libaccgmp.so (__pgi_uacc_enumerate + 0x173) [0x12e31]
========= Host Frame:/opt/pgi/17.5.0/linux86-64/17.5/lib/libaccgmp.so (__pgi_uacc_initialize + 0x9b) [0x1340d]
========= Host Frame:/opt/pgi/17.5.0/linux86-64/17.5/lib/libaccgmp.so (__pgi_uacc_dataenterstart + 0x50) [0x9de1]
========= Host Frame:main.exe [0x16a5]
========= Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xf5) [0x206e5]
========= Host Frame:main.exe [0x11c9]
=========
========= ERROR SUMMARY: 1 error
srun: error: nid03948: task 0: Exited with exit code 1
srun: Terminating job step 4066800.15
問題は、PGI コンパイラが を送信variable=A bytes=8しているため、私の送信要求を無視していることだと思いますA[0:N]。
では、PGI コンパイラを使用して C/OpenACC でグローバル動的配列を宣言するにはどうすればよいでしょうか?