[openmp][nfc] Refactor GridValues

[JonChesterfield]

Summary

Remove redundant fields and replace pointer with virtual function

Of fourteen fields, three are dead and four can be computed from the
remainder. This leaves a couple of currently dead fields in place as
they are expected to be used from the deviceRTL shortly. Two of the
fields that can be computed are only used from codegen and require a
log2() implementation so are inlined into codegen instead.

This change leaves the new methods in the same location in the struct
as the previous fields for convenience at review.

[JonChesterfield]

• reorder field

[jdoerfert]

llvm/include/llvm/Frontend/OpenMP/OMPGridValues.h

102

It should be in the device rtl then, no?

[JonChesterfield]

llvm/include/llvm/Frontend/OpenMP/OMPGridValues.h

102

This header is currently used from clang and the (amdgpu, could also be cuda if we like) host plugin. Possibly also from llvm. As of D108391 it would be used from the devicertl.

The idea is to have a single source of truth for the various magic numbers that the pieces should agree on and llvm is the common point on the dependency tree. I'm currently interested in that because I want to change some of them for gfx10 and have that magically ripple through the components. I'm not totally confident that will work out nicely for the host plugin as it has to dynamically handle different architectures but I think it'll be good enough.

It's not totally ideal to hand spin a function that is in the math support header but I also don't want to try to make various llvm headers ffreestanding-safe.

[jdoerfert]

llvm/include/llvm/Frontend/OpenMP/OMPGridValues.h

102

It's not totally ideal to hand spin a function that is in the math support header but I also don't want to try to make various llvm headers ffreestanding-safe.

The function is only needed in the device rtl. Put it in the device rtl.

[JonChesterfield]

llvm/include/llvm/Frontend/OpenMP/OMPGridValues.h

102

This particular function is only called by warpSizeLog2 which is currently only used by CGOpenMPRuntimeGPU. The deviceRTL doesn't call the function. However if this header includes the rest of llvm support then it can't call any of the others either.

I'm going to drop the log2 accessor (and this function) in favour of two calls into math support from CGOpenMPRuntime.

[JonChesterfield]

• delete log2 accessors per review comments

[JonChesterfield]

• whitespace, drop asserts

[jdoerfert]

LG

[JonChesterfield]

Failed a nvptx codegen test (maybe the change to calculate log2 at runtime), currently away from my desk but will revert when I get back unless beaten to it.

[JonChesterfield]

clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp

550

Bug is here. ~0 >> 27u == -1 (bad) and ~0u >> 27u == 31 (good). Win for exact codegen tests.

[JonChesterfield]

• require unsigned shift

clang/include/clang/Basic/TargetInfo.h

protected:
   unsigned char MaxAtomicPromoteWidth, MaxAtomicInlineWidth;
   unsigned short SimdDefaultAlign;
   std::string DataLayoutString;
   const char *UserLabelPrefix;
   const char *MCountName;
   unsigned char RegParmMax, SSERegParmMax;
   TargetCXXABI TheCXXABI;
   const LangASMap *AddrSpaceMap;
-  const llvm::omp::GV *GridValues =
-      nullptr; // target-specific GPU grid values that must be
-               // consistent between host RTL (plugin), device RTL, and clang.
 
   mutable StringRef PlatformName;
   mutable VersionTuple PlatformMinVersion;
 
   unsigned HasAlignMac68kSupport : 1;
   unsigned RealTypeUsesObjCFPRet : 3;
   unsigned ComplexLongDoubleUsesFP2Ret : 1;
 

   /// Return an AST address space which can be used opportunistically
   /// for constant global memory. It must be possible to convert pointers into
   /// this address space to LangAS::Default. If no such address space exists,
   /// this may return None, and such optimizations will be disabled.
   virtual llvm::Optional<LangAS> getConstantAddressSpace() const {
     return LangAS::Default;
   }
 
-  /// Return a target-specific GPU grid values
-  const llvm::omp::GV &getGridValue() const {
-    assert(GridValues != nullptr && "GridValues not initialized");
-    return *GridValues;
+  // access target-specific GPU grid values that must be consistent between
+  // host RTL (plugin), deviceRTL and clang.
+  virtual const llvm::omp::GV &getGridValue() const {
+    llvm_unreachable("getGridValue not implemented on this target");
   }
 
   /// Retrieve the name of the platform as it is used in the
   /// availability attribute.
   StringRef getPlatformName() const { return PlatformName; }
 
   /// Retrieve the minimum desired version of the platform, to
   /// which the program should be compiled.

clang/lib/Basic/Targets/AMDGPU.h

public:
   LangAS getCUDABuiltinAddressSpace(unsigned AS) const override {
     return LangAS::Default;
   }
 
   llvm::Optional<LangAS> getConstantAddressSpace() const override {
     return getLangASFromTargetAS(Constant);
   }
 
+  const llvm::omp::GV &getGridValue() const override {
+    return llvm::omp::AMDGPUGridValues;
+  }
+
   /// \returns Target specific vtbl ptr address space.
   unsigned getVtblPtrAddressSpace() const override {
     return static_cast<unsigned>(Constant);
   }
 
   /// \returns If a target requires an address within a target specific address
   /// space \p AddressSpace to be converted in order to be used, then return the
   /// corresponding target specific DWARF address space.

clang/lib/Basic/Targets/AMDGPU.cpp

 
 #include "AMDGPU.h"
 #include "clang/Basic/Builtins.h"
 #include "clang/Basic/CodeGenOptions.h"
 #include "clang/Basic/LangOptions.h"
 #include "clang/Basic/MacroBuilder.h"
 #include "clang/Basic/TargetBuiltins.h"
 #include "llvm/ADT/StringSwitch.h"
-#include "llvm/Frontend/OpenMP/OMPGridValues.h"
 
 using namespace clang;
 using namespace clang::targets;
 
 namespace clang {
 namespace targets {
 
 // If you edit the description strings, make sure you update
: TargetInfo(Triple),
       GPUKind(isAMDGCN(Triple) ?
               llvm::AMDGPU::parseArchAMDGCN(Opts.CPU) :
               llvm::AMDGPU::parseArchR600(Opts.CPU)),
       GPUFeatures(isAMDGCN(Triple) ?
                   llvm::AMDGPU::getArchAttrAMDGCN(GPUKind) :
                   llvm::AMDGPU::getArchAttrR600(GPUKind)) {
   resetDataLayout(isAMDGCN(getTriple()) ? DataLayoutStringAMDGCN
                                         : DataLayoutStringR600);
-  GridValues = &llvm::omp::AMDGPUGridValues;
 
   setAddressSpaceMap(Triple.getOS() == llvm::Triple::Mesa3D ||
                      !isAMDGCN(Triple));
   UseAddrSpaceMapMangling = true;
 
   HasLegalHalfType = true;
   HasFloat16 = true;
   WavefrontSize = GPUFeatures & llvm::AMDGPU::FEATURE_WAVE32 ? 32 : 64;

clang/lib/Basic/Targets/NVPTX.h

void setSupportedOpenCLOpts() override {
     Opts["__opencl_c_fp64"] = true;
     Opts["cl_khr_byte_addressable_store"] = true;
     Opts["cl_khr_global_int32_base_atomics"] = true;
     Opts["cl_khr_global_int32_extended_atomics"] = true;
     Opts["cl_khr_local_int32_base_atomics"] = true;
     Opts["cl_khr_local_int32_extended_atomics"] = true;
   }
 
+  const llvm::omp::GV &getGridValue() const override {
+    return llvm::omp::NVPTXGridValues;
+  }
+
   /// \returns If a target requires an address within a target specific address
   /// space \p AddressSpace to be converted in order to be used, then return the
   /// corresponding target specific DWARF address space.
   ///
   /// \returns Otherwise return None and no conversion will be emitted in the
   /// DWARF.
   Optional<unsigned>
   getDWARFAddressSpace(unsigned AddressSpace) const override {

clang/lib/Basic/Targets/NVPTX.cpp

 //===----------------------------------------------------------------------===//
 
 #include "NVPTX.h"
 #include "Targets.h"
 #include "clang/Basic/Builtins.h"
 #include "clang/Basic/MacroBuilder.h"
 #include "clang/Basic/TargetBuiltins.h"
 #include "llvm/ADT/StringSwitch.h"
-#include "llvm/Frontend/OpenMP/OMPGridValues.h"
 
 using namespace clang;
 using namespace clang::targets;
 
 const Builtin::Info NVPTXTargetInfo::BuiltinInfo[] = {
 #define BUILTIN(ID, TYPE, ATTRS)                                               \
   {#ID, TYPE, ATTRS, nullptr, ALL_LANGUAGES, nullptr},
 #define LIBBUILTIN(ID, TYPE, ATTRS, HEADER)                                    \
PTXVersion = llvm::StringSwitch<unsigned>(Feature)
                      .Case("+ptx40", 40)
                      .Case("+ptx32", 32)
                      .Default(32);
   }
 
   TLSSupported = false;
   VLASupported = false;
   AddrSpaceMap = &NVPTXAddrSpaceMap;
-  GridValues = &llvm::omp::NVPTXGridValues;
   UseAddrSpaceMapMangling = true;
 
   // Define available target features
   // These must be defined in sorted order!
   NoAsmVariants = true;
   GPU = CudaArch::SM_20;
 
   if (TargetPointerWidth == 32)

clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp

 };
 
 /// GPU Configuration:  This information can be derived from cuda registers,
 /// however, providing compile time constants helps generate more efficient
 /// code.  For all practical purposes this is fine because the configuration
 /// is the same for all known NVPTX architectures.
 enum MachineConfiguration : unsigned {
   /// See "llvm/Frontend/OpenMP/OMPGridValues.h" for various related target
-  /// specific Grid Values like GV_Warp_Size, GV_Warp_Size_Log2,
-  /// and GV_Warp_Size_Log2_Mask.
+  /// specific Grid Values like GV_Warp_Size, GV_Slot_Size
 
   /// Global memory alignment for performance.
   GlobalMemoryAlignment = 128,
 
   /// Maximal size of the shared memory buffer.
   SharedMemorySize = 128,
 };
 

 };
 } // anonymous namespace
 
 /// Get the id of the warp in the block.
 /// We assume that the warp size is 32, which is always the case
 /// on the NVPTX device, to generate more efficient code.
 static llvm::Value *getNVPTXWarpID(CodeGenFunction &CGF) {
   CGBuilderTy &Bld = CGF.Builder;
-  unsigned LaneIDBits = CGF.getTarget().getGridValue().GV_Warp_Size_Log2;
+  unsigned LaneIDBits = CGF.getTarget().getGridValue().warpSizeLog2();
   auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
   return Bld.CreateAShr(RT.getGPUThreadID(CGF), LaneIDBits, "nvptx_warp_id");
 }
 
 /// Get the id of the current lane in the Warp.
 /// We assume that the warp size is 32, which is always the case
 /// on the NVPTX device, to generate more efficient code.
 static llvm::Value *getNVPTXLaneID(CodeGenFunction &CGF) {
   CGBuilderTy &Bld = CGF.Builder;
   unsigned LaneIDMask =
-      CGF.getContext().getTargetInfo().getGridValue().GV_Warp_Size_Log2_Mask;
+      CGF.getContext().getTargetInfo().getGridValue().warpSizeLog2Mask();
   auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
   return Bld.CreateAnd(RT.getGPUThreadID(CGF), Bld.getInt32(LaneIDMask),
                        "nvptx_lane_id");
 }
 
 CGOpenMPRuntimeGPU::ExecutionMode
 CGOpenMPRuntimeGPU::getExecutionMode() const {
   return CurrentExecutionMode;
 }
 

llvm/include/llvm/Frontend/OpenMP/OMPGridValues.h

 struct GV {
   /// The maximum number of workers in a kernel.
   /// (THREAD_ABSOLUTE_LIMIT) - (GV_Warp_Size), might be issue for blockDim.z
   const unsigned GV_Threads;
   /// The size reserved for data in a shared memory slot.
   const unsigned GV_Slot_Size;
   /// The default value of maximum number of threads in a worker warp.
   const unsigned GV_Warp_Size;
-  /// Alternate warp size for some AMDGCN architectures. Same as GV_Warp_Size
-  /// for NVPTX.
-  const unsigned GV_Warp_Size_32;
+
   /// The number of bits required to represent the max number of threads in warp
-  const unsigned GV_Warp_Size_Log2;
-  /// GV_Warp_Size * GV_Slot_Size,
-  const unsigned GV_Warp_Slot_Size;
+  constexpr unsigned warpSizeLog2() const { return log2(GV_Warp_Size); }
+
+  constexpr unsigned warpSlotSize() const {
+    return GV_Warp_Size * GV_Slot_Size;
+  }
+
   /// the maximum number of teams.
   const unsigned GV_Max_Teams;
-  /// Global Memory Alignment
-  const unsigned GV_Mem_Align;
-  /// (~0u >> (GV_Warp_Size - GV_Warp_Size_Log2))
-  const unsigned GV_Warp_Size_Log2_Mask;
+
+  constexpr unsigned warpSizeLog2Mask() const {
+    return ~0u >> (32u - warpSizeLog2());
+  }
+
   // An alternative to the heavy data sharing infrastructure that uses global
   // memory is one that uses device __shared__ memory.  The amount of such space
   // (in bytes) reserved by the OpenMP runtime is noted here.
   const unsigned GV_SimpleBufferSize;
   // The absolute maximum team size for a working group
   const unsigned GV_Max_WG_Size;
   // The default maximum team size for a working group
   const unsigned GV_Default_WG_Size;
-  // This is GV_Max_WG_Size / GV_WarpSize. 32 for NVPTX and 16 for AMDGCN.
-  const unsigned GV_Max_Warp_Number;
-  /// The slot size that should be reserved for a working warp.
-  /// (~0u >> (GV_Warp_Size - GV_Warp_Size_Log2))
-  const unsigned GV_Warp_Size_Log2_MaskL;
+
+  constexpr unsigned maxWarpNumber() const {
+    return GV_Max_WG_Size / GV_Warp_Size;
+  }
+
+private:
+  static constexpr unsigned log2(unsigned I) {
+    // assumes I is nonzero power of 2
+    // reimplemented here for use from freestanding devicertl
+    unsigned R = 0;
+    while (I >>= 1) {
+      R++;
+    }
+    return R;
+  }
 };
 
 /// For AMDGPU GPUs
 static constexpr GV AMDGPUGridValues = {
-    448,       // GV_Threads
-    256,       // GV_Slot_Size
-    64,        // GV_Warp_Size
-    32,        // GV_Warp_Size_32
-    6,         // GV_Warp_Size_Log2
-    64 * 256,  // GV_Warp_Slot_Size
-    128,       // GV_Max_Teams
-    256,       // GV_Mem_Align
-    63,        // GV_Warp_Size_Log2_Mask
-    896,       // GV_SimpleBufferSize
-    1024,      // GV_Max_WG_Size,
-    256,       // GV_Defaut_WG_Size
-    1024 / 64, // GV_Max_WG_Size / GV_WarpSize
-    63         // GV_Warp_Size_Log2_MaskL
+    448,  // GV_Threads
+    256,  // GV_Slot_Size
+    64,   // GV_Warp_Size
+    128,  // GV_Max_Teams
+    896,  // GV_SimpleBufferSize
+    1024, // GV_Max_WG_Size,
+    256,  // GV_Default_WG_Size
 };
 
+static_assert(6 == AMDGPUGridValues.warpSizeLog2(), "");
+static_assert(64 * 256 == AMDGPUGridValues.warpSlotSize(), "");
+static_assert(63 == AMDGPUGridValues.warpSizeLog2Mask(), "");
+static_assert(1024 / 64 == AMDGPUGridValues.maxWarpNumber(), "");
+
 /// For Nvidia GPUs
 static constexpr GV NVPTXGridValues = {
-    992,               // GV_Threads
-    256,               // GV_Slot_Size
-    32,                // GV_Warp_Size
-    32,                // GV_Warp_Size_32
-    5,                 // GV_Warp_Size_Log2
-    32 * 256,          // GV_Warp_Slot_Size
-    1024,              // GV_Max_Teams
-    256,               // GV_Mem_Align
-    (~0u >> (32 - 5)), // GV_Warp_Size_Log2_Mask
-    896,               // GV_SimpleBufferSize
-    1024,              // GV_Max_WG_Size
-    128,               // GV_Defaut_WG_Size
-    1024 / 32,         // GV_Max_WG_Size / GV_WarpSize
-    31                 // GV_Warp_Size_Log2_MaskL
+    992,  // GV_Threads
+    256,  // GV_Slot_Size
+    32,   // GV_Warp_Size
+    1024, // GV_Max_Teams
+    896,  // GV_SimpleBufferSize
+    1024, // GV_Max_WG_Size
+    128,  // GV_Default_WG_Size
 };
 
+static_assert(5 == NVPTXGridValues.warpSizeLog2(), "");
+static_assert(32 * 256 == NVPTXGridValues.warpSlotSize(), "");
+static_assert((~0u >> (32 - 5)) == NVPTXGridValues.warpSizeLog2Mask(), "");
+static_assert(1024 / 32 == NVPTXGridValues.maxWarpNumber(), "");
+
 } // namespace omp
 } // namespace llvm
 
 #endif // LLVM_FRONTEND_OPENMP_OMPGRIDVALUES_H