[OpenMP] Remove implicit data sharing using device shared memory from libomptarget

[gtbercea]

Summary

This patch reverts the changes to libomptarget that were coupled with the changes to Clang code gen for data sharing using shared memory. A similar patch exists for Clang: D43625

Shared memory is meant to be used as an optimization on top of a more general scheme. So far we didn't have a global memory implementation ready so shared memory was a solution which applied to the current level of OpenMP complexity supported by trunk on GPU devices (due to the missing NVPTX backend patch this functionality has never been exercised). Now that we have a global memory solution this patch is "in the way" and needs to be removed (for now). This patch (or an equivalent version of it) will be put out for review once the global memory scheme is in place.

[gtbercea]

ping

[grokos]

As far as I am concerned, I agree that the existing data sharing scheme should be removed in favour of a simpler future patch which will introduce the more generic data sharing scheme.

@Hahnfeld: Did our private discussion answer your questions about this plan?

[gtbercea]

In D43626#1024135, @grokos wrote:

As far as I am concerned, I agree that the existing data sharing scheme should be removed in favour of a simpler future patch which will introduce the more generic data sharing scheme.

@Hahnfeld: Did our private discussion answer your questions about this plan?

I could post the combined patch, the one which combines this patch with the one that adds the new functions. The new functions will contain a naive implementation of those methods but one which would work correctly when the Clang code gen lands. I'm just putting this here as an alternative.

[Hahnfeld]

In D43626#1024135, @grokos wrote:

@Hahnfeld: Did our private discussion answer your questions about this plan?

Partly, I'm still not happy about this turns out to be required after all of this was upstreamed. However if it's really needed you should at least make sure that basic applications continue to compile, something which we only reached recently.

In any case, the result of our private discussion needs to be made public. In particular I still don't see either explanation nor reasoning in any of the patches.

[gtbercea]

In D43626#1025517, @Hahnfeld wrote:

In D43626#1024135, @grokos wrote:

@Hahnfeld: Did our private discussion answer your questions about this plan?

Partly, I'm still not happy about this turns out to be required after all of this was upstreamed. However if it's really needed you should at least make sure that basic applications continue to compile, something which we only reached recently.

Data sharing cannot be used today due to the missing LLVM NVPTX backend patch so it should not affect applications that compile right now. This patch and it's Clang counterpart need to be committed at the same time to ensure continuity.

In any case, the result of our private discussion needs to be made public. In particular I still don't see either explanation nor reasoning in any of the patches.

The shared memory scheme was never meant to cover all the cases in OpenMP and it was meant more as an optimization. For that to be true, the global memory scheme needs to be the default scheme used. The shared memory scheme can then be re-intorduced on top of the global memory scheme as an optimization.

[Hahnfeld]

In D43626#1027103, @gtbercea wrote:

In D43626#1025517, @Hahnfeld wrote:

In D43626#1024135, @grokos wrote:

@Hahnfeld: Did our private discussion answer your questions about this plan?

Partly, I'm still not happy about this turns out to be required after all of this was upstreamed. However if it's really needed you should at least make sure that basic applications continue to compile, something which we only reached recently.

Data sharing cannot be used today due to the missing LLVM NVPTX backend patch so it should not affect applications that compile right now. This patch and it's Clang counterpart need to be committed at the same time to ensure continuity.

Got that but that doesn't justify upstreaming a part and kind of reverting it a few weeks later.

In any case, the result of our private discussion needs to be made public. In particular I still don't see either explanation nor reasoning in any of the patches.

The shared memory scheme was never meant to cover all the cases in OpenMP and it was meant more as an optimization. For that to be true, the global memory scheme needs to be the default scheme used. The shared memory scheme can then be re-intorduced on top of the global memory scheme as an optimization.

Again: You should put that in the summary of all patches!

[gtbercea]

In D43626#1027237, @Hahnfeld wrote:

In D43626#1027103, @gtbercea wrote:

In D43626#1025517, @Hahnfeld wrote:

In D43626#1024135, @grokos wrote:

@Hahnfeld: Did our private discussion answer your questions about this plan?

Partly, I'm still not happy about this turns out to be required after all of this was upstreamed. However if it's really needed you should at least make sure that basic applications continue to compile, something which we only reached recently.

Data sharing cannot be used today due to the missing LLVM NVPTX backend patch so it should not affect applications that compile right now. This patch and it's Clang counterpart need to be committed at the same time to ensure continuity.

Got that but that doesn't justify upstreaming a part and kind of reverting it a few weeks later.

Reverting this is not part of the initial plan of course. I do agree with you that this is not ideal. We have been working on this problem over the past couple of weeks and have come up with a better solution to the data sharing problem using global memory. We always meant to use global memory as a more generic solution due to the fact that it doesn't have the size constraints that shared memory does. A global memory solution is also more complicated because there is a lot more legwork required by the compiler in order to share variables correctly. This is why the shared memory solution was there first since it was "easier" to rely on the properties of the GPU shared memory.

In any case, the result of our private discussion needs to be made public. In particular I still don't see either explanation nor reasoning in any of the patches.

The shared memory scheme was never meant to cover all the cases in OpenMP and it was meant more as an optimization. For that to be true, the global memory scheme needs to be the default scheme used. The shared memory scheme can then be re-intorduced on top of the global memory scheme as an optimization.

Again: You should put that in the summary of all patches!

[gtbercea]

@Hahnfeld any further questions on this?

[Hahnfeld]

In D43626#1029879, @gtbercea wrote:

@Hahnfeld any further questions on this?

In D43626#1027237, @Hahnfeld wrote:

Again: You should put that in the summary of all patches!

Can't be that hard! I won't accept this patch anyway because I'm not interested anymore in this data sharing functionality.

[gtbercea]

In D43626#1030053, @Hahnfeld wrote:

In D43626#1029879, @gtbercea wrote:

@Hahnfeld any further questions on this?

In D43626#1027237, @Hahnfeld wrote:

Again: You should put that in the summary of all patches!

Can't be that hard! I won't accept this patch anyway because I'm not interested anymore in this data sharing functionality.

Sorry for mistakenly including you in the list of reviewers.

[Hahnfeld]

In D43626#1030077, @gtbercea wrote:

In D43626#1030053, @Hahnfeld wrote:

In D43626#1029879, @gtbercea wrote:

@Hahnfeld any further questions on this?

In D43626#1027237, @Hahnfeld wrote:

Again: You should put that in the summary of all patches!

Can't be that hard! I won't accept this patch anyway because I'm not interested anymore in this data sharing functionality.

Sorry for mistakenly including you in the list of reviewers.

Can you still please update the summary of all patches as I suggested?

[grokos]

I think the patch looks good. Let's remove the obsolete code in preparation for the new data-sharing scheme.

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, void ***SharedArgs,
+EXTERN bool __kmpc_kernel_parallel(void **WorkFn,
                                    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,

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

EXTERN void __kmpc_kernel_init_params(void *Ptr) {
 
   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;