[OpenMP] Add dynamic memory function to omp.h and add documentation

[jhuber6]

Summary

This patch adds the llvm_omp_target_dynamic_shared_alloc function to
the omp.h header file so users can access it by default. Also changed
the name to keep it consistent with the other target allocators. Added
some documentation so users know how to use it. Didn't add the interface
for Fortran since there's no way to test it right now.

[jdoerfert]

openmp/docs/design/Runtimes.rst

1010

Can we link to the example?

openmp/libomptarget/src/exports

58

We might need to keep the old one for ABI compatibility. Not sure if we require that here though.

[jhuber6]

Making suggested changes.

[jdoerfert]

LG

openmp/docs/design/Runtimes.rst

 allocated using ``malloc`` and ``free`` for the CUDA plugin. This is necessary
 for some applications that allocate too much memory either through the user or
 globalization.
 
 LIBOMPTARGET_SHARED_MEMORY_SIZE
 """""""""""""""""""""""""""""""
 
 This environment variable sets the amount of dynamic shared memory in bytes used
-by the kernel once it is launched. A pointer to the dynamic memory buffer can
-currently only be accessed using the ``__kmpc_get_dynamic_shared`` device
-runtime call.
+by the kernel once it is launched. A pointer to the dynamic memory buffer can be
+accessed using the ``llvm_omp_target_dynamic_shared_alloc`` function. An example
+is shown in :ref:`libomptarget_dynamic_shared`.
 
 .. toctree::
    :hidden:
    :maxdepth: 1
 
    Offloading
 
 

 
 LLVM/OpenMP Target Device Runtime (``libomptarget-ARCH-SUBARCH.bc``)
 --------------------------------------------------------------------
 
 The target device runtime is an LLVM bitcode library that implements OpenMP
 runtime functions on the target device. It is linked with the device code's LLVM
 IR during compilation.
 
+.. _libomptarget_dynamic_shared:
+
+Dynamic Shared Memory
+^^^^^^^^^^^^^^^^^^^^^
+
+The target device runtime contains a pointer to the dynamic shared memory
+buffer. This pointer can be obtained using the
+``llvm_omp_target_dynamic_shared_alloc`` extension. If this function is called
+from the host it will simply return a null pointer. In order to use this buffer
+the kernel must be launched with an adequate amount of dynamic shared memory
+allocated. Currently this is done using the ``LIBOMPTARGET_SHARED_MEMORY_SIZE``
+environment variable. An example is given below.
+
+.. code-block:: c++
+
+    void foo() {
+      int x;
+    #pragma omp target parallel map(from : x)
+      {
+        int *buf = llvm_omp_target_dynamic_shared_alloc();
+    #pragma omp barrier
+        if (omp_get_thread_num() == 0)
+          *buf = 1;
+    #pragma omp barrier
+        if (omp_get_thread_num() == 1)
+          x = *buf;
+      }
+    }
+
+.. code-block:: console
+
+    $ clang++ -fopenmp -fopenmp-targets=nvptx64 shared.c
+    $ env LIBOMPTARGET_SHARED_MEMORY_SIZE=256 ./shared
+
 .. _libomptarget_device_debugging:
 
 Debugging
 ^^^^^^^^^
 
 The device runtime supports debugging in the runtime itself. This is configured
 at compile-time using the flag ``-fopenmp-target-debug=<N>`` rather than using a
 separate debugging build. If debugging is not enabled, the debugging paths will

openmp/libomptarget/DeviceRTL/include/Interface.h

 int omp_get_num_devices(void);
 
 int omp_get_num_teams(void);
 
 int omp_get_team_num();
 
 int omp_get_initial_device(void);
 
-void *llvm_omp_get_dynamic_shared();
+void *llvm_omp_target_dynamic_shared_alloc();
 
 /// Synchronization
 ///
 ///{
 void omp_init_lock(omp_lock_t *Lock);
 
 void omp_destroy_lock(omp_lock_t *Lock);
 

openmp/libomptarget/DeviceRTL/src/State.cpp

 
 __attribute__((noinline)) void __kmpc_free_shared(void *Ptr, uint64_t Bytes) {
   FunctionTracingRAII();
   memory::freeShared(Ptr, Bytes, "Frontend free shared");
 }
 
 void *__kmpc_get_dynamic_shared() { return memory::getDynamicBuffer(); }
 
+void *llvm_omp_target_dynamic_shared_alloc() {
+  return __kmpc_get_dynamic_shared();
+}
+
 void *llvm_omp_get_dynamic_shared() { return __kmpc_get_dynamic_shared(); }
 
 /// Allocate storage in shared memory to communicate arguments from the main
 /// thread to the workers in generic mode. If we exceed
 /// NUM_SHARED_VARIABLES_IN_SHARED_MEM we will malloc space for communication.
 constexpr uint64_t NUM_SHARED_VARIABLES_IN_SHARED_MEM = 64;
 
 [[clang::loader_uninitialized]] static void

openmp/libomptarget/include/omptarget.h

 
 /// Explicit target memory allocators
 /// Using the llvm_ prefix until they become part of the OpenMP standard.
 void *llvm_omp_target_alloc_device(size_t size, int device_num);
 void *llvm_omp_target_alloc_host(size_t size, int device_num);
 void *llvm_omp_target_alloc_shared(size_t size, int device_num);
 
 /// Dummy target so we have a symbol for generating host fallback.
-void *llvm_omp_get_dynamic_shared();
+void *llvm_omp_target_dynamic_shared_alloc();
 
 /// add the clauses of the requires directives in a given file
 void __tgt_register_requires(int64_t flags);
 
 /// adds a target shared library to the target execution image
 void __tgt_register_lib(__tgt_bin_desc *desc);
 
 /// Initialize all RTLs at once

openmp/libomptarget/src/api.cpp

 EXTERN void *llvm_omp_target_alloc_host(size_t size, int device_num) {
   return targetAllocExplicit(size, device_num, TARGET_ALLOC_HOST, __func__);
 }
 
 EXTERN void *llvm_omp_target_alloc_shared(size_t size, int device_num) {
   return targetAllocExplicit(size, device_num, TARGET_ALLOC_SHARED, __func__);
 }
 
+EXTERN void *llvm_omp_target_dynamic_shared_alloc() { return nullptr; }
 EXTERN void *llvm_omp_get_dynamic_shared() { return nullptr; }
 
 EXTERN void omp_target_free(void *device_ptr, int device_num) {
   TIMESCOPE();
   DP("Call to omp_target_free for device %d and address " DPxMOD "\n",
      device_num, DPxPTR(device_ptr));
 
   if (!device_ptr) {

openmp/libomptarget/src/exports

global:
     omp_target_is_present;
     omp_target_memcpy;
     omp_target_memcpy_rect;
     omp_target_associate_ptr;
     omp_target_disassociate_ptr;
     llvm_omp_target_alloc_host;
     llvm_omp_target_alloc_shared;
     llvm_omp_target_alloc_device;
-    llvm_omp_get_dynamic_shared;
+    llvm_omp_target_dynamic_shared_alloc;
     __tgt_set_info_flag;
     __tgt_print_device_info;
     omp_get_interop_ptr;
     omp_get_interop_str;
     omp_get_interop_int;
     omp_get_interop_name;
     omp_get_interop_type_desc;
     omp_get_interop_rc_desc;
     __tgt_interop_init;
     __tgt_interop_use;
     __tgt_interop_destroy;
   local:
     *;
 };
 

openmp/libomptarget/test/api/omp_dynamic_shared_memory.c

 // RUN: %libomptarget-compile-nvptx64-nvidia-cuda
 // RUN: env LIBOMPTARGET_SHARED_MEMORY_SIZE=256 \
 // RUN:   %libomptarget-run-nvptx64-nvidia-cuda | %fcheck-nvptx64-nvidia-cuda
 // REQUIRES: nvptx64-nvidia-cuda
 
 #include <omp.h>
 #include <stdio.h>
 
-void *llvm_omp_get_dynamic_shared();
-
 int main() {
   int x;
 #pragma omp target parallel map(from : x)
   {
-    int *buf = llvm_omp_get_dynamic_shared() + 252;
+    int *buf = llvm_omp_target_dynamic_shared_alloc() + 252;
 #pragma omp barrier
     if (omp_get_thread_num() == 0)
       *buf = 1;
 #pragma omp barrier
     if (omp_get_thread_num() == 1)
       x = *buf;
   }
 
   // CHECK: PASS
-  if (x == 1 && llvm_omp_get_dynamic_shared() == NULL)
+  if (x == 1 && llvm_omp_target_dynamic_shared_alloc() == NULL)
     printf("PASS\n");
 }

openmp/runtime/src/include/omp.h.var

#   if defined(_OPENMP) && _OPENMP >= 201811
     #pragma omp begin declare variant match(device={kind(nohost)})
     static inline int omp_is_initial_device(void) { return 0; }
     #pragma omp end declare variant
 #   endif
 
     /* OpenMP 5.2 */
     extern int __KAI_KMPC_CONVENTION omp_in_explicit_task(void);
 
+    /* LLVM Extensions */
+    extern void *llvm_omp_target_dynamic_shared_alloc();
+
 #   undef __KAI_KMPC_CONVENTION
 #   undef __KMP_IMP
 
     /* Warning:
        The following typedefs are not standard, deprecated and will be removed in a future release.
     */
     typedef int     omp_int_t;
     typedef double  omp_wtime_t;
 
 #   ifdef __cplusplus
     }
 #   endif
 
 #endif /* __OMP_H */