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

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.

Reviewed By: jdoerfert

Differential Revision: https://reviews.llvm.org/D123246

Author: jhuber6

openmp/docs/design/Runtimes.rst

@@ -1006,9 +1006,9 @@ 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:
@@ -1104,6 +1104,40 @@ 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

openmp/libomptarget/DeviceRTL/include/Interface.h

@@ -132,7 +132,7 @@ 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
 ///

openmp/libomptarget/DeviceRTL/src/State.cpp

@@ -530,6 +530,10 @@ __attribute__((noinline)) void __kmpc_free_shared(void *Ptr, uint64_t Bytes) {
 
 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

openmp/libomptarget/include/omptarget.h

@@ -226,7 +226,7 @@ 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);

openmp/libomptarget/src/api.cpp

@@ -53,6 +53,7 @@ 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) {

openmp/libomptarget/src/exports

@@ -40,7 +40,7 @@ VERS1.0 {
     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;

openmp/libomptarget/test/api/omp_dynamic_shared_memory.c

@@ -6,13 +6,11 @@
 #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;
@@ -22,6 +20,6 @@ int main() {
   }
 
   // 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

@@ -496,6 +496,9 @@
     /* 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