[OpenMP][WIP] `spir64` device runtime library template

[spoutn1k]

openmp/libomptarget/DeviceRTL/src/State.cpp

69–86

This is only defined for AMDGCN with a TODO, in a way that seems to overwrite the default malloc. Should I do the same for this implementation ?

openmp/libomptarget/DeviceRTL/src/Utils.cpp

87–91

Found shuffle below in the spirv interface but not pack. Should I start using bitwise operation the way it is done for AMD ?

[tianshilei1992]

Can you update the patch with context?

[jdoerfert]

openmp/libomptarget/DeviceRTL/src/State.cpp

69–86

On AMDGPU there is no "default malloc". On NVPTX there is, on Intel we need to check. You can probably go with the AMD route for now.

openmp/libomptarget/DeviceRTL/src/Utils.cpp

87–91

There is a default implementation for Pack and Unpack, you can just not overwrite them. I think that is what we do for AMD.

[spoutn1k]

Added context

[spoutn1k]

Add pack/unpack and malloc/free for spir64 target runtime implementation.

[tianshilei1992]

llvm/lib/Frontend/OpenMP/OMPContext.cpp

55

SPIR-V can be for host as well right? OpenCL CPU runtime can also take SPIR-V.

openmp/libomptarget/DeviceRTL/src/LibC.cpp

37

I'm thinking potentially we can have a generic implementation of omp_vprintf that just returns -1 and having the NVPTX version calls vprintf. In this way we don't need to do it for every target.

[spoutn1k]

llvm/lib/Frontend/OpenMP/OMPContext.cpp

55

It can ! The objective of this work is to use it with intel gpus. I put it there in order to have it understand the variants in DeviceRTL. Is there a way to signify both host and device ?

openmp/libomptarget/DeviceRTL/src/LibC.cpp

37

How do I go about doing this ? Declare omp_vprintf as weak ?

[spoutn1k]

Added contents of DeviceRTL/src/Mapping.cpp; ignored activemask for the time being.

[tianshilei1992]

openmp/libomptarget/DeviceRTL/src/LibC.cpp

37

just something like:

int32_t omp_vprintf(const char *Format, void *Arguments, uint32_t) {
  return -1;
}
#pragma omp begin declare variant match(device = {arch(nvptx64)})
int32_t omp_vprintf(const char *Format, void *Arguments, uint32_t) {
  return vprintf(...);
}
#pragma omp end declare variant

In this way, when the arch is NVPTX64, that variant will replace the default one.

openmp/libomptarget/DeviceRTL/src/Mapping.cpp

162

Do we have these builtins in LLVM right now?

[spoutn1k]

openmp/libomptarget/DeviceRTL/src/LibC.cpp

37

Oh I see. I was under the impression this would create a symbol conflict. Will do !

openmp/libomptarget/DeviceRTL/src/Mapping.cpp

162

That was my question but Johannes told me to go ahead and use them.

[spoutn1k]

Defined spirv methods as extern and defined spirv enums.

clang/lib/Frontend/CompilerInvocation.cpp

for (unsigned i = 0; i < A->getNumValues(); ++i) {
       llvm::Triple TT(A->getValue(i));
 
       if (TT.getArch() == llvm::Triple::UnknownArch ||
           !(TT.getArch() == llvm::Triple::aarch64 || TT.isPPC() ||
             TT.getArch() == llvm::Triple::nvptx ||
             TT.getArch() == llvm::Triple::nvptx64 ||
             TT.getArch() == llvm::Triple::amdgcn ||
             TT.getArch() == llvm::Triple::x86 ||
-            TT.getArch() == llvm::Triple::x86_64))
+            TT.getArch() == llvm::Triple::x86_64 ||
+            TT.getArch() == llvm::Triple::spir64))
         Diags.Report(diag::err_drv_invalid_omp_target) << A->getValue(i);
       else if (getArchPtrSize(T) != getArchPtrSize(TT))
         Diags.Report(diag::err_drv_incompatible_omp_arch)
             << A->getValue(i) << T.str();
       else
         Opts.OMPTargetTriples.push_back(TT);
     }
   }

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

static constexpr GV NVPTXGridValues = {
     32,        // GV_Warp_Size
     (1 << 16), // GV_Max_Teams
     3200,      // GV_Default_Num_Teams
     896,       // GV_SimpleBufferSize
     1024,      // GV_Max_WG_Size
     128,       // GV_Default_WG_Size
 };
 
+/// For Intel GPUs
+static constexpr GV SPIR64GridValues64 = {
+    256,  // GV_Slot_Size
+    64,   // GV_Warp_Size
+    (1 << 16), // GV_Max_Teams
+    440,  // GV_Default_Num_Teams
+    896,  // GV_SimpleBufferSize
+    1024, // GV_Max_WG_Size,
+    256,  // GV_Default_WG_Size
+};
+
+
 } // namespace omp
 } // namespace llvm
 
 #endif // LLVM_FRONTEND_OPENMP_OMPGRIDVALUES_H

llvm/include/llvm/Frontend/OpenMP/OMPKinds.def

 __OMP_TRAIT_PROPERTY(device, arch, ppcle)
 __OMP_TRAIT_PROPERTY(device, arch, ppc64)
 __OMP_TRAIT_PROPERTY(device, arch, ppc64le)
 __OMP_TRAIT_PROPERTY(device, arch, x86)
 __OMP_TRAIT_PROPERTY(device, arch, x86_64)
 __OMP_TRAIT_PROPERTY(device, arch, amdgcn)
 __OMP_TRAIT_PROPERTY(device, arch, nvptx)
 __OMP_TRAIT_PROPERTY(device, arch, nvptx64)
+__OMP_TRAIT_PROPERTY(device, arch, spir64)
 
 __OMP_TRAIT_SET(implementation)
 
 __OMP_TRAIT_SELECTOR(implementation, vendor, true)
 
 __OMP_TRAIT_PROPERTY(implementation, vendor, amd)
 __OMP_TRAIT_PROPERTY(implementation, vendor, arm)
 __OMP_TRAIT_PROPERTY(implementation, vendor, bsc)

llvm/lib/Frontend/OpenMP/OMPContext.cpp

OMPContext::OMPContext(bool IsDeviceCompilation, Triple TargetTriple) {
   case Triple::ppc64le:
   case Triple::x86:
   case Triple::x86_64:
     ActiveTraits.set(unsigned(TraitProperty::device_kind_cpu));
     break;
   case Triple::amdgcn:
   case Triple::nvptx:
   case Triple::nvptx64:
+  case Triple::spir64:
     ActiveTraits.set(unsigned(TraitProperty::device_kind_gpu));
     break;
   default:
     break;
   }
 
   // Add the appropriate device architecture trait based on the triple.
 #define OMP_TRAIT_PROPERTY(Enum, TraitSetEnum, TraitSelectorEnum, Str)         \
   if (TraitSelector::TraitSelectorEnum == TraitSelector::device_arch) {        \

openmp/libomptarget/DeviceRTL/CMakeLists.txt

if("${gpu_arch}" IN_LIST all_amdgpu_architectures)
     compileDeviceRTLLibrary(${gpu_arch} amdgpu amdgcn-amd-amdhsa)
   elseif("${gpu_arch}" IN_LIST all_nvptx_architectures)
     compileDeviceRTLLibrary(${gpu_arch} nvptx nvptx64-nvidia-cuda --cuda-feature=+ptx61)
   else()
     libomptarget_error_say("Unknown GPU architecture '${gpu_arch}'")
   endif()
 endforeach()
 
+add_custom_target(omptarget.devicertl.spir64)
+# TODO Adapt names/versions here
+compileDeviceRTLLibrary(spir64 spir64 spir64-intel-l0 -fopenmp-targets=spir64)
+
 # Archive all the object files generated above into a static library
 add_library(omptarget.devicertl STATIC)
 set_target_properties(omptarget.devicertl PROPERTIES LINKER_LANGUAGE CXX)
 target_link_libraries(omptarget.devicertl PRIVATE omptarget.devicertl.all_objs)
 
 install(TARGETS omptarget.devicertl ARCHIVE DESTINATION ${OPENMP_INSTALL_LIBDIR})

openmp/libomptarget/DeviceRTL/include/Types.h

 using LaneMaskTy = uint64_t;
 #pragma omp end declare variant
 
 #pragma omp begin declare variant match(                                       \
     device = {arch(amdgcn)}, implementation = {extension(match_none)})
 using LaneMaskTy = uint64_t;
 #pragma omp end declare variant
 
+#pragma omp begin declare variant match(device = {arch(spir64)})
+// TODO verify this
+using LaneMaskTy = uint64_t;
+#pragma omp end declare variant
+
 namespace lanes {
 enum : LaneMaskTy { All = ~(LaneMaskTy)0 };
 } // namespace lanes
 
 /// The ident structure that describes a source location. The struct is
 /// identical to the one in the kmp.h file. We maintain the same data structure
 /// for compatibility.
 struct IdentTy {

openmp/libomptarget/DeviceRTL/src/LibC.cpp

 #pragma omp begin declare variant match(device = {arch(amdgcn)})
 namespace impl {
 int32_t omp_vprintf(const char *Format, void *Arguments, uint32_t) {
   return -1;
 }
 } // namespace impl
 #pragma omp end declare variant
 
+/// Intel implementation
+#pragma omp begin declare variant match(device = {arch(spir64)})
+namespace impl {
+int32_t omp_vprintf(const char *Format, void *Arguments, uint32_t) {
+  return -1;
+}
+} // namespace impl
+#pragma omp end declare variant
+
 extern "C" {
 
 int memcmp(const void *lhs, const void *rhs, size_t count) {
   auto *L = reinterpret_cast<const unsigned char *>(lhs);
   auto *R = reinterpret_cast<const unsigned char *>(rhs);
 
   for (size_t I = 0; I < count; ++I)
     if (L[I] != R[I])

openmp/libomptarget/DeviceRTL/src/Mapping.cpp

 uint32_t getNumberOfWarpsInBlock() {
   return (mapping::getBlockSize() + mapping::getWarpSize() - 1) /
          mapping::getWarpSize();
 }
 
 #pragma omp end declare variant
 ///}
 
+/// Intel implementation
+///
+///{
+#pragma omp begin declare variant match(device = {arch(spir64)})
+
+extern size_t __spirv_BuiltInGlobalSize(int);
+extern size_t __spirv_BuiltInLocalInvocationId(int);
+extern size_t __spirv_BuiltInWorkgroupId(int);
+extern size_t __spirv_BuiltInWorkgroupSize(int);
+
+const llvm::omp::GV &getGridValue() { return llvm::omp::SPIR64GridValues64; }
+
+uint32_t getNumHardwareThreadsInBlock() {
+  return __spirv_BuiltInWorkgroupSize(0);
+}
+
+LaneMaskTy activemask() { return -1; }
+
+LaneMaskTy lanemaskLT() {
+  uint32_t Lane = mapping::getThreadIdInWarp();
+  int64_t Ballot = mapping::activemask();
+  uint64_t Mask = ((uint64_t)1 << Lane) - (uint64_t)1;
+  return Mask & Ballot;
+}
+
+LaneMaskTy lanemaskGT() {
+  uint32_t Lane = mapping::getThreadIdInWarp();
+  if (Lane == (mapping::getWarpSize() - 1))
+    return 0;
+  int64_t Ballot = mapping::activemask();
+  uint64_t Mask = (~((uint64_t)0)) << (Lane + 1);
+  return Mask & Ballot;
+}
+
+uint32_t getThreadIdInWarp() {
+  return impl::getThreadIdInBlock() & (mapping::getWarpSize() - 1);
+}
+
+uint32_t getThreadIdInBlock() { return __spirv_BuiltInLocalInvocationId(0); }
+
+uint32_t getKernelSize() { return __spirv_BuiltInGlobalSize(0); }
+
+uint32_t getBlockId() { return __spirv_BuiltInWorkgroupId(0); }
+
+uint32_t getNumberOfBlocks() {
+  return __spirv_BuiltInGlobalSize(0) / __spirv_BuiltInWorkgroupSize(0);
+}
+
+uint32_t getWarpId() {
+  return impl::getThreadIdInBlock() / mapping::getWarpSize();
+}
+
+uint32_t getNumberOfWarpsInBlock() {
+  return mapping::getBlockSize() / mapping::getWarpSize();
+}
+
+#pragma omp end declare variant
+///}
+
 uint32_t getWarpSize() { return getGridValue().GV_Warp_Size; }
 
 } // namespace impl
 } // namespace ompx
 
 /// We have to be deliberate about the distinction of `mapping::` and `impl::`
 /// below to avoid repeating assumptions or including irrelevant ones.
 ///{

openmp/libomptarget/DeviceRTL/src/Misc.cpp

 double getWTime() {
   unsigned long long nsecs;
   asm volatile("mov.u64  %0, %%globaltimer;" : "=l"(nsecs));
   return (double)nsecs * getWTick();
 }
 
 #pragma omp end declare variant
 
+/// Intel implementation
+#pragma omp begin declare variant match(device = {arch(spir64)})
+
+double getWTick() { return ((double)1E-9); }
+
+double getWTime() {
+  // TODO
+  return 0;
+}
+
+#pragma omp end declare variant
+
 } // namespace impl
 } // namespace ompx
 
 /// Interfaces
 ///
 ///{
 
 extern "C" {

openmp/libomptarget/DeviceRTL/src/State.cpp

 }
 
 void free(void *Ptr) {}
 }
 
 #pragma omp end declare variant
 ///}
 
+/// Intel implementations of the shuffle sync idiom.
+///
+///{
+#pragma omp begin declare variant match(device = {arch(spir64)})
+
+extern "C" {
+void *malloc(uint64_t Size) {
+  // TODO: Use some preallocated space for dynamic malloc.
+  return nullptr;
+}
+
+void free(void *Ptr) {}
+}
+
+#pragma omp end declare variant
+///}
+
 /// A "smart" stack in shared memory.
 ///
 /// The stack exposes a malloc/free interface but works like a stack internally.
 /// In fact, it is a separate stack *per warp*. That means, each warp must push
 /// and pop symmetrically or this breaks, badly. The implementation will (aim
 /// to) detect non-lock-step warps and fallback to malloc/free. The same will
 /// happen if a warp runs out of memory. The master warp in generic memory is
 /// special and is given more memory than the rest.
 ///

openmp/libomptarget/DeviceRTL/src/Synchronization.cpp

for (;;) {
       }
     }
   } // wait for 0 to be the read value
 }
 
 #pragma omp end declare variant
 ///}
 
+/// Intel Implementation
+///
+///{
+#pragma omp begin declare variant match(device = {arch(spir64)})
+
+extern int __spirv_AtomicIAdd(int *, int, int, int);
+extern void __spirv_MemoryBarrier(int, int);
+extern void __spirv_ControlBarrier(int, int, int);
+
+typedef enum {
+  CrossDevice = 0,
+  Device = 1,
+  Workgroup = 2,
+  Subgroup = 3,
+  Invocation = 4
+} Scope_t;
+
+typedef enum {
+  Relaxed = 0x0,
+  Acquire = 0x2,
+  Release = 0x4,
+  AcquireRelease = 0x8,
+  SequentiallyConsistent = 0x10,
+  UniformMemory = 0x40,
+  SubgroupMemory = 0x80,
+  WorkgroupMemory = 0x100,
+  CrossWorkgroupMemory = 0x200,
+  AtomicCounterMemory = 0x400,
+  ImageMemory = 0x800
+} MemorySemantics_t;
+
+uint32_t atomicInc(uint32_t *uA, uint32_t V, atomic::OrderingTy Ordering) {
+  int *A = (int *)A;
+  switch (Ordering) {
+  default:
+    __builtin_unreachable();
+  case atomic::relaxed:
+    return __spirv_AtomicIAdd(A, Scope_t::Device, MemorySemantics_t::Relaxed,
+                              V);
+  case atomic::aquire:
+    return __spirv_AtomicIAdd(A, Scope_t::Device, MemorySemantics_t::Acquire,
+                              V);
+  case atomic::release:
+    return __spirv_AtomicIAdd(A, Scope_t::Device, MemorySemantics_t::Release,
+                              V);
+  case atomic::acq_rel:
+    return __spirv_AtomicIAdd(A, Scope_t::Device,
+                              MemorySemantics_t::AcquireRelease, V);
+  case atomic::seq_cst:
+    return __spirv_AtomicIAdd(A, Scope_t::Device,
+                              MemorySemantics_t::SequentiallyConsistent, V);
+  }
+}
+
+void namedBarrierInit() { __builtin_trap(); } // TODO
+
+void namedBarrier() { __builtin_trap(); } // TODO
+
+void fenceTeam(atomic::OrderingTy Ordering) {
+  switch (Ordering) {
+  default:
+    __builtin_unreachable();
+  case atomic::aquire:
+    return __spirv_MemoryBarrier(Scope_t::Workgroup,
+                                 MemorySemantics_t::Acquire);
+  case atomic::release:
+    return __spirv_MemoryBarrier(Scope_t::Workgroup,
+                                 MemorySemantics_t::Release);
+  case atomic::acq_rel:
+    return __spirv_MemoryBarrier(Scope_t::Workgroup,
+                                 MemorySemantics_t::AcquireRelease);
+  case atomic::seq_cst:
+    return __spirv_MemoryBarrier(Scope_t::Workgroup,
+                                 MemorySemantics_t::SequentiallyConsistent);
+  }
+}
+void fenceKernel(atomic::OrderingTy Ordering) {
+  switch (Ordering) {
+  default:
+    __builtin_unreachable();
+  case atomic::aquire:
+    return __spirv_MemoryBarrier(Scope_t::Invocation,
+                                 MemorySemantics_t::Acquire);
+  case atomic::release:
+    return __spirv_MemoryBarrier(Scope_t::Invocation,
+                                 MemorySemantics_t::Release);
+  case atomic::acq_rel:
+    return __spirv_MemoryBarrier(Scope_t::Invocation,
+                                 MemorySemantics_t::AcquireRelease);
+  case atomic::seq_cst:
+    return __spirv_MemoryBarrier(Scope_t::Invocation,
+                                 MemorySemantics_t::SequentiallyConsistent);
+  }
+}
+void fenceSystem(atomic::OrderingTy Ordering) {
+  switch (Ordering) {
+  default:
+    __builtin_unreachable();
+  case atomic::aquire:
+    return __spirv_MemoryBarrier(Scope_t::Device, MemorySemantics_t::Acquire);
+  case atomic::release:
+    return __spirv_MemoryBarrier(Scope_t::Device, MemorySemantics_t::Release);
+  case atomic::acq_rel:
+    return __spirv_MemoryBarrier(Scope_t::Device,
+                                 MemorySemantics_t::AcquireRelease);
+  case atomic::seq_cst:
+    return __spirv_MemoryBarrier(Scope_t::Device,
+                                 MemorySemantics_t::SequentiallyConsistent);
+  }
+}
+
+void syncWarp(__kmpc_impl_lanemask_t) {
+  __spirv_ControlBarrier(Scope_t::Subgroup, Scope_t::Invocation,
+                         MemorySemantics_t::Acquire);
+}
+
+void syncThreads() {
+  __spirv_ControlBarrier(Scope_t::Workgroup, Scope_t::Invocation,
+                         MemorySemantics_t::Acquire);
+}
+
+void syncThreadsAligned() { syncThreads(); }
+
+void unsetLock(omp_lock_t *) { __builtin_trap(); }
+int testLock(omp_lock_t *) { __builtin_trap(); }
+void initLock(omp_lock_t *) { __builtin_trap(); }
+void destroyLock(omp_lock_t *) { __builtin_trap(); }
+void setLock(omp_lock_t *) { __builtin_trap(); }
+
+#pragma omp end declare variant
+///}
+
 } // namespace impl
 
 void synchronize::init(bool IsSPMD) {
   if (!IsSPMD)
     impl::namedBarrierInit();
 }
 
 void synchronize::warp(LaneMaskTy Mask) { impl::syncWarp(Mask); }

openmp/libomptarget/DeviceRTL/src/Utils.cpp

uint64_t Pack(uint32_t LowBits, uint32_t HighBits) {
   uint64_t Val;
   asm("mov.b64 %0, {%1,%2};" : "=l"(Val) : "r"(LowBits), "r"(HighBits));
   return Val;
 }
 
 #pragma omp end declare variant
 ///}
 
+/// Intel Implementation
+///
+///{
+#pragma omp begin declare variant match(device = {arch(spir64)})
+
+void Unpack(uint64_t Val, uint32_t *LowBits, uint32_t *HighBits) {
+  static_assert(sizeof(unsigned long) == 8, "");
+  *LowBits = (uint32_t)(Val & 0x00000000FFFFFFFFUL);
+  *HighBits = (uint32_t)((Val & 0xFFFFFFFF00000000UL) >> 32);
+}
+
+uint64_t Pack(uint32_t LowBits, uint32_t HighBits) {
+  return (((uint64_t)HighBits) << 32) | (uint64_t)LowBits;
+}
+
+#pragma omp end declare variant
+///}
+
 int32_t shuffle(uint64_t Mask, int32_t Var, int32_t SrcLane);
 int32_t shuffleDown(uint64_t Mask, int32_t Var, uint32_t LaneDelta,
                     int32_t Width);
 
 /// AMDGCN Implementation
 ///
 ///{
 #pragma omp begin declare variant match(device = {arch(amdgcn)})
 
 int32_t shuffle(uint64_t Mask, int32_t Var, int32_t SrcLane) {
   int Width = mapping::getWarpSize();
   int Self = mapping::getThreadIdInWarp();
   int Index = SrcLane + (Self & ~(Width - 1));
   return __builtin_amdgcn_ds_bpermute(Index << 2, Var);
 }
 
 int32_t shuffleDown(uint64_t Mask, int32_t Var, uint32_t LaneDelta,
                     int32_t Width) {
   int Self = mapping::getThreadIdInWarp();
   int Index = Self + LaneDelta;
   Index = (int)(LaneDelta + (Self & (Width - 1))) >= Width ? Self : Index;
   return __builtin_amdgcn_ds_bpermute(Index << 2, Var);
 }
 
-bool isSharedMemPtr(const void * Ptr) {
-  return __builtin_amdgcn_is_shared((const __attribute__((address_space(0))) void *)Ptr);
+bool isSharedMemPtr(const void *Ptr) {
+  return __builtin_amdgcn_is_shared(
+      (const __attribute__((address_space(0))) void *)Ptr);
 }
 #pragma omp end declare variant
 ///}
 
 /// NVPTX Implementation
 ///
 ///{
 #pragma omp begin declare variant match(                                       \
     device = {arch(nvptx, nvptx64)}, implementation = {extension(match_any)})
 
 int32_t shuffle(uint64_t Mask, int32_t Var, int32_t SrcLane) {
   return __nvvm_shfl_sync_idx_i32(Mask, Var, SrcLane, 0x1f);
 }
 
 int32_t shuffleDown(uint64_t Mask, int32_t Var, uint32_t Delta, int32_t Width) {
   int32_t T = ((mapping::getWarpSize() - Width) << 8) | 0x1f;
   return __nvvm_shfl_sync_down_i32(Mask, Var, Delta, T);
 }
 
 bool isSharedMemPtr(const void *Ptr) { return __nvvm_isspacep_shared(Ptr); }
 
 #pragma omp end declare variant
 ///}
+
+/// Intel Implementation
+///
+///{
+#pragma omp begin declare variant match(device = {arch(spir64)})
+
+int32_t shuffle(uint64_t Mask, int32_t Var, int32_t SrcLane) { // TODO
+  __builtin_trap();
+}
+
+int32_t shuffleDown(uint64_t Mask, int32_t Var, uint32_t Delta,
+                    int32_t Width) { // TODO
+  __builtin_trap();
+}
+
+bool isSharedMemPtr(const void *Ptr) { // TODO
+  __builtin_trap();
+}
+
+#pragma omp end declare variant
+///}
 } // namespace impl
 
 uint64_t utils::pack(uint32_t LowBits, uint32_t HighBits) {
   return impl::Pack(LowBits, HighBits);
 }
 
 void utils::unpack(uint64_t Val, uint32_t &LowBits, uint32_t &HighBits) {
   impl::Unpack(Val, &LowBits, &HighBits);