[Openmp-commits] [PATCH] D112227: [libomptarget] Build DeviceRTL for amdgpu

[Jon Chesterfield via Phabricator via Openmp-commits]

JonChesterfield added inline comments.


================
Comment at: openmp/libomptarget/DeviceRTL/src/Configuration.cpp:23
 
-extern uint32_t __omp_rtl_debug_kind;
+// extern uint32_t __omp_rtl_debug_kind;
 
----------------
Otherwise the missing symbols prevents linking, not clear why it works on nvptx64


================
Comment at: openmp/libomptarget/DeviceRTL/src/Synchronization.cpp:71
 
-uint32_t atomicInc(uint32_t *Address, uint32_t Val, int Ordering) {
-  return __builtin_amdgcn_atomic_inc32(Address, Val, Ordering, "");
+uint32_t atomicInc(uint32_t *A, uint32_t V, int Ordering) {
+  // builtin_amdgcn_atomic_inc32 should expand to this switch when
----------------
This is not good, need to revise sema checking on these intrinsics and add some lowering in clang/llvm that builds the switch. Written longhand here to get things running.


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D112227/new/

https://reviews.llvm.org/D112227