[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–87

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

180

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

 endif()
 
 if (LLVM_DIR)
   # Builds that use pre-installed LLVM have LLVM_DIR set.
   find_program(CLANG_TOOL clang PATHS ${LLVM_TOOLS_BINARY_DIR} NO_DEFAULT_PATH)
   find_program(LINK_TOOL llvm-link PATHS ${LLVM_TOOLS_BINARY_DIR}
     NO_DEFAULT_PATH)
   find_program(OPT_TOOL opt PATHS ${LLVM_TOOLS_BINARY_DIR} NO_DEFAULT_PATH)
-  libomptarget_say("Building Device RTL. Using clang: ${CLANG_TOOL}")
+  libomptarget_say("Building DeviceRTL. Using clang: ${CLANG_TOOL}")
 elseif (LLVM_TOOL_CLANG_BUILD AND NOT CMAKE_CROSSCOMPILING AND NOT OPENMP_STANDALONE_BUILD)
   # LLVM in-tree builds may use CMake target names to discover the tools.
   set(CLANG_TOOL $<TARGET_FILE:clang>)
   set(LINK_TOOL $<TARGET_FILE:llvm-link>)
   set(OPT_TOOL $<TARGET_FILE:opt>)
   libomptarget_say("Building DeviceRTL. Using clang from in-tree build")
 else()
   libomptarget_say("Not building DeviceRTL. No appropriate clang found")

 
 set(devicertl_base_directory ${CMAKE_CURRENT_SOURCE_DIR})
 set(include_directory ${devicertl_base_directory}/include)
 set(source_directory ${devicertl_base_directory}/src)
 
 set(all_capabilities 35 37 50 52 53 60 61 62 70 72 75 80)
 
 set(LIBOMPTARGET_NVPTX_COMPUTE_CAPABILITIES ${all_capabilities} CACHE STRING
-  "List of CUDA Compute Capabilities to be used to compile the NVPTX device RTL.")
+  "List of CUDA Compute Capabilities to be used to compile the NVPTX DeviceRTL.")
 string(TOLOWER ${LIBOMPTARGET_NVPTX_COMPUTE_CAPABILITIES} LIBOMPTARGET_NVPTX_COMPUTE_CAPABILITIES)
 
 if (LIBOMPTARGET_NVPTX_COMPUTE_CAPABILITIES STREQUAL "all")
   set(nvptx_sm_list ${all_capabilities})
 elseif(LIBOMPTARGET_NVPTX_COMPUTE_CAPABILITIES STREQUAL "auto")
   if (NOT LIBOMPTARGET_DEP_CUDA_FOUND)
     libomptarget_error_say("[NVPTX] Cannot auto detect compute capability as CUDA not found.")
   endif()
   set(nvptx_sm_list ${LIBOMPTARGET_DEP_CUDA_ARCH})
 else()
   string(REPLACE "," ";" nvptx_sm_list "${LIBOMPTARGET_NVPTX_COMPUTE_CAPABILITIES}")
 endif()
 
 # If user set LIBOMPTARGET_NVPTX_COMPUTE_CAPABILITIES to empty, we disable the
 # build.
 if (NOT nvptx_sm_list)
-  libomptarget_say("Not building CUDA offloading device RTL: empty compute capability list")
+  libomptarget_say("Not building CUDA offloading DeviceRTL: empty compute capability list")
   return()
 endif()
 
 # 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)
+function(compileDeviceRTLLibrary target_cpu target_name)
+  set(target_bc_flags ${ARGN})
 
-# 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")
   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}
+  add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/linked_${bclib_name}
       COMMAND ${LINK_TOOL}
-        -o ${CMAKE_CURRENT_BINARY_DIR}/${bclib_name} ${bc_files}
+        -o ${CMAKE_CURRENT_BINARY_DIR}/linked_${bclib_name} ${bc_files}
       DEPENDS ${bc_files}
       COMMENT "Linking LLVM bitcode ${bclib_name}"
   )
 
-  add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/${bclib_name}_opt
-      COMMAND ${OPT_TOOL} ${link_opt_flags} ${CMAKE_CURRENT_BINARY_DIR}/${bclib_name}
+  add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/${bclib_name}
+      COMMAND ${OPT_TOOL} ${link_opt_flags} ${CMAKE_CURRENT_BINARY_DIR}/linked_${bclib_name}
                       -o ${CMAKE_CURRENT_BINARY_DIR}/${bclib_name}
-      DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/${bclib_name}
+      DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/linked_${bclib_name}
       COMMENT "Optimizing LLVM bitcode ${bclib_name}"
   )
 
   # Add a file-level dependency to ensure that llvm-link and opt are up-to-date.
   # By default, add_custom_command only builds the tool if the executable is missing
   if("${LINK_TOOL}" STREQUAL "$<TARGET_FILE:llvm-link>")
-    add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/${bclib_name}
+    add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/linked_${bclib_name}
       DEPENDS llvm-link
       APPEND)
   endif()
   if("${OPT_TOOL}" STREQUAL "$<TARGET_FILE:opt>")
-    add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/${bclib_name}_opt
+    add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/${bclib_name}
       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)
+  add_custom_target(${bclib_target_name} ALL DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/${bclib_name})
 
   # Copy library to destination.
   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}")
+endfunction()
+
+# Generate a Bitcode library for all the compute capabilities the user requested
+foreach(sm ${nvptx_sm_list})
+  compileDeviceRTLLibrary(sm_${sm} nvptx -target nvptx64 -Xclang -target-feature -Xclang +ptx61 "-D__CUDA_ARCH__=${sm}0")
+endforeach()
+
+foreach(mcpu ${amdgpu_mcpus})
+  compileDeviceRTLLibrary(${mcpu} amdgpu -target amdgcn-amd-amdhsa "-D__AMDGCN__" -fvisibility=default -nogpulib)
 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);

openmp/libomptarget/plugins/amdgpu/CMakeLists.txt

execute_process(COMMAND "${AMDGPU_ARCH_TOOL}" RESULT_VARIABLE amdgpu_arch_result
   OUTPUT_VARIABLE amdgpu_arch_output)
 if (${amdgpu_arch_result})
   libomptarget_say("Not generating amdgcn test targets as amdgpu-arch exited with ${amdgpu_arch_result}")
   return()
 endif()
 
 # Report to the parent scope that we are building a plugin for amdgpu
 set(LIBOMPTARGET_SYSTEM_TARGETS "${LIBOMPTARGET_SYSTEM_TARGETS} amdgcn-amd-amdhsa" PARENT_SCOPE)
+set(LIBOMPTARGET_SYSTEM_TARGETS "${LIBOMPTARGET_SYSTEM_TARGETS} amdgcn-amd-amdhsa-newRTL" PARENT_SCOPE)

openmp/libomptarget/test/mapping/data_member_ref.cpp

 // RUN: %libomptarget-compilexx-run-and-check-generic
 
 // amdgcn does not have printf definition
 // XFAIL: amdgcn-amd-amdhsa
+// XFAIL: amdgcn-amd-amdhsa-newRTL
 
 #include <stdio.h>
 
 struct View {
   int Data;
 };
 
 struct ViewPtr {

openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers.cpp

 // RUN: %libomptarget-compilexx-run-and-check-generic
 
 // amdgcn does not have printf definition
 // XFAIL: amdgcn-amd-amdhsa
+// XFAIL: amdgcn-amd-amdhsa-newRTL
 
 #include <cstdio>
 #include <cstdlib>
 
 typedef struct {
   int a;
   double *b;
 } C1;

openmp/libomptarget/test/mapping/declare_mapper_nested_mappers.cpp

 // RUN: %libomptarget-compilexx-run-and-check-generic
 
 // amdgcn does not have printf definition
 // XFAIL: amdgcn-amd-amdhsa
+// XFAIL: amdgcn-amd-amdhsa-newRTL
 
 #include <cstdio>
 #include <cstdlib>
 
 typedef struct {
   int a;
   double *b;
 } C;

openmp/libomptarget/test/mapping/delete_inf_refcount.c

 // RUN: %libomptarget-compile-run-and-check-generic
 
 // fails with error message 'Unable to generate target entries' on amdgcn
 // XFAIL: amdgcn-amd-amdhsa
+// XFAIL: amdgcn-amd-amdhsa-newRTL
 
 #include <stdio.h>
 #include <omp.h>
 
 #pragma omp declare target
 int isHost;
 #pragma omp end declare target
 

openmp/libomptarget/test/mapping/lambda_by_value.cpp

 // RUN: %libomptarget-compilexx-run-and-check-generic
 
 // amdgcn does not have printf definition
 // XFAIL: amdgcn-amd-amdhsa
+// XFAIL: amdgcn-amd-amdhsa-newRTL
 
 #include <stdio.h>
 #include <stdint.h>
 
 // CHECK: before: [[V1:111]] [[V2:222]] [[PX:0x[^ ]+]] [[PY:0x[^ ]+]]
 // CHECK: lambda: [[V1]] [[V2]] [[PX_TGT:0x[^ ]+]] 0x{{.*}}
 // CHECK: tgt   : [[V2]] [[PX_TGT]] 1
 // CHECK: out   : [[V2]] [[V2]] [[PX]] [[PY]]

openmp/libomptarget/test/mapping/ompx_hold/struct.c

 // RUN: %libomptarget-compile-generic -fopenmp-extensions
 // RUN: %libomptarget-run-generic | %fcheck-generic -strict-whitespace
 
 // amdgcn does not have printf definition
 // XFAIL: amdgcn-amd-amdhsa
+// XFAIL: amdgcn-amd-amdhsa-newRTL
 
 #include <omp.h>
 #include <stdio.h>
 
 #define CHECK_PRESENCE(Var1, Var2, Var3)                                       \
   printf("    presence of %s, %s, %s: %d, %d, %d\n",                           \
          #Var1, #Var2, #Var3,                                                  \
          omp_target_is_present(&(Var1), omp_get_default_device()),             \

openmp/libomptarget/test/mapping/ptr_and_obj_motion.c

 // RUN: %libomptarget-compile-run-and-check-generic
 
 // amdgcn does not have printf definition
 // XFAIL: amdgcn-amd-amdhsa
+// XFAIL: amdgcn-amd-amdhsa-newRTL
 
 #include <stdio.h>
 
 typedef struct {
   double *dataptr;
   int dummy1;
   int dummy2;
 } DV;

openmp/libomptarget/test/mapping/reduction_implicit_map.cpp

 // RUN: %libomptarget-compilexx-run-and-check-generic
 
 // amdgcn does not have printf definition
 // UNSUPPORTED: amdgcn-amd-amdhsa
+// UNSUPPORTED: amdgcn-amd-amdhsa-newRTL
 
 #include <stdio.h>
 
 void sum(int* input, int size, int* output)
 {
 #pragma omp target teams distribute parallel for reduction(+:output[0]) \
                                                  map(to:input[0:size])
   for (int i = 0; i < size; i++)

openmp/libomptarget/test/offloading/bug49021.cpp

 // RUN: %libomptarget-compilexx-generic -O3 && %libomptarget-run-generic
 
 // Wrong results on amdgcn
 // UNSUPPORTED: amdgcn-amd-amdhsa
+// UNSUPPORTED: amdgcn-amd-amdhsa-newRTL
 
 #include <iostream>
 
 template <typename T> int test_map() {
   std::cout << "map(complex<>)" << std::endl;
   T a(0.2), a_check;
 #pragma omp target map(from : a_check)
   { a_check = a; }

openmp/libomptarget/test/offloading/bug49334.cpp

 // RUN: %libomptarget-compilexx-run-and-check-generic
 
 // Currently hangs on amdgpu
 // UNSUPPORTED: amdgcn-amd-amdhsa
-
+// UNSUPPORTED: amdgcn-amd-amdhsa-newRTL
 // UNSUPPORTED: x86_64-pc-linux-gnu
 
 #include <cassert>
 #include <iostream>
 #include <memory>
 #include <vector>
 
 class BlockMatrix {

openmp/libomptarget/test/offloading/bug50022.cpp

 // RUN: %libomptarget-compilexx-and-run-generic
 
 // UNSUPPORTED: amdgcn-amd-amdhsa
+// UNSUPPORTED: amdgcn-amd-amdhsa-newRTL
 
 #include <cassert>
 #include <iostream>
 #include <stdexcept>
 
 int main(int argc, char *argv[]) {
   int a = 0;
   std::cout << "outside a = " << a << " addr " << &a << std::endl;

openmp/libomptarget/test/offloading/global_constructor.cpp

 // RUN: %libomptarget-compilexx-generic && %libomptarget-run-generic | %fcheck-generic
 
 // Fails in DAGToDAG on an address space problem
 // UNSUPPORTED: amdgcn-amd-amdhsa
+// UNSUPPORTED: amdgcn-amd-amdhsa-newRTL
 
 #include <cmath>
 #include <cstdio>
 
 const double Host = log(2.0) / log(2.0);
 #pragma omp declare target
 const double Device = log(2.0) / log(2.0);
 #pragma omp end declare target

openmp/libomptarget/test/offloading/host_as_target.c

 // Check that specifying device as omp_get_initial_device():
 // - Doesn't cause the runtime to fail.
 // - Offloads code to the host.
 // - Doesn't transfer data.  In this case, just check that neither host data nor
 //   default device data are affected by the specified transfers.
 // - Works whether it's specified directly or as the default device.
 
 // RUN: %libomptarget-compile-run-and-check-generic
 
 // amdgcn does not have printf definition
 // XFAIL: amdgcn-amd-amdhsa
+// XFAIL: amdgcn-amd-amdhsa-newRTL
 
 #include <stdio.h>
 #include <omp.h>
 
 static void check(char *X, int Dev) {
   printf("  host X = %c\n", *X);
   #pragma omp target device(Dev)
   printf("device X = %c\n", *X);

openmp/libomptarget/test/unified_shared_memory/api.c

 // RUN: %libomptarget-compile-run-and-check-generic
 // XFAIL: nvptx64-nvidia-cuda
 // XFAIL: nvptx64-nvidia-cuda-newRTL
 
 // Fails on amdgcn with error: GPU Memory Error
 // XFAIL: amdgcn-amd-amdhsa
+// XFAIL: amdgcn-amd-amdhsa-newRTL
 
 #include <stdio.h>
 #include <omp.h>
 
 // ---------------------------------------------------------------------------
 // Various definitions copied from OpenMP RTL
 
 extern void __tgt_register_requires(int64_t);

openmp/libomptarget/test/unified_shared_memory/close_enter_exit.c

 // RUN: %libomptarget-compile-run-and-check-generic
 
 // REQUIRES: unified_shared_memory
 // UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
 
 // Fails on amdgcn with error: GPU Memory Error
 // XFAIL: amdgcn-amd-amdhsa
+// XFAIL: amdgcn-amd-amdhsa-newRTL
 
 #include <omp.h>
 #include <stdio.h>
 
 #pragma omp requires unified_shared_memory
 
 #define N 1024
 

openmp/libomptarget/test/unified_shared_memory/close_modifier.c

 // RUN: %libomptarget-compile-run-and-check-generic
 
 // REQUIRES: unified_shared_memory
 // UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
 
 // amdgcn does not have printf definition
 // XFAIL: amdgcn-amd-amdhsa
+// XFAIL: amdgcn-amd-amdhsa-newRTL
 
 #include <omp.h>
 #include <stdio.h>
 
 #pragma omp requires unified_shared_memory
 
 #define N 1024
 

openmp/libomptarget/test/unified_shared_memory/shared_update.c

 // RUN: %libomptarget-compile-run-and-check-generic
 
 // REQUIRES: unified_shared_memory
 
 // amdgcn does not have printf definition
 // XFAIL: amdgcn-amd-amdhsa
+// XFAIL: amdgcn-amd-amdhsa-newRTL
 
 #include <stdio.h>
 #include <omp.h>
 
 // ---------------------------------------------------------------------------
 // Various definitions copied from OpenMP RTL
 
 extern void __tgt_register_requires(int64_t);