[OpenMP][AMDGCN] Support OpenMP offloading for AMDGCN architecture - Part 2

[saiislam]

Summary

New file include to support platform dependent grid constants. It will be
used by clang, libomptarget plugins, and deviceRTLs to access constant
values consistently and with fast access in the deviceRTLs.

Originally authored by Greg Rodgers (@gregrodgers).

[arsenm]

clang/include/clang/Basic/OpenMPGridValues.h

96

What about wave32?

clang/include/clang/Basic/TargetInfo.h

1315

long long is the worst C type and should never be used. int64_t?

clang/lib/Basic/Targets/AMDGPU.cpp

290

Why do these need casts?

[jdoerfert]

clang/include/clang/Basic/OpenMPGridValues.h

134

Can we move this into llvm core (subfolder Frontend/OpenMP/)? We probably need this during OpenMPOpt, and for Flang as well.

[saiislam]

1. Added wave32 as an additional alternate wave size.
2. Replaced long long with uint64_t and int with unsigned.
3. Removed unnecessary casts.
4. Moved OpenMPGridValues.h to Frontend/OpenMP.
5. Refactored OpenMPGridValues.h as OMPGridValues.h to follow style of neighboring files.

[arsenm]

clang/lib/Basic/Targets/AMDGPU.cpp

290

Don't need parentheses

clang/lib/Basic/Targets/NVPTX.cpp

66

Don't need parentheses

[jdoerfert]

clang/include/clang/Basic/TargetInfo.h

202

I would recommend adding a descriptive comment, even though that seems not to be the norm here.
Why do we split these into "normal" and "long" grid values again?

= nullptr and an assert in the getters?

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

20 ↗

(On Diff #268264)

I would recommend omp not GPU, or omp before gpu.

[saiislam]

1. Renamed the GPU namespace as omp.
2. Added assert in getter functions.
3. Added comments.
4. Removed unnecassary parantheses.

[saiislam]

clang/include/clang/Basic/TargetInfo.h

202

Realized no practical use of separate long values. Folded into regular enum and arrays.

[jdoerfert]

clang/include/clang/Basic/TargetInfo.h

1317

I somehow believe !GridValues is the wrong way around. You should also initialize it directly to null so this makes more sense.

The return value and the type of gridvalues don't match, pick int or unsigned ;)

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

108 ↗

(On Diff #269314)

Nits:
-This is
I don't understand the fixme, do we need that?

[saiislam]

1. Corrected return type of getGridValue to unsigned
2. Fixed GridValues assert issue
3. Edited comments

[saiislam]

Ping.

[jdoerfert]

I am fine with this. @arsenm feel free to accept.

---

In D80917#2085138, @saiislam wrote:

Ping.

Afaik, pings should be reserved to ~1 week of inactivity. I know this can be frustrating, but we also need to consider noise for reviewers.

[saiislam]

In D80917#2085377, @jdoerfert wrote:

I am fine with this. @arsenm feel free to accept.

---

In D80917#2085138, @saiislam wrote:

Ping.

Afaik, pings should be reserved to ~1 week of inactivity. I know this can be frustrating, but we also need to consider noise for reviewers.

Ok. I will keep in mind next time. Thanks :-)

[JonChesterfield]

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

66 ↗

(On Diff #269927)

What's the point of warp_size_32? It's always set to 32 and seems redundant with warp_size

68 ↗

(On Diff #269927)

log2 of a compile time constant would be computed at compile time, we don't need this field

76 ↗

(On Diff #269927)

As above, this can be computed at compile time, e.g. from that expression

clang/include/clang/Basic/OpenMPGridValues.h

+//===--- OpenMPGridValues.h - Language-specific address spaces --*- C++ -*-===//
+//
+//                     The LLVM Compiler Infrastructure
+//
+// This file is distributed under the University of Illinois Open Source
+// License. See LICENSE.TXT for details.
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// \brief Provides definitions for Target specific Grid Values
+///
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_CLANG_BASIC_OPENMPGRIDVALUES_H
+#define LLVM_CLANG_BASIC_OPENMPGRIDVALUES_H
+
+namespace clang {
+
+namespace GPU {
+
+/// \brief Defines various target-specific Gpu grid values that must be
+///        consistent between host RTL (plugin), device RTL, and clang.
+///        By adding this to TargetInfo in clang, we can change grid values
+///        for a "fat" binary so that different passes get the correct values
+///        when generating code for a multi-target binary. Both amdgcn
+///        and nvptx values are stored in this file. In the future, should
+///        there be differences between GPUs of the same architecture,
+///        then simply make a different array and use the new array name.
+///
+/// Example usage in clang:
+///   const unsigned slot_size = ctx.GetTargetInfo().getGridValue(GV_Warp_Size);
+///
+/// Example usage in libomptarget/deviceRTLs:
+///   #include "OpenMPGridValues.h"
+///   #ifdef __AMDGPU__
+///     #define GRIDVAL AMDGPUGpuGridValues
+///   #else
+///     #define GRIDVAL NVPTXGpuGridValues
+///   #endif
+///   ... Then use this reference for GV_Warp_Size in the deviceRTL source.
+///   GRIDVAL[GV_Warp_Size]
+///
+/// Example usage in libomptarget hsa plugin:
+///   #include "OpenMPGridValues.h"
+///   #define GRIDVAL AMDGPUGpuGridValues
+///   ... Then use this reference to access GV_Warp_Size in the hsa plugin.
+///   GRIDVAL[GV_Warp_Size]
+///
+/// Example usage in libomptarget cuda plugin:
+///    #include "OpenMPGridValues.h"
+///    #define GRIDVAL NVPTXGpuGridValues
+///   ... Then use this reference to access GV_Warp_Size in the cuda plugin.
+///    GRIDVAL[GV_Warp_Size]
+///
+enum GVIDX {
+  /// The maximum number of workers in a kernel.
+  /// (THREAD_ABSOLUTE_LIMIT) - (GV_Warp_Size), might be issue for blockDim.z
+  GV_Threads,
+  /// The size reserved for data in a shared memory slot.
+  GV_Slot_Size,
+  /// The maximum number of threads in a worker warp.
+  GV_Warp_Size,
+  /// The number of bits required to represent the max number of threads in warp
+  GV_Warp_Size_Log2,
+  /// GV_Warp_Size * GV_Slot_Size,
+  GV_Warp_Slot_Size,
+  /// the maximum number of teams.
+  GV_Max_Teams,
+  /// Global Memory Alignment
+  GV_Mem_Align,
+  /// (~0u >> (GV_Warp_Size - GV_Warp_Size_Log2))
+  GV_Warp_Size_Log2_Mask,
+  // 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.
+  GV_SimpleBufferSize,
+  // The absolute maximum team size for a working group
+  GV_Max_WG_Size,
+  // The default maximum team size for a working group
+  GV_Default_WG_Size,
+  // This is GV_Max_WG_Size / GV_WarpSize. 32 for Nvidia and 16 for AMD.
+  GV_Max_Warp_Number
+};
+
+enum GVLIDX {
+  /// The slot size that should be reserved for a working warp.
+  /// (~0u >> (GV_Warp_Size - GV_Warp_Size_Log2))
+  GV_Warp_Size_Log2_MaskL
+};
+
+/// For AMDGPU GPUs
+static constexpr int AMDGPUGpuGridValues[] = {
+    448,      // GV_Threads   FIXME:  How can we make this bigger?
+    256,      // GV_Slot_Size
+    64,       // GV_Warp_Size
+    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 // This is GV_Max_WG_Size / GV_WarpSize
+
+};
+static constexpr long long AMDGPUGpuLongGridValues[] = {
+    63 // GV_Warp_Size_Log2_MaskL
+};
+
+/// For Nvidia GPUs
+static constexpr int NVPTXGpuGridValues[] = {
+    992,               // GV_Threads
+    256,               // GV_Slot_Size
+    32,                // GV_Warp_Size
+    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
+};
+
+static constexpr long long NVPTXGpuLongGridValues[] = {
+    31 // GV_Warp_Size_Log2_MaskL
+};
+
+} // namespace GPU
+} // namespace clang
+#endif

clang/include/clang/Basic/TargetInfo.h

 /// Defines the clang::TargetInfo interface.
 ///
 //===----------------------------------------------------------------------===//
 
 #ifndef LLVM_CLANG_BASIC_TARGETINFO_H
 #define LLVM_CLANG_BASIC_TARGETINFO_H
 
 #include "clang/Basic/AddressSpaces.h"
+#include "clang/Basic/CodeGenOptions.h"
 #include "clang/Basic/LLVM.h"
 #include "clang/Basic/LangOptions.h"
+#include "clang/Basic/OpenMPGridValues.h"
 #include "clang/Basic/Specifiers.h"
 #include "clang/Basic/TargetCXXABI.h"
 #include "clang/Basic/TargetOptions.h"
 #include "llvm/ADT/APFloat.h"
 #include "llvm/ADT/APInt.h"
 #include "llvm/ADT/ArrayRef.h"
 #include "llvm/ADT/IntrusiveRefCntPtr.h"
 #include "llvm/ADT/Optional.h"
protected:
 
   unsigned char MaxAtomicPromoteWidth, MaxAtomicInlineWidth;
   unsigned short SimdDefaultAlign;
   std::unique_ptr<llvm::DataLayout> DataLayout;
   const char *MCountName;
   unsigned char RegParmMax, SSERegParmMax;
   TargetCXXABI TheCXXABI;
   const LangASMap *AddrSpaceMap;
+  const int *GridValues;
+  const long long int *LongGridValues;
 
   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;
   }
 
+  int getGridValue(GPU::GVIDX gv) const { return GridValues[gv]; }
+
+  long long getLongGridValue(GPU::GVLIDX gv) const {
+    return LongGridValues[gv];
+  }
+
   /// 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.
   VersionTuple getPlatformMinVersion() const { return PlatformMinVersion; }
 

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/OpenMPGridValues.h"
 #include "clang/Basic/TargetBuiltins.h"
 #include "llvm/ADT/StringSwitch.h"
 #include "llvm/IR/DataLayout.h"
 
 using namespace clang;
 using namespace clang::targets;
 
 namespace clang {
: TargetInfo(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);
   assert(DataLayout->getAllocaAddrSpace() == Private);
+  GridValues = (const int *)&(GPU::AMDGPUGpuGridValues[0]);
+  LongGridValues = (const long long *)&(GPU::AMDGPUGpuLongGridValues[0]);
 
   setAddressSpaceMap(Triple.getOS() == llvm::Triple::Mesa3D ||
                      !isAMDGCN(Triple));
   UseAddrSpaceMapMangling = true;
 
   HasLegalHalfType = true;
   HasFloat16 = true;
 

clang/lib/Basic/Targets/NVPTX.cpp

 //===--- NVPTX.cpp - Implement NVPTX target feature support ---------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This file implements NVPTX TargetInfo objects.
 //
 //===----------------------------------------------------------------------===//
 
 #include "NVPTX.h"
 #include "Targets.h"
 #include "clang/Basic/Builtins.h"
 #include "clang/Basic/MacroBuilder.h"
+#include "clang/Basic/OpenMPGridValues.h"
 #include "clang/Basic/TargetBuiltins.h"
 #include "llvm/ADT/StringSwitch.h"
 
 using namespace clang;
 using namespace clang::targets;
 
 const Builtin::Info NVPTXTargetInfo::BuiltinInfo[] = {
 #define BUILTIN(ID, TYPE, ATTRS)                                               \
PTXVersion = llvm::StringSwitch<unsigned>(Feature)
                      .Case("+ptx40", 40)
                      .Case("+ptx32", 32)
                      .Default(32);
   }
 
   TLSSupported = false;
   VLASupported = false;
   AddrSpaceMap = &NVPTXAddrSpaceMap;
+  GridValues = (const int *)&(GPU::NVPTXGpuGridValues[0]);
+  LongGridValues = (const long long *)&(GPU::NVPTXGpuLongGridValues[0]);
   UseAddrSpaceMapMangling = true;
 
   // Define available target features
   // These must be defined in sorted order!
   NoAsmVariants = true;
   GPU = CudaArch::SM_20;
 
   if (TargetPointerWidth == 32)