[Openmp-commits] [PATCH] D65340: [OpenMP][libomptarget] Add support for close map modifier

[Jonas Hahnfeld via Phabricator via Openmp-commits]

Hahnfeld added a comment.

In D65340#1623236 <https://reviews.llvm.org/D65340#1623236>, @gtbercea wrote:

> @Hahnfeld 
>  I would like to clarify something about tgt_register_requires call.
>
> This call is special and in a compiler that supports handling of requires clauses (like unified_shared_memory clause) ( i.e. Clang 9.0 onwards),  this function is called before any main function code. Here is a printout with LIBOMPTARGET_DEBUG=1 to realize the difference:
>
>   ===== WHERE COMPILER PLACES CALL TO __tgt_register_requires() ========
>   Libomptarget --> Loading RTLs... <<<< THIS IS THE VERY FIRST PRINT FROM THE RUNTIME. The call to __tgt_register_requires is just before that.
>   Libomptarget --> Loading library 'libomptarget.rtl.ppc64.so'...
>   Libomptarget --> Successfully loaded library 'libomptarget.rtl.ppc64.so'!
>   Libomptarget --> Registering RTL libomptarget.rtl.ppc64.so supporting 4 devices!
>   Libomptarget --> Loading library 'libomptarget.rtl.x86_64.so'...
>   Libomptarget --> Unable to load library 'libomptarget.rtl.x86_64.so': libomptarget.rtl.x86_64.so: cannot open shared object file: No such file or directory!
>   Libomptarget --> Loading library 'libomptarget.rtl.cuda.so'...
>   Target CUDA RTL --> Start initializing CUDA
>   Libomptarget --> Successfully loaded library 'libomptarget.rtl.cuda.so'!
>   Libomptarget --> Registering RTL libomptarget.rtl.cuda.so supporting 4 devices!
>   Libomptarget --> Loading library 'libomptarget.rtl.aarch64.so'...
>   Libomptarget --> Unable to load library 'libomptarget.rtl.aarch64.so': libomptarget.rtl.aarch64.so: cannot open shared object file: No such file or directory!
>   Libomptarget --> RTLs loaded!
>   Libomptarget --> Image 0x0000000010020080 is compatible with RTL libomptarget.rtl.ppc64.so!
>   Libomptarget --> RTL 0x000000004747b3e0 has index 0!
>   Libomptarget --> Registering image 0x0000000010020080 with RTL libomptarget.rtl.ppc64.so!
>   Libomptarget --> Done registering entries!
>   ===== MANUAL CALL TO __tgt_register_requires() IN MAIN (top of MAIN function!) ========
>
>
> Manually calling __tgt_register_requires can never actually work correctly because any user space call to this function will be after the binary has been loaded and the flags already decided. The compiler just has to support the handling of the requires clauses.


I disagree, the information about `requires` directives only needs to be present when initializing the device, not when loading a library. I've posted D66019 <https://reviews.llvm.org/D66019> to solve this and fix the test I mentioned in D65001 <https://reviews.llvm.org/D65001>.



================
Comment at: libomptarget/test/unified_shared_memory/close_manual.c:46-52
+// For the initalization of the device.
+#pragma omp target map(tofrom : device_alloc)
+  { device_alloc = 10; }
+
+  // CHECK: device_alloc updated on the device
+  if (device_alloc == 10)
+    printf("device_alloc updated on the device\n");
----------------
gtbercea wrote:
> Hahnfeld wrote:
> > Why do we need this?
> To initialize the device.
This should happen in `__tgt_target_data_begin` which can be a legitimate first directive in a program.


Repository:
  rOMP OpenMP

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D65340/new/

https://reviews.llvm.org/D65340