[libomptarget][AMDGPU] Remove MaxParallelLevel

[pdhaliwal]

Summary

Removes MaxParallelLevel references from rtl.cpp and drops
resulting dead code.

[JonChesterfield]

Very good, thank you.

Some context. Nested parallel is implemented by a transform to sequential execution by one parallel region. This transform uses linked call frames, allocated by malloc from the gpu within deviceRTL.

A previous version of amdgpu's deviceRTL had a buggy malloc implementation. The code that this patch deletes worked around that by allocating the gpu memory from the host instead.

That work around is now gone from the downstream fork (because malloc is no longer buggy), thus also removed from here.

openmp/libomptarget/plugins/amdgpu/src/rtl.cpp

 
 /// Use a single entity to encode a kernel and a set of flags
 struct KernelTy {
   // execution mode of kernel
   // 0 - SPMD mode (without master warp)
   // 1 - Generic mode (with master warp)
   int8_t ExecutionMode;
   int16_t ConstWGSize;
-  int8_t MaxParLevel;
   int32_t device_id;
   void *CallStackAddr;
   const char *Name;
 
-  KernelTy(int8_t _ExecutionMode, int16_t _ConstWGSize, int8_t _MaxParLevel,
-           int32_t _device_id, void *_CallStackAddr, const char *_Name,
+  KernelTy(int8_t _ExecutionMode, int16_t _ConstWGSize, int32_t _device_id,
+           void *_CallStackAddr, const char *_Name,
            uint32_t _kernarg_segment_size)
       : ExecutionMode(_ExecutionMode), ConstWGSize(_ConstWGSize),
-        MaxParLevel(_MaxParLevel), device_id(_device_id),
-        CallStackAddr(_CallStackAddr), Name(_Name) {
+        device_id(_device_id), CallStackAddr(_CallStackAddr), Name(_Name) {
     DP("Construct kernelinfo: ExecMode %d\n", ExecutionMode);
 
     std::string N(_Name);
     if (KernelArgPoolMap.find(N) == KernelArgPoolMap.end()) {
       KernelArgPoolMap.insert(
           std::make_pair(N, std::unique_ptr<KernelArgPool>(
                                 new KernelArgPool(_kernarg_segment_size))));
     }
for (__tgt_offload_entry *e = HostBegin; e != HostEnd; ++e) {
     }
 
     // default value GENERIC (in case symbol is missing from cubin file)
     int8_t ExecModeVal = ExecutionModeType::GENERIC;
 
     // get flat group size if present, else Default_WG_Size
     int16_t WGSizeVal = RTLDeviceInfoTy::Default_WG_Size;
 
-    // Max parallel level
-    int16_t MaxParLevVal = 0;
-
     // get Kernel Descriptor if present.
     // Keep struct in sync wih getTgtAttributeStructQTy in CGOpenMPRuntime.cpp
     struct KernDescValType {
       uint16_t Version;
       uint16_t TSize;
       uint16_t WG_Size;
       uint8_t Mode;
       uint8_t HostServices;
-      uint8_t MaxParallelLevel;
     };
     struct KernDescValType KernDescVal;
     std::string KernDescNameStr(e->name);
     KernDescNameStr += "_kern_desc";
     const char *KernDescName = KernDescNameStr.c_str();
 
     void *KernDescPtr;
     uint32_t KernDescSize;
if (err == ATMI_STATUS_SUCCESS) {
            sizeof(KernDescVal), KernDescVal.TSize, KernDescName);
 
       DP("After loading global for %s KernDesc \n", KernDescName);
       DP("KernDesc: Version: %d\n", KernDescVal.Version);
       DP("KernDesc: TSize: %d\n", KernDescVal.TSize);
       DP("KernDesc: WG_Size: %d\n", KernDescVal.WG_Size);
       DP("KernDesc: Mode: %d\n", KernDescVal.Mode);
       DP("KernDesc: HostServices: %x\n", KernDescVal.HostServices);
-      DP("KernDesc: MaxParallelLevel: %x\n", KernDescVal.MaxParallelLevel);
-
-      // gather location of callStack and size of struct
-      MaxParLevVal = KernDescVal.MaxParallelLevel;
-      if (MaxParLevVal > 0) {
-        uint32_t varsize;
-        const char *CsNam = "omptarget_nest_par_call_stack";
-        err = atmi_interop_hsa_get_symbol_info(place, CsNam, &CallStackAddr,
-                                               &varsize);
-        if (err != ATMI_STATUS_SUCCESS) {
-          fprintf(stderr, "Addr of %s failed\n", CsNam);
-          return NULL;
-        }
-        void *StructSizePtr;
-        const char *SsNam = "omptarget_nest_par_call_struct_size";
-        err = interop_get_symbol_info((char *)image->ImageStart, img_size,
-                                      SsNam, &StructSizePtr, &varsize);
-        if ((err != ATMI_STATUS_SUCCESS) ||
-            (varsize != sizeof(TgtStackItemSize))) {
-          fprintf(stderr, "Addr of %s failed\n", SsNam);
-          return NULL;
-        }
-        memcpy(&TgtStackItemSize, StructSizePtr, sizeof(TgtStackItemSize));
-        DP("Size of our struct is %d\n", TgtStackItemSize);
-      }
 
       // Get ExecMode
       ExecModeVal = KernDescVal.Mode;
       DP("ExecModeVal %d\n", ExecModeVal);
       if (KernDescVal.WG_Size == 0) {
         KernDescVal.WG_Size = RTLDeviceInfoTy::Default_WG_Size;
         DP("Setting KernDescVal.WG_Size to default %d\n", KernDescVal.WG_Size);
       }
if (err == ATMI_STATUS_SUCCESS) {
         DP("Warning: Loading WGSize '%s' - symbol not found, "
            "using default value %d\n",
            WGSizeName, WGSizeVal);
       }
 
       check("Loading WGSize computation property", err);
     }
 
-    KernelsList.push_back(KernelTy(ExecModeVal, WGSizeVal, MaxParLevVal,
-                                   device_id, CallStackAddr, e->name,
+    KernelsList.push_back(KernelTy(ExecModeVal, WGSizeVal, device_id,
+                                   CallStackAddr, e->name,
                                    kernarg_segment_size));
     __tgt_offload_entry entry = *e;
     entry.addr = (void *)&KernelsList.back();
     DeviceInfo.addOffloadEntry(device_id, entry);
     DP("Entry point %ld maps to %s\n", e - HostBegin, e->name);
   }
 
   return DeviceInfo.getOffloadEntriesTable(device_id);
if (print_kernel_trace > 1) {
     fprintf(stderr, "threadsPerGroup: %d\n", threadsPerGroup);
     fprintf(stderr, "num_groups: %d\n", num_groups);
     fprintf(stderr, "loop_tripcount: %ld\n", loop_tripcount);
   }
   DP("Final %d num_groups and %d threadsPerGroup\n", num_groups,
      threadsPerGroup);
 }
 
-static void *AllocateNestedParallelCallMemory(int MaxParLevel, int NumGroups,
-                                              int ThreadsPerGroup,
-                                              int device_id,
-                                              void *CallStackAddr, int SPMD) {
-  if (print_kernel_trace > 1)
-    fprintf(stderr, "MaxParLevel %d SPMD %d NumGroups %d NumThrds %d\n",
-            MaxParLevel, SPMD, NumGroups, ThreadsPerGroup);
-  // Total memory needed is Teams * Threads * ParLevels
-  size_t NestedMemSize =
-      MaxParLevel * NumGroups * ThreadsPerGroup * TgtStackItemSize * 4;
-
-  if (print_kernel_trace > 1)
-    fprintf(stderr, "NestedMemSize %ld \n", NestedMemSize);
-  assert(device_id < DeviceInfo.NumberOfDevices && "Device ID too large");
-  void *TgtPtr = NULL;
-  atmi_status_t err =
-      atmi_malloc(&TgtPtr, NestedMemSize, get_gpu_mem_place(device_id));
-  err = DeviceInfo.freesignalpool_memcpy_h2d(CallStackAddr, &TgtPtr,
-                                             sizeof(void *), device_id);
-  if (print_kernel_trace > 2)
-    fprintf(stderr, "CallSck %lx TgtPtr %lx *TgtPtr %lx \n",
-            (long)CallStackAddr, (long)&TgtPtr, (long)TgtPtr);
-  if (err != ATMI_STATUS_SUCCESS) {
-    fprintf(stderr, "Mem not wrtten to target, err %d\n", err);
-  }
-  return TgtPtr; // we need to free this after kernel.
-}
-
 static uint64_t acquire_available_packet_id(hsa_queue_t *queue) {
   uint64_t packet_id = hsa_queue_add_write_index_relaxed(queue, 1);
   bool full = true;
   while (full) {
     full =
         packet_id >= (queue->size + hsa_queue_load_read_index_scacquire(queue));
   }
   return packet_id;
int32_t __tgt_rtl_run_target_team_region(int32_t device_id, void *tgt_entry_ptr,
   DeviceInfo.load_run_lock.unlock_shared();
   return res;
 }
 
 int32_t __tgt_rtl_run_target_team_region_locked(
     int32_t device_id, void *tgt_entry_ptr, void **tgt_args,
     ptrdiff_t *tgt_offsets, int32_t arg_num, int32_t num_teams,
     int32_t thread_limit, uint64_t loop_tripcount) {
-  static pthread_mutex_t nested_parallel_mutex = PTHREAD_MUTEX_INITIALIZER;
-
   // Set the context we are using
   // update thread limit content in gpu memory if un-initialized or specified
   // from host
 
   DP("Run target team region thread_limit %d\n", thread_limit);
 
   // All args are references.
   std::vector<void *> args(arg_num);
int32_t __tgt_rtl_run_target_team_region_locked(
   getLaunchVals(threadsPerGroup, num_groups, KernelInfo->ConstWGSize,
                 KernelInfo->ExecutionMode, DeviceInfo.EnvTeamLimit,
                 DeviceInfo.EnvNumTeams,
                 num_teams,     // From run_region arg
                 thread_limit,  // From run_region arg
                 loop_tripcount // From run_region arg
   );
 
-  void *TgtCallStack = NULL;
-  if (KernelInfo->MaxParLevel > 0) {
-    pthread_mutex_lock(&nested_parallel_mutex);
-    TgtCallStack = AllocateNestedParallelCallMemory(
-        KernelInfo->MaxParLevel, num_groups, threadsPerGroup,
-        KernelInfo->device_id, KernelInfo->CallStackAddr,
-        KernelInfo->ExecutionMode);
-  }
   if (print_kernel_trace > 0)
     // enum modes are SPMD, GENERIC, NONE 0,1,2
     fprintf(stderr,
             "DEVID:%2d SGN:%1d ConstWGSize:%-4d args:%2d teamsXthrds:(%4dX%4d) "
             "reqd:(%4dX%4d) n:%s\n",
             device_id, KernelInfo->ExecutionMode, KernelInfo->ConstWGSize,
             arg_num, num_groups, threadsPerGroup, num_teams, thread_limit,
             KernelInfo->Name);
while (hsa_signal_wait_scacquire(packet->completion_signal,
       ;
 
     assert(ArgPool);
     ArgPool->deallocate(packet->kernarg_address);
     DeviceInfo.FreeSignalPool.push(packet->completion_signal);
   }
 
   DP("Kernel completed\n");
-  // Free call stack for nested
-  if (TgtCallStack) {
-    pthread_mutex_unlock(&nested_parallel_mutex);
-    atmi_free(TgtCallStack);
-  }
-
   return OFFLOAD_SUCCESS;
 }
 
 int32_t __tgt_rtl_run_target_region(int32_t device_id, void *tgt_entry_ptr,
                                     void **tgt_args, ptrdiff_t *tgt_offsets,
                                     int32_t arg_num) {
   // use one team and one thread
   // fix thread num