[OpenMP][NVPTX] Replace void** buffer by byte-wise buffer

[jdoerfert]

Summary

This commit implements the existing void buffer used to share
arguments between threads in a team with a byte-wise buffer. For now,
the void buffer is kept for compatibility.

The byte-wise buffer, if used directly, allows to save memory when small
arguments are shared between team threads. It does also allow to track
an additional offset that differentiates two distinct back-to-back
memory regions, e.g., for shared (copy in & out) and firstprivate (copy
in only) variables.

This is a preparation patch for https://reviews.llvm.org/D59319

[jdoerfert]

Fix the set/release use case

[jdoerfert]

openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h

73

What is this buffer used for? Transferring pointers to the shread variables to the parallel regions? If so, it must be handled by the compiler. There are several reasons to do this:

1. You're using malloc/free functions for large buffers. The fact is that the size of this buffer is known at the compile time and compiler can generate the fixed size buffer in the global memory if required. We already have similar implementation for target regions, globalized variables etc. You can take a look and adapt it for your purpose.
2. Malloc/free are not very fast on the GPU, so it will get an additional performance with the preallocated buffers.
3. Another one problem with malloc/free is that they are using preallocated memory and the size of this memory is limited by 8Mb (if I do recall correctly). This memory is required for the correct support of the local variables globalization and we alredy ran into the situation when malloc could not allocate enough memory for it with some previous implementations.
4. You can reused the shared memory buffers already generated by the compiler and save shared memory.

[Quote by ABataev copied from https://reviews.llvm.org/D59319?id=190767#inline-525900 after the patch was split.]

This buffer is supposed to be used to communicate variables in shared and firstprivate clauses between threads in a team. In this patch it is simply used to implement the old void** buffer. How, when, if we use it is part of the interface implementation. For now, this buffer simply serves the users of the omptarget_nvptx_globalArgs global.

If you want to provide compiler allocated memory to avoid the buffer use, no problem,
the __kmpc_target_region_kernel_parallel function allows to do so, see the SharedMemPointers flag. I wouldn't want to put the logic to generate these buffers in the front-end though.

[ABataev]

openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h

73

Why?

[jdoerfert]

openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h

73

Why what?

[ABataev]

openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h

73

Why you don't want to put the buffers generation to the compiler?

[jdoerfert]

openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h

73

Why you don't want to put the buffers generation to the compiler?

I did never say that. I explicitly explained how the new interface allows you to do exactly that. Maybe you confuse it with me not wanting the generation to be part of the front-end (=Clang). That is because it is an implementation choice for performance, as such it should not be done by Clang if it might obstruct analyses later on, or could be done better with more analysis support.

If we conclude we actually need to share values, after we tried to eliminate the need for sharing altogether, we can decide if a global buffer is preferable. If so, we can check if
there is already one that is unused or if we would need to create a new one.

Regarding this patch: It is not supposed to change the behavior at all. This patch just introduces a more general buffer which is then used to implement the old buffer. How/when the buffer is used is not affected.

[ABataev]

openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h

73

Only Clang can do better analysis for the shared variables, runtime is not. This part must be implemented in Clang, not in the runtime. I had no time to fix the existing implementation of parameters passing in non-SPMD mode. But if you working on this, you should definetely implement this in the compiler, not in the library.
It does not brea the analysis or anything else. It really produces better code with the better performance and less memory use.
If you're going to implement it, you need to implement it in the best possible way. Who else is going to fix this later, when we ran into the problems with the shared memory and/or malloced memory?

[jdoerfert]

openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h

73

Only Clang can do better analysis for the shared variables, runtime is not. This part must be implemented in Clang, not in the runtime. I had no time to fix the existing implementation of parameters passing in non-SPMD mode. But if you working on this, you should definetely implement this in the compiler, not in the library.

Clang is not the right place for analysis. That is what the patch set is all about.

I do not propose to only use the library implementation to handle sharing.

The initially generated code will utilize a library based solution that potentially calls malloc and memcpy. However, the later LLVM passes can change the behavior by providing a buffer that is accessible by all threads in the team, eliminating the need for malloc and memcpy.
This is a *generalization* of the scheme you currently employ, eliminating the need for mallocs completely (which is not the case right now).

It does not brea the analysis or anything else. It really produces better code with the better performance and less memory use.

Adding (global) state makes analysis results rarely better but almost always worse. If you do not believe me, please feel free to ask a compiler developer that you trust.

If you're going to implement it, you need to implement it in the best possible way.

I totally agree.

Who else is going to fix this later, when we ran into the problems with the shared memory and/or malloced memory?

Please read my comments again. I explain multiple times that:

A) This patch has nothing to do with your problem.
B) How the proposed TRegion interface tackles your problem.

[jdoerfert]

ping.

[gtbercea]

Could you check if there is any change in the number of registers new scheme vs. old scheme?

openmp/libomptarget/deviceRTLs/nvptx/src/omp_data.cu

 // Scratchpad for teams reduction.
 ////////////////////////////////////////////////////////////////////////////////
 __device__ __shared__ void *ReductionScratchpadPtr;
 
 ////////////////////////////////////////////////////////////////////////////////
 // Data sharing related variables.
 ////////////////////////////////////////////////////////////////////////////////
 __device__ __shared__ omptarget_nvptx_SharedArgs omptarget_nvptx_globalArgs;
+
+////////////////////////////////////////////////////////////////////////////////
+/// Pointer to share memory between team threads.
+////////////////////////////////////////////////////////////////////////////////
+__device__ __shared__ shared_bytes_buffer _shared_bytes_buffer_memory;
+

openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h

 #define __SHFL_DOWN_SYNC(mask, var, delta, width)                              \
   __shfl_down((var), (delta), (width))
 #define __ACTIVEMASK() __ballot(1)
 #endif
 
 #define __SYNCTHREADS_N(n) asm volatile("bar.sync %0;" : : "r"(n) : "memory");
 #define __SYNCTHREADS() __SYNCTHREADS_N(0)
 
+/// Helper structure to manage the memory shared by the threads in a team.
+///
+/// This buffer can manage two adjacent byte-wise objects by tracking the
+/// beginning of the second, as an offset, in addition to the beginning of the
+/// first, as a pointer.
+///
+/// Note: Only the team master is allowed to call non-const functions!
+struct shared_bytes_buffer {
+
+  INLINE void init() {
+    _ptr = &_data[0];
+    _size = PRE_SHARED_BYTES;
+    _offset = 0;
+  }
+
+  /// Release any dynamic allocated memory.
+  INLINE void release() {
+    if (_size != PRE_SHARED_BYTES)
+      SafeFree(_ptr, (char *)"free shared dynamic buffer");
+    // Always perform an init, it is cheap and required after a set call was
+    // performed during the last use of the buffer.
+    init();
+  }
+
+  INLINE void set(void *ptr, size_t offset) {
+    // Note that release will set _size to PRE_SHARED_BYTES, thereby avoiding
+    // the next release call from freeing the associated memory.
+    release();
+    _ptr = (char *)ptr;
+    _offset = offset;
+  }
+
+  INLINE void resize(size_t size, size_t offset) {
+    _offset = offset;
+
+    if (size <= _size)
+      return;
+
+    if (_size != PRE_SHARED_BYTES)
+      SafeFree(_ptr, (char *)"free shared dynamic buffer");
+
+    _size = size;
+    _ptr = (char *)SafeMalloc(_size, (char *)"new shared buffer");
+  }
+
+  // Called by all threads.
+  INLINE void *begin() const { return _ptr; };
+  INLINE size_t size() const { return _size; };
+  INLINE size_t get_offset() const { return _offset; };
+
+private:
+  // Pre-allocated space that holds PRE_SHARED_BYTES many bytes.
+  char _data[PRE_SHARED_BYTES];
+
+  // Pointer to the currently used buffer.
+  char *_ptr;
+
+  // Size of the currently used buffer.
+  uint32_t _size;
+
+  // Offset into the currently used buffer.
+  uint32_t _offset;
+};
+
+extern __device__ __shared__ shared_bytes_buffer _shared_bytes_buffer_memory;
+
 // arguments needed for L0 parallelism only.
+//
+// NOTE: Deprecated, use shared_byte_buffer instead.
 class omptarget_nvptx_SharedArgs {
 public:
   // All these methods must be called by the master thread only.
-  INLINE void Init() {
-    args  = buffer;
-    nArgs = MAX_SHARED_ARGS;
-  }
-  INLINE void DeInit() {
-    // Free any memory allocated for outlined parallel function with a large
-    // number of arguments.
-    if (nArgs > MAX_SHARED_ARGS) {
-      SafeFree(args, (char *)"new extended args");
-      Init();
-    }
-  }
+  INLINE void Init() { _shared_bytes_buffer_memory.init(); }
+  INLINE void DeInit() { _shared_bytes_buffer_memory.release(); }
   INLINE void EnsureSize(size_t size) {
-    if (size > nArgs) {
-      if (nArgs > MAX_SHARED_ARGS) {
-        SafeFree(args, (char *)"new extended args");
-      }
-      args = (void **) SafeMalloc(size * sizeof(void *),
-                                  (char *)"new extended args");
-      nArgs = size;
-    }
+    _shared_bytes_buffer_memory.resize(size * sizeof(void *), 0);
   }
   // Called by all threads.
-  INLINE void **GetArgs() const { return args; };
-private:
-  // buffer of pre-allocated arguments.
-  void *buffer[MAX_SHARED_ARGS];
-  // pointer to arguments buffer.
-  // starts off as a pointer to 'buffer' but can be dynamically allocated.
-  void **args;
-  // starts off as MAX_SHARED_ARGS but can increase in size.
-  uint32_t nArgs;
+  INLINE void **GetArgs() const {
+    return (void **)_shared_bytes_buffer_memory.begin();
+  };
 };
 
 extern __device__ __shared__ omptarget_nvptx_SharedArgs
     omptarget_nvptx_globalArgs;
 
 // Data sharing related quantities, need to match what is used in the compiler.
 enum DATA_SHARING_SIZES {
   // The maximum number of workers in a kernel.

openmp/libomptarget/deviceRTLs/nvptx/src/option.h

 #define MAX_THREADS_PER_TEAM 1024
 
 #define WARPSIZE 32
 
 // The named barrier for active parallel threads of a team in an L1 parallel
 // region to synchronize with each other.
 #define L1_BARRIER (1)
 
-// Maximum number of preallocated arguments to an outlined parallel/simd function.
-// Anything more requires dynamic memory allocation.
-#define MAX_SHARED_ARGS 20
+// Maximum number of preallocated bytes that can be passed to an outlined
+// parallel/simd function before dynamic memory allocation is required.
+#define PRE_SHARED_BYTES 128
 
 // Maximum number of omp state objects per SM allocated statically in global
 // memory.
 #if __CUDA_ARCH__ >= 700
 #define OMP_STATE_COUNT 32
 #define MAX_SM 84
 #elif __CUDA_ARCH__ >= 600
 #define OMP_STATE_COUNT 32