[libomptarget] [amdgpu] Fix default setting of max flat workgroup size

[dhruvachak]

Summary

When max flat workgroup size is not specified, it is set to the default
workgroup size. This prevents kernel launch with a workgroup size larger
than the default. The fix is to ignore a size of 0 and treat it as
unspecified.

[JonChesterfield]

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

1922

I guess this avoids a case where threadsPerGroup is set to zero by a zero ConstWGSize.

We've had a few bugs in this area now. The function is a map from a few integers to a few integers containing a lot of domain knowledge about deriving good launch parameters. It's only tested from the OpenMP level, where we can't easily perturb these values, so we miss stuff.

I think it's time to move all the launch value manipulation logic into this function, add a boolean to disable printing, and exhaustively test that function at the unit level.

We can either do that by making the function constexpr and writing a bunch of static asserts in this file or by declaring the prototype in a header and adding a standalone file that runs a bunch of values through it. Either way, we'd end up with a clear list of the cases it's intended to handle and much better odds of eliminating bugs in the edge cases.

[JonChesterfield]

Desire for tests notwithstanding, still want this point fix. Thanks!

[dhruvachak]

I guess this avoids a case where threadsPerGroup is set to zero by a zero ConstWGSize.

The scenario I am fixing is what happens in the default case today. No flat max-work-group size is intended, so it starts off as 0 but then we set ConstWGSize to 256 if the former is 0 (line 1710 above). So getLaunchVals limits threadsPerGroup to 256 because ConstWGSize is 256. My fix allows setting threadsPerGroup beyond 256 but limited to 1024.

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

if (err == HSA_STATUS_SUCCESS) {
       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);
 
       // 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 KernDescVal.WG_Size is 0, it is equivalent to not
+      // specified. Hence, max_flat_workgroup_size is filtered out in
+      // getLaunchVals
       WGSizeVal = KernDescVal.WG_Size;
       DP("WGSizeVal %d\n", WGSizeVal);
       check("Loading KernDesc computation property", err);
     } else {
       DP("Warning: Loading KernDesc '%s' - symbol not found, ", KernDescName);
 
       // Generic
       std::string ExecModeNameStr(e->name);
if (ExecutionMode == GENERIC) { // Add master warp for GENERIC
       DP("Adding master wavefront: +%d threads\n", RTLDeviceInfoTy::Warp_Size);
     }
     if (threadsPerGroup > RTLDeviceInfoTy::Max_WG_Size) { // limit to max
       threadsPerGroup = RTLDeviceInfoTy::Max_WG_Size;
       DP("Setting threads per block to maximum %d\n", threadsPerGroup);
     }
   }
   // check flat_max_work_group_size attr here
-  if (threadsPerGroup > ConstWGSize) {
+  if (ConstWGSize > 0 && threadsPerGroup > ConstWGSize) {
     threadsPerGroup = ConstWGSize;
     DP("Reduced threadsPerGroup to flat-attr-group-size limit %d\n",
        threadsPerGroup);
   }
   if (print_kernel_trace & STARTUP_DETAILS)
     fprintf(stderr, "threadsPerGroup: %d\n", threadsPerGroup);
   DP("Preparing %d threads\n", threadsPerGroup);