A very useful feature in OpenMP is the possibility to declare allocatable arrays as THREADPRIVATE, in such a way that once initially allocated by the varius threads in a parallel region of the program, they can be privately accessed within any other parallel part of the code activated later by OMP PARALLEL - OMP END PARALLEL. This allows for saving much time by avoiding re-allocation and re-deallocation of private space within different parallel portions of the code, when those private arrays are needed from the beginning to the end of the program. I wonder if a similar possibility exists for allocatable, device arrays, to be initially allocated in a GPU by different threads in an OpenMP parallel region. I use CUDA Fortran (but this is unimportant), and it seems that on exit from the OMP parallel region responsible of allocation of the memory space in the GPU, the arrays are no longer accessible from subsequent new parallel regions of the code. Thankyou in advance for the help :)
Hi Arnaldo,
Just a question, have you declared the variables used to alloc memory in the device (typically XXX_d) as threadprivate?
Best!
Hi Eduardo,
thankyou for your reply :)
Yes, I did, but actually it seems that arrays declared with the attribute 'device', cannot be declared as threadprivate at the same time. If you try to do so, you get complaints from the compiler. Here is an example with an array called 'w_dev':
pgf95 -O4 -fastsse -Minline -Munroll -mp -Mcuda -c module_args_exchange.f90
PGF90-S-1011-Device variable cannot be THREADPRIVATE - w_dev (module_args_exchange.f90) 0 inform, 0 warnings, 1 severes, 0 fatal for nn_args_exchange
Hi again,
Ok, my point is that you should be able to do it for variables declared in the host and used to allocate memory in the device... for example:
float *A, *A_d;
cudaMalloc(*A_d....)
cudaMemcpy(*A_d, *A ....)
kernel(A_d .... )
A_d could be made treadprivate (a least theoretically), and that should be enough for you, shouldn´t it?
Hi Eduardo,
yes, in principle the one you indicate is the solution, but in practice it seems to be forbidden, at least in CUDA Fortran.
Ideed, in CUDA Fortran, the array A_d must be declared as 'device', in order to inform the compiler that it must be allocated in the GPU, and not in the host.
For instance:
Real, Allocatable :: A(:) ! this will be allocated in the host
Real, Allocatable, device :: A_d(:) ! this will be allocated in the GPU
If you declare A as threadprivate with:
!$OMP THREADPRIVATE(A)
everything is OK, but if you also try to define A_d as threadprivate with:
!$OMP THREADPRIVATE(A_d)
The compiler stops and gives an error. It seems that device arrays are not treated the same way as ordinary host arrays.
Hi,
I realized you were working in Fortran after my 2nd answer
No idea how to manage this directly in Fortran :(
I'll follow this question.
Good luck!
- Badges
- Science topic
- Similar topics
- Physical Sciences
- Waves