[OpenMP][libomptarget] Add data sharing support in libomptarget

[gtbercea]

Summary

This patch extends the libomptarget functionality in patch D14254 with support for the data sharing scheme for supporting implicitly shared variables. The runtime therefore maintains a list of references to shared variables.

[Hahnfeld]

libomptarget/deviceRTLs/nvptx/src/interface.h

472–477

We have to agree on an order here: In D41012 I chose to put IsOMPRuntimeInitialized last which matches the other functions.

[gtbercea]

libomptarget/deviceRTLs/nvptx/src/interface.h

472–477

Is it ok if we go with the order in this patch?

[Hahnfeld]

libomptarget/deviceRTLs/nvptx/src/interface.h

472–477

In the end I don't really care, but IMO it's more consistent to have IsOMPRuntimeInitialized last in the list of arguments. That's the reason I submitted D41012 as it currently is.

[gtbercea]

libomptarget/deviceRTLs/nvptx/src/interface.h

472–477

Great! Let's use yours then and I can change the order of the args in this patch. Can you make sure you push the other patch and then once this one lands it should all work.

[gtbercea]

ping

[Hahnfeld]

Please rebase on top of D14254 and run clang-format.

libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu

35

I'm not sure if this needs to be defined here... Hopefully this lightens up after rebasing ;-)

[Hahnfeld]

libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu

35

Yeah, I think this needs be be declared extern here and defined in omp_data.cu.

[Hahnfeld]

libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu

35

Actually, all of these extern declarations should go away because they already exist in omptarget-nvptx.h! We don't need to reiterate them here...

[gtbercea]

Update against latest master branch.

[grokos]

Looks good. Thanks for submitting! I'll take care of duplicate extern declarations in another patch.

[Hahnfeld]

libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu

35

I still think this won't work: If there are multiple translation units with target regions, the linker should complain about duplicate symbol definitions... I think that's why there is a separate file omp_data.cu.

[Hahnfeld]

libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu

35

(obviously only when building and using the bclib which we don't yet have in trunk; but I think we shouldn't introduce the problem if we can easily avoid it)

[gtbercea]

libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu

35

Do you have a failing example for this?

Also I'm not sure what you want me to move where. The comments keep going back and forth.

[Hahnfeld]

libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu

35

Compiling and linking multiple source files with target regions should fail when using the bclib.

Just move the definition (ie without extern) to omp_data.cu. (Terminology taken from https://stackoverflow.com/a/1410632)

[gtbercea]

Move definition of data sharing variables.

[gtbercea]

@Hahnfeld any further comments?

[Hahnfeld]

In D41485#1000702, @gtbercea wrote:

@Hahnfeld any further comments?

No, look good!

libomptarget/deviceRTLs/nvptx/src/interface.h

 
 // non standard
 EXTERN void __kmpc_kernel_init_params(void *ReductionScratchpadPtr);
 EXTERN void __kmpc_kernel_init(int ThreadLimit, int16_t RequiresOMPRuntime);
 EXTERN void __kmpc_kernel_deinit(int16_t IsOMPRuntimeInitialized);
 EXTERN void __kmpc_spmd_kernel_init(int ThreadLimit, int16_t RequiresOMPRuntime,
                                     int16_t RequiresDataSharing);
 EXTERN void __kmpc_spmd_kernel_deinit();
 EXTERN void __kmpc_kernel_prepare_parallel(void *WorkFn,
+                                           void ***SharedArgs, int32_t nArgs,
                                            int16_t IsOMPRuntimeInitialized);
-EXTERN bool __kmpc_kernel_parallel(void **WorkFn,
+EXTERN bool __kmpc_kernel_parallel(void **WorkFn, void ***SharedArgs,
                                    int16_t IsOMPRuntimeInitialized);
 EXTERN void __kmpc_kernel_end_parallel();
 EXTERN bool __kmpc_kernel_convergent_parallel(void *buffer, uint32_t Mask,
                                               bool *IsFinal,
                                               int32_t *LaneSource);
 EXTERN void __kmpc_kernel_end_convergent_parallel(void *buffer);
 EXTERN bool __kmpc_kernel_convergent_simd(void *buffer, uint32_t Mask,
                                           bool *IsFinal, int32_t *LaneSource,
                                           int32_t *LaneId, int32_t *NumLanes);
 EXTERN void __kmpc_kernel_end_convergent_simd(void *buffer);

libomptarget/deviceRTLs/nvptx/src/omp_data.cu

 // Data sharing state
 ////////////////////////////////////////////////////////////////////////////////
 __device__ __shared__ DataSharingStateTy DataSharingState;
 
 ////////////////////////////////////////////////////////////////////////////////
 // Scratchpad for teams reduction.
 ////////////////////////////////////////////////////////////////////////////////
 __device__ __shared__ void *ReductionScratchpadPtr;
+
+////////////////////////////////////////////////////////////////////////////////
+// Data sharing related variables.
+////////////////////////////////////////////////////////////////////////////////
+__device__ __shared__ omptarget_nvptx_SharedArgs omptarget_nvptx_sharedArgs;

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

 #else
 #define __SHFL_SYNC(mask, var, srcLane) __shfl((var), (srcLane))
 #define __SHFL_DOWN_SYNC(mask, var, delta, width)                              \
   __shfl_down((var), (delta), (width))
 #define __BALLOT_SYNC(mask, predicate) __ballot((predicate))
 #define __ACTIVEMASK() __ballot(1)
 #endif
 
+// arguments needed for L0 parallelism only.
+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 EnsureSize(int 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;
+    }
+  }
+  // Called by all threads.
+  INLINE void **GetArgs() { 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;
+};
+
+extern __device__ __shared__ omptarget_nvptx_SharedArgs omptarget_nvptx_sharedArgs;
+
 // 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.
   DS_Max_Worker_Threads = 992,
   // The size reserved for data in a shared memory slot.
   DS_Slot_Size = 256,
   // The slot size that should be reserved for a working warp.
   DS_Worker_Warp_Slot_Size = WARPSIZE * DS_Slot_Size,

libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu

 //
 // The team master sets the outlined function and its arguments in these
 // variables to communicate with the workers.  Since they are in shared memory,
 // there is one copy of these variables for each kernel, instance, and team.
 //
 extern volatile __device__ __shared__ omptarget_nvptx_WorkFn
     omptarget_nvptx_workFn;
 extern __device__ __shared__ uint32_t execution_param;
 
 ////////////////////////////////////////////////////////////////////////////////
 // init entry points
 ////////////////////////////////////////////////////////////////////////////////
 
 INLINE unsigned smid() {
   unsigned id;
   asm("mov.u32 %0, %%smid;" : "=r"(id));
   return id;
 }
 
 EXTERN void __kmpc_kernel_init_params(void *Ptr) {
   PRINT(LD_IO, "call to __kmpc_kernel_init_params with version %f\n",
         OMPTARGET_NVPTX_VERSION);
 
   SetTeamsReductionScratchpadPtr(Ptr);
 }
 
 EXTERN void __kmpc_kernel_init(int ThreadLimit, int16_t RequiresOMPRuntime) {
   PRINT(LD_IO, "call to __kmpc_kernel_init with version %f\n",
         OMPTARGET_NVPTX_VERSION);
 
+  // init parallel work arguments
+  omptarget_nvptx_sharedArgs.Init();
+
   if (!RequiresOMPRuntime) {
     // If OMP runtime is not required don't initialize OMP state.
     setExecutionParameters(Generic, RuntimeUninitialized);
     return;
   }
   setExecutionParameters(Generic, RuntimeInitialized);
 
   int threadIdInBlock = GetThreadIdInBlock();

 #else
     int slot = smid() % MAX_SM;
 #endif
     omptarget_nvptx_device_State[slot].Enqueue(
         omptarget_nvptx_threadPrivateContext);
   }
   // Done with work.  Kill the workers.
   omptarget_nvptx_workFn = 0;
+
+  // Deinit parallel work arguments
+  omptarget_nvptx_sharedArgs.DeInit();
 }
 
 EXTERN void __kmpc_spmd_kernel_init(int ThreadLimit, int16_t RequiresOMPRuntime,
                                     int16_t RequiresDataSharing) {
   PRINT0(LD_IO, "call to __kmpc_spmd_kernel_init\n");
 
   if (!RequiresOMPRuntime) {
     // If OMP runtime is not required don't initialize OMP state.

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 omp state objects per SM allocated statically in global
 // memory.
 #if __CUDA_ARCH__ >= 600
 #define OMP_STATE_COUNT 32
 #define MAX_SM 56
 #else
 #define OMP_STATE_COUNT 16
 #define MAX_SM 16

libomptarget/deviceRTLs/nvptx/src/parallel.cu

 //        min (thread_limit*numLanes; blockDim.x)
 //      } else { // no thread_limit, no num_threads, use all cuda threads
 //        blockDim.x;
 //      }
 //    }
 //
 // This routine is always called by the team master..
 EXTERN void __kmpc_kernel_prepare_parallel(void *WorkFn,
+                                           void ***SharedArgs, int32_t nArgs,
                                            int16_t IsOMPRuntimeInitialized) {
   PRINT0(LD_IO, "call to __kmpc_kernel_prepare_parallel\n");
   omptarget_nvptx_workFn = WorkFn;
 
+  if (nArgs > 0) {
+    omptarget_nvptx_sharedArgs.EnsureSize(nArgs);
+    *SharedArgs = omptarget_nvptx_sharedArgs.GetArgs();
+  }
+
   if (!IsOMPRuntimeInitialized)
     return;
 
   // This routine is only called by the team master.  The team master is
   // the first thread of the last warp.  It always has the logical thread
   // id of 0 (since it is a shadow for the first worker thread).
   int threadId = 0;
   omptarget_nvptx_TaskDescr *currTaskDescr =

 }
 
 // All workers call this function.  Deactivate those not needed.
 // Fn - the outlined work function to execute.
 // returns True if this thread is active, else False.
 //
 // Only the worker threads call this routine.
 EXTERN bool __kmpc_kernel_parallel(void **WorkFn,
+                                   void ***SharedArgs,
                                    int16_t IsOMPRuntimeInitialized) {
   PRINT0(LD_IO | LD_PAR, "call to __kmpc_kernel_parallel\n");
 
   // Work function and arguments for L1 parallel region.
   *WorkFn = omptarget_nvptx_workFn;
+  *SharedArgs = omptarget_nvptx_sharedArgs.GetArgs();
 
   if (!IsOMPRuntimeInitialized)
     return true;
 
   // If this is the termination signal from the master, quit early.
   if (!*WorkFn)
     return false;