[libomptarget][WIP] Patch amdgpu DeviceRTL until it compiles

[JonChesterfield]

Summary

Requires cmake changes in D111983, D111987

Mixture of typo fixes, stub implementations and a workaround for
runtime values for memory ordering

[JonChesterfield]

clang/lib/Driver/ToolChains/AMDGPUOpenMP.cpp

255

Naming is somewhat inconsistent. I think 'amdgcn' refers to 'graphics core next', but the gfx10 cards are branded 'radeon dna'. Changing to a new library seems a relatively good point to rename it.

openmp/libomptarget/DeviceRTL/include/Synchronization.h

48

^ would like to call this 'load'

openmp/libomptarget/DeviceRTL/src/Configuration.cpp

23

What is this? Symbol is undefined on nvptx as far as I can tell. Amdgpu fails to link without a definition.

[JonChesterfield]

openmp/libomptarget/DeviceRTL/src/Synchronization.cpp

87

Suggestion offline is that delaying the check in clang for these builtins until template instantiation would make them usable from a function template, worthwhile usability improvement

[JonChesterfield]

Needs manual rebase on trunk + clang patch to delay checking the arguments to amdgcn intrinsics until template instantiation time

[JonChesterfield]

openmp/libomptarget/DeviceRTL/src/Synchronization.cpp

92

The assert in clang was unrelated to amdgpu, posted a diff at D112159 with a fix.

The other atomic builtins codegen with a switch on order when passed a variable. That seems like the right thing to do for the amdgpu-specific atomic builtins.

[JonChesterfield]

Going to discard this and recreate as a new patch which includes the two cmake ones (D111983 and D111987)

clang/lib/Driver/ToolChains/AMDGPUOpenMP.cpp

void AMDGPUOpenMPToolChain::addClangTargetOptions(
   CC1Args.push_back("-fcuda-is-device");
 
   if (DriverArgs.hasArg(options::OPT_nogpulib))
     return;
 
   std::string BitcodeSuffix;
   if (DriverArgs.hasFlag(options::OPT_fopenmp_target_new_runtime,
                          options::OPT_fno_openmp_target_new_runtime, false))
-    BitcodeSuffix = "new-amdgcn-" + GPUArch;
+    BitcodeSuffix = "new-amdgpu-" + GPUArch;
   else
     BitcodeSuffix = "amdgcn-" + GPUArch;
 
   addOpenMPDeviceRTL(getDriver(), DriverArgs, CC1Args, BitcodeSuffix,
                      getTriple());
 }
 
 llvm::opt::DerivedArgList *AMDGPUOpenMPToolChain::TranslateArgs(

openmp/libomptarget/DeviceRTL/include/Synchronization.h

 /// Memory fence with \p Ordering semantics for the system.
 void system(int Ordering);
 
 } // namespace fence
 
 namespace atomic {
 
 /// Atomically read \p Addr with \p Ordering semantics.
 uint32_t read(uint32_t *Addr, int Ordering);
 
 /// Atomically store \p V to \p Addr with \p Ordering semantics.
-uint32_t store(uint32_t *Addr, uint32_t V, int Ordering);
+void store(uint32_t *Addr, uint32_t V, int Ordering);
 
 /// Atomically store \p V to \p Addr with \p Ordering semantics.
-uint64_t store(uint64_t *Addr, uint64_t V, int Ordering);
+void store(uint64_t *Addr, uint64_t V, int Ordering);
 
 /// Atomically increment \p *Addr and wrap at \p V with \p Ordering semantics.
 uint32_t inc(uint32_t *Addr, uint32_t V, int Ordering);
 
 /// Atomically add \p V to \p *Addr with \p Ordering semantics.
 uint32_t add(uint32_t *Addr, uint32_t V, int Ordering);
 
 /// Atomically add \p V to \p *Addr with \p Ordering semantics.
 uint64_t add(uint64_t *Addr, uint64_t V, int Ordering);
 
 } // namespace atomic
 
 } // namespace _OMP
 
 #endif

openmp/libomptarget/DeviceRTL/src/Configuration.cpp

 #include "DeviceEnvironment.h"
 #include "State.h"
 #include "Types.h"
 
 using namespace _OMP;
 
 #pragma omp declare target
 
-extern uint32_t __omp_rtl_debug_kind;
+// extern uint32_t __omp_rtl_debug_kind;
 
 // TOOD: We want to change the name as soon as the old runtime is gone.
 DeviceEnvironmentTy CONSTANT(omptarget_device_environment)
     __attribute__((used));
 
 uint32_t config::getDebugKind() {
-  return __omp_rtl_debug_kind & omptarget_device_environment.DebugKind;
+  return /*__omp_rtl_debug_kind &*/ omptarget_device_environment.DebugKind;
 }
 
 uint32_t config::getNumDevices() {
   return omptarget_device_environment.NumDevices;
 }
 
 uint32_t config::getDeviceNum() {
   return omptarget_device_environment.DeviceNum;

openmp/libomptarget/DeviceRTL/src/Mapping.cpp

 namespace impl {
 
 /// AMDGCN Implementation
 ///
 ///{
 #pragma omp begin declare variant match(device = {arch(amdgcn)})
 
 constexpr const llvm::omp::GV &getGridValue() {
-  return llvm::omp::AMDGPUGridValues;
+  return llvm::omp::getAMDGPUGridValues<__AMDGCN_WAVEFRONT_SIZE>();
 }
 
 uint32_t getGridDim(uint32_t n, uint16_t d) {
   uint32_t q = n / d;
   return q + (n > q * d);
 }
 
 uint32_t getWorkgroupDim(uint32_t group_id, uint32_t grid_size,

openmp/libomptarget/DeviceRTL/src/Synchronization.cpp

 
 /// Atomics
 ///
 ///{
 /// NOTE: This function needs to be implemented by every target.
 uint32_t atomicInc(uint32_t *Address, uint32_t Val, int Ordering);
 
 uint32_t atomicRead(uint32_t *Address, int Ordering) {
+  // TODO: amdgpu can load from memory. nvptx probably can too
   return __atomic_fetch_add(Address, 0U, __ATOMIC_SEQ_CST);
 }
 
+void atomicStore(uint32_t *Address, uint32_t Val, int Ordering) {
+  __atomic_store_n(Address, Val, Ordering);
+}
+
+void atomicStore(uint64_t *Address, uint64_t Val, int Ordering) {
+  __atomic_store_n(Address, Val, Ordering);
+}
+
 uint32_t atomicAdd(uint32_t *Address, uint32_t Val, int Ordering) {
   return __atomic_fetch_add(Address, Val, Ordering);
 }
 uint32_t atomicMax(uint32_t *Address, uint32_t Val, int Ordering) {
   return __atomic_fetch_max(Address, Val, Ordering);
 }
 
 uint32_t atomicExchange(uint32_t *Address, uint32_t Val, int Ordering) {

 }
 ///}
 
 /// AMDGCN Implementation
 ///
 ///{
 #pragma omp begin declare variant match(device = {arch(amdgcn)})
 
-uint32_t atomicInc(uint32_t *Address, uint32_t Val, int Ordering) {
-  return __builtin_amdgcn_atomic_inc32(Address, Val, Ordering, "");
+#define FENCE_ATOMIC_CASES()                                                   \
+  CASE(__ATOMIC_ACQUIRE);                                                      \
+  CASE(__ATOMIC_RELEASE);                                                      \
+  CASE(__ATOMIC_ACQ_REL);                                                      \
+  CASE(__ATOMIC_SEQ_CST)
+#define ALL_ATOMIC_CASES()                                                     \
+  FENCE_ATOMIC_CASES();                                                        \
+  CASE(__ATOMIC_RELAXED);                                                      \
+  CASE(__ATOMIC_CONSUME)
+
+#if 0
+// Can't spell the dispatch from runtime ordering like:
+template <int Ordering>
+static uint32_t atomicInc(uint32_t *Address, uint32_t Val) {
+  return __builtin_amdgcn_atomic_inc32(Address, Val, Ordering , "");
 }
+// clang/lib/AST/ExprConstant.cpp:14739: bool clang::Expr::EvaluateAsInt(clang::Expr::EvalResult&, const clang::ASTContext&, clang::Expr::SideEffectsKind, bool) const: Assertion `!isValueDependent() && "Expression evaluator can't be called on a dependent expression."' failed.
 
-uint32_t SHARD(namedBarrierTracker);
+// Can spell it with raw macros or struct dispatch, or raise ordering to a
+// template parameter on the Synchonization API. Choosing struct dispatch.
+#endif
+
+
+namespace {
+template <int Ordering> struct atomicOpTy;
+#define CASE(X)                                                                \
+  template <> struct atomicOpTy<X> {                                           \
+    static uint32_t atomicInc(uint32_t *Address, uint32_t Val) {               \
+      return __builtin_amdgcn_atomic_inc32(Address, Val, X, "");               \
+    }                                                                          \
+  };
+ALL_ATOMIC_CASES();
+#undef CASE
+}
+
+uint32_t atomicInc(uint32_t *Address, uint32_t Val, int Ordering) {
+  switch (Ordering) {
+  default:
+    __builtin_unreachable();
+#define CASE(X)                                                                \
+  case X:                                                                      \
+    return atomicOpTy<X>::atomicInc(Address, Val)
+    FENCE_ATOMIC_CASES();
+#undef CASE
+  }
+}
+
+uint32_t SHARED(namedBarrierTracker);
 
 void namedBarrierInit() {
   // Don't have global ctors, and shared memory is not zero init
   atomic::store(&namedBarrierTracker, 0u, __ATOMIC_RELEASE);
 }
 
 void namedBarrier() {
   uint32_t NumThreads = omp_get_num_threads();
   // assert(NumThreads % 32 == 0);
 
-  uint32_t WarpSize = maping::getWarpSize();
+  uint32_t WarpSize = mapping::getWarpSize();
   uint32_t NumWaves = NumThreads / WarpSize;
 
   fence::team(__ATOMIC_ACQUIRE);
 
   // named barrier implementation for amdgcn.
   // Uses two 16 bit unsigned counters. One for the number of waves to have
   // reached the barrier, and one to count how many times the barrier has been
   // passed. These are packed in a single atomically accessed 32 bit integer.
if ((load & 0x0000ffffu) == (NumWaves - 1)) {
       load &= 0xffff0000u; // because bits zeroed second
 
       // Reset the wave counter and release the waiting waves
       atomic::store(&namedBarrierTracker, load, __ATOMIC_RELAXED);
     } else {
       // more waves still to go, spin until generation counter changes
       do {
         __builtin_amdgcn_s_sleep(0);
-        load = atomi::load(&namedBarrierTracker, __ATOMIC_RELAXED);
+        load = atomic::read(&namedBarrierTracker, __ATOMIC_RELAXED);
       } while ((load & 0xffff0000u) == generation);
     }
   }
   fence::team(__ATOMIC_RELEASE);
 }
 
+namespace {
+template <int Ordering> struct atomicFenceTy;
+#define CASE(X)                                                                \
+  template <> struct atomicFenceTy<X> {                                        \
+    static void fenceTeam() { __builtin_amdgcn_fence(X, "workgroup"); }        \
+    static void fenceKernel() { __builtin_amdgcn_fence(X, "agent"); }          \
+    static void fenceSystem() { __builtin_amdgcn_fence(X, ""); }               \
+  };
+FENCE_ATOMIC_CASES();
+#undef CASE
+}
+
+void fenceTeam(int Ordering) {
+  switch (Ordering) {
+  default:
+    __builtin_unreachable();
+#define CASE(X)                                                                \
+  case X:                                                                      \
+    atomicFenceTy<X>::fenceTeam()
+    FENCE_ATOMIC_CASES();
+#undef CASE
+  }
+}
+
+void fenceKernel(int Ordering) {
+  switch (Ordering) {
+  default:
+    __builtin_unreachable();
+#define CASE(X)                                                                \
+  case X:                                                                      \
+    atomicFenceTy<X>::fenceKernel()
+    FENCE_ATOMIC_CASES();
+#undef CASE
+  }
+}
+
+void fenceSystem(int Ordering) {
+  switch (Ordering) {
+  default:
+    __builtin_unreachable();
+#define CASE(X)                                                                \
+  case X:                                                                      \
+    atomicFenceTy<X>::fenceSystem()
+    FENCE_ATOMIC_CASES();
+#undef CASE
+  }
+}
+
 void syncWarp(__kmpc_impl_lanemask_t) {
   // AMDGCN doesn't need to sync threads in a warp
 }
 
 void syncThreads() { __builtin_amdgcn_s_barrier(); }
 
-void fenceTeam(int Ordering) { __builtin_amdgcn_fence(Ordering, "workgroup"); }
-
-void fenceKernel(int Ordering) { __builtin_amdgcn_fence(Ordering, "agent"); }
-
-void fenceSystem(int Ordering) { __builtin_amdgcn_fence(Ordering, ""); }
+// TODO: Don't have wavefront lane locks. Possibly can't have them.
+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
 ///}
 
 /// NVPTX Implementation
 ///
 ///{
 #pragma omp begin declare variant match(                                       \

 }
 
 int testLock(omp_lock_t *Lock) {
   return atomicAdd((uint32_t *)Lock, 0u, __ATOMIC_SEQ_CST);
 }
 
 void initLock(omp_lock_t *Lock) { unsetLock(Lock); }
 
-void destoryLock(omp_lock_t *Lock) { unsetLock(Lock); }
+void destroyLock(omp_lock_t *Lock) { unsetLock(Lock); }
 
 void setLock(omp_lock_t *Lock) {
   // TODO: not sure spinning is a good idea here..
   while (atomicCAS((uint32_t *)Lock, UNSET, SET, __ATOMIC_SEQ_CST) != UNSET) {
     int32_t start = __nvvm_read_ptx_sreg_clock();
     int32_t now;
     for (;;) {
       now = __nvvm_read_ptx_sreg_clock();

 void fence::kernel(int Ordering) { impl::fenceKernel(Ordering); }
 
 void fence::system(int Ordering) { impl::fenceSystem(Ordering); }
 
 uint32_t atomic::read(uint32_t *Addr, int Ordering) {
   return impl::atomicRead(Addr, Ordering);
 }
 
+void atomic::store(uint32_t *Addr, uint32_t V, int Ordering) {
+   impl::atomicStore(Addr, V, Ordering);
+}
+
+void atomic::store(uint64_t *Addr, uint64_t V, int Ordering) {
+  impl::atomicStore(Addr, V, Ordering);
+}
+
 uint32_t atomic::inc(uint32_t *Addr, uint32_t V, int Ordering) {
   return impl::atomicInc(Addr, V, Ordering);
 }
 
 uint32_t atomic::add(uint32_t *Addr, uint32_t V, int Ordering) {
   return impl::atomicAdd(Addr, V, Ordering);
 }
 

 }
 
 void __kmpc_end_critical(IdentTy *Loc, int32_t TId, CriticalNameTy *Name) {
   omp_unset_lock(reinterpret_cast<omp_lock_t *>(Name));
 }
 
 void omp_init_lock(omp_lock_t *Lock) { impl::initLock(Lock); }
 
-void omp_destroy_lock(omp_lock_t *Lock) { impl::destoryLock(Lock); }
+void omp_destroy_lock(omp_lock_t *Lock) { impl::destroyLock(Lock); }
 
 void omp_set_lock(omp_lock_t *Lock) { impl::setLock(Lock); }
 
 void omp_unset_lock(omp_lock_t *Lock) { impl::unsetLock(Lock); }
 
 int omp_test_lock(omp_lock_t *Lock) { return impl::testLock(Lock); }
 } // extern "C"
 
 #pragma omp end declare target

openmp/libomptarget/DeviceRTL/src/Utils.cpp

 namespace impl {
 
 /// AMDGCN Implementation
 ///
 ///{
 #pragma omp begin declare variant match(device = {arch(amdgcn)})
 
 void Unpack(uint64_t Val, uint32_t *LowBits, uint32_t *HighBits) {
-  *LowBits = (uint32_t)(Val & UINT64_C(0x00000000FFFFFFFF));
-  *HighBits = (uint32_t)((Val & UINT64_C(0xFFFFFFFF00000000)) >> 32);
+  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
 

 
 /// 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::getgetThreadIdInWarp();
+  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;