[libomptarget] Build DeviceRTL for amdgpu

[JonChesterfield]

Summary

Passes same tests as the current deviceRTL. Includes cmake change from D111987.

[JonChesterfield]

openmp/libomptarget/DeviceRTL/src/Configuration.cpp

23

Otherwise the missing symbols prevents linking, not clear why it works on nvptx64

openmp/libomptarget/DeviceRTL/src/Synchronization.cpp

71

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.

[jdoerfert]

openmp/libomptarget/DeviceRTL/src/Configuration.cpp

23

linking what? Clang emits the symbol, maybe just not for amdgpu.

[JonChesterfield]

openmp/libomptarget/DeviceRTL/src/Configuration.cpp

23

Where? The only reference I can find to it is here, and it's marked extern.

[JonChesterfield]

Subscribed some AMD people to this. I wanted to apply this patch as-is to amd-stg-open to feed it to the internal testing, but it doesn't apply because Driver/ToolChains/AMDGPUOpenMP.cpp in rocm is significantly different to trunk (in particular the call to addOpenMPDeviceRTL is commented out)

[JonChesterfield]

• rebase

[JonChesterfield]

• Enable tests on amdgpu, with same ones marked xfail/unsupported as on the old runtime

[JonChesterfield]

• drop outdated comment

[JonChesterfield]

• delete the extern debug_kind variable

[JonChesterfield]

I think this is good enough for now. It drops the not yet used debug variable and writes out the lowering for runtime values of memory ordering manually. The latter will be simplified once clang learns to emit the switch instead of error. Omp lock is a problem I don't have a good solution to.

clang/lib/Driver/ToolChains/AMDGPUOpenMP.cpp

255

'amdgcn' appears to be a subset of 'amdgpu', so this seems a reasonable point to rename it.

openmp/libomptarget/DeviceRTL/CMakeLists.txt

182

rearranging the naming here - the llvm-link file is now prefixed linked_, with the optimised library left without a prefix. Updated depends / output clauses to match.

openmp/libomptarget/DeviceRTL/src/Configuration.cpp

23

I think the cuda toolchain treats unresolved references as 'just use zero', in which case deleting this is a no-op on nvptx. Maybe it's intended to be patched by cuda rtl.cpp in the future? If so can reintroduce it then

[JonChesterfield]

• rebase

[JonChesterfield]

• rebase

[JonChesterfield]

• rebase, having landed D111987 this time

[JonChesterfield]

Problem with missing symbol for __omp_rtl_debug_kind was a local error. I did the initial testing of this with a jury rigged clang that linked the new bitcode and ignored the old. The generation of this integer is guarded by which runtime clang thinks it is compiling for. Thus, my local clang compiled for the old runtime and linked with the new, which it turns out does not work.

I was under the impression that pointing clang at the new runtime with libomptarget-amdgcn-bc-path was sufficient, but evidently not. Doesn't matter for this patch now that the variable is reinstated.

[JonChesterfield]

Requires either D112544 or disabling bug51982 on newRTL before landing

[jdoerfert]

LG, I'll land my fix and with this we can switch over to the new RT.

[JonChesterfield]

Didn't fare very well under CI, investigating. Ten failures at https://lab.llvm.org/buildbot/#/builders/193/builds/915, but they all pass locally.

[JonChesterfield]

openmp/libomptarget/DeviceRTL/src/Synchronization.cpp

196–201

Error here - syncThreadsAligned is deleted but should not be

[JonChesterfield]

• Reintroduce syncThreadsAligned, dropped in git merge conflict

[JonChesterfield]

• fix new+old test enabling

[JonChesterfield]

Tagging Ron as this is current stuck on the mystery of passing locally and failing CI

[JonChesterfield]

Landed a slightly modified version of this - code and test changes are included, but the tests are not run by default. I'm hopeful this will help the process of working out why ~10 are failing on CI and passing locally.

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/CMakeLists.txt

 
 # Check all SM values
 foreach(sm ${nvptx_sm_list})
   if (NOT ${sm} IN_LIST all_capabilities)
     libomptarget_warning_say("[NVPTX] Compute capability ${sm} is not supported. Make sure clang can work with it.")
   endif()
 endforeach()
 
+set(amdgpu_mcpus gfx700 gfx701 gfx801 gfx803 gfx900 gfx902 gfx906 gfx908 gfx90a gfx1010 gfx1030 gfx1031)
+if (DEFINED LIBOMPTARGET_AMDGCN_GFXLIST)
+  set(amdgpu_mcpus ${LIBOMPTARGET_AMDGCN_GFXLIST})
+endif()
+
+
 # Activate RTL message dumps if requested by the user.
 set(LIBOMPTARGET_DEVICE_DEBUG FALSE CACHE BOOL
   "Activate DeviceRTL debug messages.")
 
 
 set(src_files
   ${source_directory}/Configuration.cpp
   ${source_directory}/Debug.cpp

 
 # Prepend -I to each list element
 set (LIBOMPTARGET_LLVM_INCLUDE_DIRS_DEVICERTL "${LIBOMPTARGET_LLVM_INCLUDE_DIRS}")
 list(TRANSFORM LIBOMPTARGET_LLVM_INCLUDE_DIRS_DEVICERTL PREPEND "-I")
 
 # Set flags for LLVM Bitcode compilation.
 set(bc_flags -S -x c++ -std=c++17
               ${clang_opt_flags}
-             -target nvptx64
              -Xclang -emit-llvm-bc
              -Xclang -aux-triple -Xclang ${aux_triple}
              -fopenmp -fopenmp-cuda-mode -Xclang -fopenmp-is-device
-             -Xclang -target-feature -Xclang +ptx61
              -I${include_directory}
              -I${devicertl_base_directory}/../include
              ${LIBOMPTARGET_LLVM_INCLUDE_DIRS_DEVICERTL}
 )
 
 if(${LIBOMPTARGET_DEVICE_DEBUG})
   list(APPEND bc_flags -DOMPTARGET_DEBUG=-1)
 else()
   list(APPEND bc_flags -DOMPTARGET_DEBUG=0)
 endif()
 
-# Create target to build all Bitcode libraries.
-add_custom_target(omptarget-new-nvptx-bc)
-add_dependencies(omptarget-new-nvptx-bc opt llvm-link)
 
-# Generate a Bitcode library for all the compute capabilities the user requested
-foreach(sm ${nvptx_sm_list})
-  # TODO: replace this with declare variant and isa selector.
-  set(cuda_flags -Xclang -target-cpu -Xclang sm_${sm} "-D__CUDA_ARCH__=${sm}0")
+macro(instantiate_DeviceRTL)
+  # parameters target_cpu, target_name, target_bc_flags
+
   set(bc_files "")
   foreach(src ${src_files})
     get_filename_component(infile ${src} ABSOLUTE)
     get_filename_component(outfile ${src} NAME)
-    set(outfile "${outfile}-sm_${sm}.bc")
+    set(outfile "${outfile}-${target_cpu}.bc")
 
     add_custom_command(OUTPUT ${outfile}
-      COMMAND ${CLANG_TOOL} ${bc_flags}
-        ${cuda_flags} ${infile} -o ${outfile}
+      COMMAND ${CLANG_TOOL}
+      ${bc_flags}
+      -Xclang -target-cpu -Xclang ${target_cpu}
+      ${target_bc_flags}
+      ${infile} -o ${outfile}
       DEPENDS ${infile}
       IMPLICIT_DEPENDS CXX ${infile}
       COMMENT "Building LLVM bitcode ${outfile}"
       VERBATIM
     )
     if("${CLANG_TOOL}" STREQUAL "$<TARGET_FILE:clang>")
       # Add a file-level dependency to ensure that clang is up-to-date.
       # By default, add_custom_command only builds clang if the
       # executable is missing.
       add_custom_command(OUTPUT ${outfile}
         DEPENDS clang
         APPEND
       )
     endif()
     set_property(DIRECTORY APPEND PROPERTY ADDITIONAL_MAKE_CLEAN_FILES ${outfile})
 
     list(APPEND bc_files ${outfile})
   endforeach()
 
-  set(bclib_name "libomptarget-new-nvptx-sm_${sm}.bc")
+  set(bclib_name "libomptarget-new-${target_name}-${target_cpu}.bc")
+
 
   # Link to a bitcode library.
   add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/${bclib_name}
       COMMAND ${LINK_TOOL}
         -o ${CMAKE_CURRENT_BINARY_DIR}/${bclib_name} ${bc_files}
       DEPENDS ${bc_files}
       COMMENT "Linking LLVM bitcode ${bclib_name}"
   )
 
   add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/${bclib_name}_opt
macro(instantiate_DeviceRTL)
   if("${OPT_TOOL}" STREQUAL "$<TARGET_FILE:opt>")
     add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/${bclib_name}_opt
       DEPENDS opt
       APPEND)
   endif()
 
   set_property(DIRECTORY APPEND PROPERTY ADDITIONAL_MAKE_CLEAN_FILES ${bclib_name})
 
-  set(bclib_target_name "omptarget-new-nvptx-sm_${sm}-bc")
+  set(bclib_target_name "omptarget-new-${target_name}-${target_cpu}-bc")
 
   add_custom_target(${bclib_target_name} ALL DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/${bclib_name}_opt)
-  add_dependencies(omptarget-new-nvptx-bc ${bclib_target_name})
-  add_dependencies(${bclib_target_name} opt llvm-link)
 
   # Copy library to destination.
+  # Note: This is acting on the llvm-link'ed library, not the opt'ed one
   add_custom_command(TARGET ${bclib_target_name} POST_BUILD
                     COMMAND ${CMAKE_COMMAND} -E copy ${CMAKE_CURRENT_BINARY_DIR}/${bclib_name}
                     ${LIBOMPTARGET_LIBRARY_DIR})
 
   # Install bitcode library under the lib destination folder.
   install(FILES ${CMAKE_CURRENT_BINARY_DIR}/${bclib_name} DESTINATION "${OPENMP_INSTALL_LIBDIR}")
+endmacro()
+
+# Generate a Bitcode library for all the compute capabilities the user requested
+foreach(sm ${nvptx_sm_list})
+  set(target_cpu sm_${sm})
+  set(target_name nvptx)
+  
+  # TODO: replace CUDA_ARCH with declare variant and isa selector.
+  set(target_bc_flags
+    -target nvptx64
+    -Xclang -target-feature
+    -Xclang +ptx61
+    "-D__CUDA_ARCH__=${sm}0")
+
+  instantiate_DeviceRTL()
+endforeach()
+
+foreach(mcpu ${amdgpu_mcpus})
+  set(target_cpu ${mcpu})
+  set(target_name amdgpu)
+
+  set(target_bc_flags
+    -target amdgcn-amd-amdhsa
+    "-D__AMDGCN__"
+    -fvisibility=default
+    -nogpulib)
+
+  instantiate_DeviceRTL()
 endforeach()

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/Synchronization.cpp

 }
 ///}
 
 /// 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, "");
+uint32_t atomicInc(uint32_t *A, uint32_t V, int Ordering) {
+  // builtin_amdgcn_atomic_inc32 should expand to this switch when
+  // passed a runtime value, but does not do so yet. Workaround here.
+  switch (Ordering) {
+  default:
+    __builtin_unreachable();
+  case __ATOMIC_RELAXED:
+    return __builtin_amdgcn_atomic_inc32(A, V, __ATOMIC_RELAXED, "");
+  case __ATOMIC_ACQUIRE:
+    return __builtin_amdgcn_atomic_inc32(A, V, __ATOMIC_ACQUIRE, "");
+  case __ATOMIC_RELEASE:
+    return __builtin_amdgcn_atomic_inc32(A, V, __ATOMIC_RELEASE, "");
+  case __ATOMIC_ACQ_REL:
+    return __builtin_amdgcn_atomic_inc32(A, V, __ATOMIC_ACQ_REL, "");
+  case __ATOMIC_SEQ_CST:
+    return __builtin_amdgcn_atomic_inc32(A, V, __ATOMIC_SEQ_CST, "");
+  }
 }
 
 uint32_t SHARED(namedBarrierTracker);
 
 void namedBarrierInit() {
   // Don't have global ctors, and shared memory is not zero init
   atomic::store(&namedBarrierTracker, 0u, __ATOMIC_RELEASE);
 }
if ((load & 0x0000ffffu) == (NumWaves - 1)) {
         __builtin_amdgcn_s_sleep(0);
         load = atomic::load(&namedBarrierTracker, __ATOMIC_RELAXED);
       } while ((load & 0xffff0000u) == generation);
     }
   }
   fence::team(__ATOMIC_RELEASE);
 }
 
+// sema checking of amdgcn_fence is aggressive. Intention is to patch clang
+// so that it is usable within a template environment and so that a runtime
+// value of the memory order is expanded to this switch within clang/llvm.
+void fenceTeam(int Ordering) {
+  switch (Ordering) {
+  default:
+    __builtin_unreachable();
+  case __ATOMIC_ACQUIRE:
+    return __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "workgroup");
+  case __ATOMIC_RELEASE:
+    return __builtin_amdgcn_fence(__ATOMIC_RELEASE, "workgroup");
+  case __ATOMIC_ACQ_REL:
+    return __builtin_amdgcn_fence(__ATOMIC_ACQ_REL, "workgroup");
+  case __ATOMIC_SEQ_CST:
+    return __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "workgroup");
+  }
+}
+void fenceKernel(int Ordering) {
+  switch (Ordering) {
+  default:
+    __builtin_unreachable();
+  case __ATOMIC_ACQUIRE:
+    return __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "agent");
+  case __ATOMIC_RELEASE:
+    return __builtin_amdgcn_fence(__ATOMIC_RELEASE, "agent");
+  case __ATOMIC_ACQ_REL:
+    return __builtin_amdgcn_fence(__ATOMIC_ACQ_REL, "agent");
+  case __ATOMIC_SEQ_CST:
+    return __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "agent");
+  }
+}
+void fenceSystem(int Ordering) {
+  switch (Ordering) {
+  default:
+    __builtin_unreachable();
+  case __ATOMIC_ACQUIRE:
+    return __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "");
+  case __ATOMIC_RELEASE:
+    return __builtin_amdgcn_fence(__ATOMIC_RELEASE, "");
+  case __ATOMIC_ACQ_REL:
+    return __builtin_amdgcn_fence(__ATOMIC_ACQ_REL, "");
+  case __ATOMIC_SEQ_CST:
+    return __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "");
+  }
+}
+
 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(                                       \

 
 void fence::system(int Ordering) { impl::fenceSystem(Ordering); }
 
 uint32_t atomic::load(uint32_t *Addr, int Ordering) {
   return impl::atomicLoad(Addr, Ordering);
 }
 
 void atomic::store(uint32_t *Addr, uint32_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);