OpenMP target "use_device_ptr" versus "use_device_addr"

[caplanr]

Hi,

I am testing our code POT3D (github.com/predsci/pot3d) on the Stampede3 system at TACC on the MAX 1550 GPUs.

The code uses Fortran "do concurrent" for offload, as well as OpenMP target for data movement and device selection.

The code is designed to run on multiple GPUs using GPU-aware MPI calls.

In order to do this, I am setting I_MPI_OFFLOAD=1 and using OpenMP target directives around the MPI calls.

Currently, the code is using "use_device_ptr()" to tell the compiler to use the GPU arrays in the MPI call - for example:   

!$omp target data use_device_ptr(a0)
  call MPI_Allreduce(MPI_IN_PLACE,a0,n,ntype_real, MPI_SUM,comm_phi,ierr)
!$omp end target data

The code works correctly, but the compiler shows this remark:

pot3d.F90(5432): remark #8980: A list item in an is_device_ptr or a use_device_ptr clause must be of type C_PTR. The support is deprecated otherwise. [A0]
!$omp target data use_device_ptr(a0)
---------------------------------^

Since my arrays are declared as allocatable or automatic arrays and not declared as C_PTR (or even Fortran pointers), it would seem the alternative would be to use "use_device_addr()"?

However, using "use_device_addr()" would be a problem for portability since the NVIDIA NV compiler requires one to use "use_device_ptr()".   I have discussed this with others and apparently the OpenMP spec is a bit ambiguous about which should/could be used?

Could you please clarify when and where to use use_device_ptr() versus "use_device_addr()" in the context of a Fortran code calling external GPU-aware libraries like MPI?

How long will it be before the "support is deprecated otherwise" shown above will stop working with the code?  Any chance the deprecation could be put n hold and the use of use_device_ptr() be allowed for regular Fortran arrays?

Thanks!

 - Ron

[TobiasK · Accepted solution]

@caplanr the OpenMP spec is quite clear on that.

The usage is deprecated with 5.2 and removed in 6.0:

https://www.openmp.org/wp-content/uploads/OpenMP-API-Specification-5-2.pdf 5.2, page 122:

Fortran
1 If a list item that appears in a use_device_ptr clause is of type C_PTR and points to a data
2 entity that is mapped to the device data environment, references to the list item in the structured
3 block that is associated with the construct on which the clause appears are converted into references
4 to a device pointer that is local to the structured block and that refers to the device address of the
5 corresponding entity. If a list item of type C_PTR does not point to a mapped object, it must
6 contain a valid device address for the target device, and the list item references are instead
7 converted to references to a local device pointer that refers to this device address. If a list item in a
8 use_device_ptr clause is not of type C_PTR, the behavior is as if the list item appeared in a
9 use_device_addr clause. Support for such list items in a use_device_ptr clause is
10 deprecated.
Fortran

https://www.openmp.org/wp-content/uploads/OpenMP-API-Specification-6-0.pdf 6.0, page 236:

26 Restrictions
27 Restrictions to the use_device_ptr clause are as follows:
28 • Each list item must be a C pointer for which the value is the address of an object that has
29 corresponding storage or is accessible on the target device.

The correct way is to use use_device_addr as you already mentioned.

[Ron_Green]

@caplanr as for the question "How long will it be before the "support is deprecated otherwise" shown above will stop working with the code"?

I can't answer that.  At the moment we have no plans to reject or break this deprecated usage.  But as the OpenMP compilation evolves along with evolution of the OpenMP Runtime it is possible some change, directly addressing this usage OR some adjacent usage or optimization, in the future could cause a breakage.  But there is no way to estimate if or when that would happen.  

Were it my code, I would leave the deprecated code as-is for now.  Rather, I document in comments this deprecated usage and why (along with a date for this note). And put in the OMP 6.0 compliant code under comments so it's ready to uncomment when ready. Then ask Nvidia about their future plans to address the OMP 6.0 deprecation and schedule a code change around their change.  <sigh> of course, a new DEFINE and IFDEFs would be one way to address this.  But I hate IFDEFs unless absolutely necessary. 

[caplanr]

Thanks!

I will take your suggestion to leave it as-is for now and relay the issue to the NVIDIA message boards (I think there is a ticket for this already).

(I also try to avoid IFDEFS wherever and whenever possible )

 - Ron

[caplanr]

Hi,

FYI - it looks like NVIDIA will have a fix to allow the use of `device_addr` instead of `device_ptr` in the next release (or possibly the one after that) so the issue should be synchronized well soon.

https://forums.developer.nvidia.com/t/omp-target-data-use-device-ptr-vs-use-device-addr/290317/12

Thanks!

 - Ron