Wie deklariere ich ein globales dynamisches Array mit C/OpenACC mit PGI Compiler

Question

Ich versuche einen einfachen Testfall auszuführen, bei dem ein dynamisch zugewiesenes Array A extern definiert ist und mit OpenACC auf die GPU hochgeladen wird. Alles mit dem PGI-Compiler.

Meine header.h Datei:

 extern int *A; 
    #pragma acc declare create(A) 

Dann meine header.c Umsetzung:

int *A; 
    #pragma acc declare copyin(A) 

Dann in main.c Ich habe

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

ich den Code mit den folgenden Befehlen kompilieren:

$ 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 

Meine PGI Compiler Version ist:

$ cc -v 
Export PGI=/opt/pgi/17.5.0 

um den Code auszuführen:

$ 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 

ich glaube, das Problem ist, dass die PGI Compiler variable=A bytes=8 sendet, also meine Anfrage zu ignorieren A[0:N] Senden .

Also, wie ein globales dynamisches Array mit C/OpenACC mit PGI-Compiler zu deklarieren?

Answer 1

Wenn Sie "declare" mit einem Zeiger verwenden, erstellen Sie einen globalen Gerätezeiger, aber nicht das Array, auf das der Zeiger zeigt. Daher, wenn Sie versuchen, das Array zu aktualisieren, existiert es nicht und warum die Laufzeitfehler.

Um zu beheben, müssen Sie auch das Array zu einem Datenbereich wie die "Daten eingeben" -Direktive hinzufügen, wie ich unten zeige. Wenn Sie das Array in einen Datenbereich setzen, wird die Laufzeit neben dem Platz für das Array zurückgehen und es an "A" "anhängen", d. H. Die Gerätekopie von "A" mit dem korrekten Gerätezeigerwert ausfüllen.

Sie möchten auch dem Compiler mitteilen, dass "A" bereits auf dem Gerät vorhanden ist, indem Sie ein "Present (A)" auf die Rechenbereiche setzen.

Beachten Sie, dass das zweite "Copyin deklarieren" nicht benötigt wird. Mit "create" werden die Gerätedaten nicht initialisiert, während "copyin" die Variable mit dem Host-Wert initialisiert. Aber da der Host-Wert ein Host-Zeiger ist, ist es immer noch Müll auf dem Gerät. Also nicht unbedingt falsch, einfach nicht nötig.

% cat header.h 

#include <stdio.h> 
#include <stdlib.h> 

extern int *A; 
#pragma acc declare create(A) 

% cat header.c 
#include <header.h> 
int *A; 

% cat test.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)); 
     #pragma acc enter data create(A[0:N]) 

     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 present(A) reduction(+:sum) 
      for(int i=0;i<N;i++){ 
       sum+=A[i]; 
      } 
     } 
     printf("sum = %d\n",sum); 
     #pragma acc exit data delete(A) 
     free(A); 
     exit(0); 
    } 
% pgcc -I./ test.c header.c -ta=tesla:cc60 -Minfo=accel 
test.c: 
main: 
    13, Generating enter data create(A[:N]) 
    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 */ 
    27, Generating exit data delete(A[:1]) 
header.c: 
% setenv PGI_ACC_TIME 1 
% a.out 1024 
main() start 
N =1024 
almost data region 
inside data region 
sum = 523776 

Accelerator Kernel Timing data 
test.c 
    main NVIDIA devicenum=0 
    time(us): 124 
    13: upload reached 1 time 
     13: data copyin transfers: 1 
      device time(us): total=33 max=33 min=33 avg=33 
    13: data region reached 1 time 
     13: data copyin transfers: 1 
      device time(us): total=9 max=9 min=9 avg=9 
    17: data region reached 2 times 
     17: data copyin transfers: 1 
      device time(us): total=33 max=33 min=33 avg=33 
     26: data copyout transfers: 1 
      device time(us): total=22 max=22 min=22 avg=22 
    21: update directive reached 1 time 
     21: data copyin transfers: 1 
      device time(us): total=10 max=10 min=10 avg=10 
    21: compute region reached 1 time 
     21: kernel launched 1 time 
      grid: [8] block: [128] 
      device time(us): total=4 max=4 min=4 avg=4 
      elapsed time(us): total=589 max=589 min=589 avg=589 
     21: reduction kernel launched 1 time 
      grid: [1] block: [256] 
      device time(us): total=4 max=4 min=4 avg=4 
      elapsed time(us): total=27 max=27 min=27 avg=27 
    27: data region reached 1 time 
     27: data copyin transfers: 1 
      device time(us): total=9 max=9 min=9 avg=9