[Openmp-commits] [openmp] 6c7b203 - Revert "[libomptarget] Build DeviceRTL for amdgpu"

Jon Chesterfield via Openmp-commits openmp-commits at lists.llvm.org
Wed Oct 27 17:02:31 PDT 2021


Author: Jon Chesterfield
Date: 2021-10-28T01:01:53+01:00
New Revision: 6c7b203d1d7000269215ab5b3d329ab03dc85e42

URL: https://github.com/llvm/llvm-project/commit/6c7b203d1d7000269215ab5b3d329ab03dc85e42
DIFF: https://github.com/llvm/llvm-project/commit/6c7b203d1d7000269215ab5b3d329ab03dc85e42.diff

LOG: Revert "[libomptarget] Build DeviceRTL for amdgpu"
 - more tests failing on CI than failed locally when writing this patch

This reverts commit 33427fdb7b52b79ce5e25b7e14e0f1a44d876bd2.

Added: 
    

Modified: 
    clang/lib/Driver/ToolChains/AMDGPUOpenMP.cpp
    openmp/libomptarget/DeviceRTL/CMakeLists.txt
    openmp/libomptarget/DeviceRTL/src/Configuration.cpp
    openmp/libomptarget/DeviceRTL/src/Synchronization.cpp
    openmp/libomptarget/plugins/amdgpu/CMakeLists.txt
    openmp/libomptarget/test/mapping/data_member_ref.cpp
    openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers.cpp
    openmp/libomptarget/test/mapping/declare_mapper_nested_mappers.cpp
    openmp/libomptarget/test/mapping/delete_inf_refcount.c
    openmp/libomptarget/test/mapping/lambda_by_value.cpp
    openmp/libomptarget/test/mapping/ompx_hold/struct.c
    openmp/libomptarget/test/mapping/ptr_and_obj_motion.c
    openmp/libomptarget/test/mapping/reduction_implicit_map.cpp
    openmp/libomptarget/test/offloading/bug49021.cpp
    openmp/libomptarget/test/offloading/bug49334.cpp
    openmp/libomptarget/test/offloading/bug50022.cpp
    openmp/libomptarget/test/offloading/global_constructor.cpp
    openmp/libomptarget/test/offloading/host_as_target.c
    openmp/libomptarget/test/unified_shared_memory/api.c
    openmp/libomptarget/test/unified_shared_memory/close_enter_exit.c
    openmp/libomptarget/test/unified_shared_memory/close_modifier.c
    openmp/libomptarget/test/unified_shared_memory/shared_update.c

Removed: 
    


################################################################################
diff  --git a/clang/lib/Driver/ToolChains/AMDGPUOpenMP.cpp b/clang/lib/Driver/ToolChains/AMDGPUOpenMP.cpp
index b138000f8cf2..5400e2617729 100644
--- a/clang/lib/Driver/ToolChains/AMDGPUOpenMP.cpp
+++ b/clang/lib/Driver/ToolChains/AMDGPUOpenMP.cpp
@@ -252,7 +252,7 @@ void AMDGPUOpenMPToolChain::addClangTargetOptions(
   std::string BitcodeSuffix;
   if (DriverArgs.hasFlag(options::OPT_fopenmp_target_new_runtime,
                          options::OPT_fno_openmp_target_new_runtime, false))
-    BitcodeSuffix = "new-amdgpu-" + GPUArch;
+    BitcodeSuffix = "new-amdgcn-" + GPUArch;
   else
     BitcodeSuffix = "amdgcn-" + GPUArch;
 

diff  --git a/openmp/libomptarget/DeviceRTL/CMakeLists.txt b/openmp/libomptarget/DeviceRTL/CMakeLists.txt
index 419c64d38116..a4f9862fb09b 100644
--- a/openmp/libomptarget/DeviceRTL/CMakeLists.txt
+++ b/openmp/libomptarget/DeviceRTL/CMakeLists.txt
@@ -226,5 +226,6 @@ foreach(sm ${nvptx_sm_list})
 endforeach()
 
 foreach(mcpu ${amdgpu_mcpus})
-  compileDeviceRTLLibrary(${mcpu} amdgpu -target amdgcn-amd-amdhsa -D__AMDGCN__ -fvisibility=default -nogpulib)
+  # require D112227 or similar to enable the compilation for amdgpu
+  # compileDeviceRTLLibrary(${mcpu} amdgpu -target amdgcn-amd-amdhsa -D__AMDGCN__ -fvisibility=default -nogpulib)
 endforeach()

diff  --git a/openmp/libomptarget/DeviceRTL/src/Configuration.cpp b/openmp/libomptarget/DeviceRTL/src/Configuration.cpp
index f7c61dc013cf..2b6f20fb1732 100644
--- a/openmp/libomptarget/DeviceRTL/src/Configuration.cpp
+++ b/openmp/libomptarget/DeviceRTL/src/Configuration.cpp
@@ -20,9 +20,9 @@ using namespace _OMP;
 
 #pragma omp declare target
 
-extern uint32_t __omp_rtl_debug_kind; // defined by CGOpenMPRuntimeGPU
+extern uint32_t __omp_rtl_debug_kind;
 
-// TODO: We want to change the name as soon as the old runtime is gone.
+// TOOD: We want to change the name as soon as the old runtime is gone.
 DeviceEnvironmentTy CONSTANT(omptarget_device_environment)
     __attribute__((used));
 

diff  --git a/openmp/libomptarget/DeviceRTL/src/Synchronization.cpp b/openmp/libomptarget/DeviceRTL/src/Synchronization.cpp
index 931dffcaa131..46e7701a4872 100644
--- a/openmp/libomptarget/DeviceRTL/src/Synchronization.cpp
+++ b/openmp/libomptarget/DeviceRTL/src/Synchronization.cpp
@@ -68,23 +68,8 @@ uint64_t atomicAdd(uint64_t *Address, uint64_t Val, int Ordering) {
 ///{
 #pragma omp begin declare variant match(device = {arch(amdgcn)})
 
-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 atomicInc(uint32_t *Address, uint32_t Val, int Ordering) {
+  return __builtin_amdgcn_atomic_inc32(Address, Val, Ordering, "");
 }
 
 uint32_t SHARED(namedBarrierTracker);
@@ -141,52 +126,6 @@ void namedBarrier() {
   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
 }
@@ -194,12 +133,13 @@ void syncWarp(__kmpc_impl_lanemask_t) {
 void syncThreads() { __builtin_amdgcn_s_barrier(); }
 void syncThreadsAligned() { syncThreads(); }
 
-// 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(); }
+void syncThreadsAligned() { syncThreads(); }
+
+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, ""); }
 
 #pragma omp end declare variant
 ///}

diff  --git a/openmp/libomptarget/plugins/amdgpu/CMakeLists.txt b/openmp/libomptarget/plugins/amdgpu/CMakeLists.txt
index fb1e0dd2c105..0b830f631e90 100644
--- a/openmp/libomptarget/plugins/amdgpu/CMakeLists.txt
+++ b/openmp/libomptarget/plugins/amdgpu/CMakeLists.txt
@@ -122,4 +122,3 @@ 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)

diff  --git a/openmp/libomptarget/test/mapping/data_member_ref.cpp b/openmp/libomptarget/test/mapping/data_member_ref.cpp
index dff5987775eb..ec238907efc1 100644
--- a/openmp/libomptarget/test/mapping/data_member_ref.cpp
+++ b/openmp/libomptarget/test/mapping/data_member_ref.cpp
@@ -2,7 +2,6 @@
 
 // amdgcn does not have printf definition
 // XFAIL: amdgcn-amd-amdhsa
-// XFAIL: amdgcn-amd-amdhsa-newRTL
 
 #include <stdio.h>
 

diff  --git a/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers.cpp b/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers.cpp
index 7825d98c05c1..7edd7db880fb 100644
--- a/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers.cpp
+++ b/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers.cpp
@@ -2,7 +2,6 @@
 
 // amdgcn does not have printf definition
 // XFAIL: amdgcn-amd-amdhsa
-// XFAIL: amdgcn-amd-amdhsa-newRTL
 
 #include <cstdio>
 #include <cstdlib>

diff  --git a/openmp/libomptarget/test/mapping/declare_mapper_nested_mappers.cpp b/openmp/libomptarget/test/mapping/declare_mapper_nested_mappers.cpp
index bf2adddfccfb..c8986dd66f2c 100644
--- a/openmp/libomptarget/test/mapping/declare_mapper_nested_mappers.cpp
+++ b/openmp/libomptarget/test/mapping/declare_mapper_nested_mappers.cpp
@@ -2,7 +2,6 @@
 
 // amdgcn does not have printf definition
 // XFAIL: amdgcn-amd-amdhsa
-// XFAIL: amdgcn-amd-amdhsa-newRTL
 
 #include <cstdio>
 #include <cstdlib>

diff  --git a/openmp/libomptarget/test/mapping/delete_inf_refcount.c b/openmp/libomptarget/test/mapping/delete_inf_refcount.c
index c6d2bda187a9..cd67dddc664b 100644
--- a/openmp/libomptarget/test/mapping/delete_inf_refcount.c
+++ b/openmp/libomptarget/test/mapping/delete_inf_refcount.c
@@ -2,7 +2,6 @@
 
 // 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>

diff  --git a/openmp/libomptarget/test/mapping/lambda_by_value.cpp b/openmp/libomptarget/test/mapping/lambda_by_value.cpp
index 9cd38339b5cf..6e353244315d 100644
--- a/openmp/libomptarget/test/mapping/lambda_by_value.cpp
+++ b/openmp/libomptarget/test/mapping/lambda_by_value.cpp
@@ -2,7 +2,6 @@
 
 // amdgcn does not have printf definition
 // XFAIL: amdgcn-amd-amdhsa
-// XFAIL: amdgcn-amd-amdhsa-newRTL
 
 #include <stdio.h>
 #include <stdint.h>

diff  --git a/openmp/libomptarget/test/mapping/ompx_hold/struct.c b/openmp/libomptarget/test/mapping/ompx_hold/struct.c
index fc63e8626d01..2a0626b5fbae 100644
--- a/openmp/libomptarget/test/mapping/ompx_hold/struct.c
+++ b/openmp/libomptarget/test/mapping/ompx_hold/struct.c
@@ -3,7 +3,6 @@
 
 // amdgcn does not have printf definition
 // XFAIL: amdgcn-amd-amdhsa
-// XFAIL: amdgcn-amd-amdhsa-newRTL
 
 #include <omp.h>
 #include <stdio.h>

diff  --git a/openmp/libomptarget/test/mapping/ptr_and_obj_motion.c b/openmp/libomptarget/test/mapping/ptr_and_obj_motion.c
index 485256134041..ddea2fb65cba 100644
--- a/openmp/libomptarget/test/mapping/ptr_and_obj_motion.c
+++ b/openmp/libomptarget/test/mapping/ptr_and_obj_motion.c
@@ -2,7 +2,6 @@
 
 // amdgcn does not have printf definition
 // XFAIL: amdgcn-amd-amdhsa
-// XFAIL: amdgcn-amd-amdhsa-newRTL
 
 #include <stdio.h>
 

diff  --git a/openmp/libomptarget/test/mapping/reduction_implicit_map.cpp b/openmp/libomptarget/test/mapping/reduction_implicit_map.cpp
index 24b97bda7d57..040accd2eb4b 100644
--- a/openmp/libomptarget/test/mapping/reduction_implicit_map.cpp
+++ b/openmp/libomptarget/test/mapping/reduction_implicit_map.cpp
@@ -2,7 +2,6 @@
 
 // amdgcn does not have printf definition
 // UNSUPPORTED: amdgcn-amd-amdhsa
-// UNSUPPORTED: amdgcn-amd-amdhsa-newRTL
 
 #include <stdio.h>
 

diff  --git a/openmp/libomptarget/test/offloading/bug49021.cpp b/openmp/libomptarget/test/offloading/bug49021.cpp
index 1e456af7d1ef..521adf230bed 100644
--- a/openmp/libomptarget/test/offloading/bug49021.cpp
+++ b/openmp/libomptarget/test/offloading/bug49021.cpp
@@ -2,7 +2,6 @@
 
 // Wrong results on amdgcn
 // UNSUPPORTED: amdgcn-amd-amdhsa
-// UNSUPPORTED: amdgcn-amd-amdhsa-newRTL
 
 #include <iostream>
 

diff  --git a/openmp/libomptarget/test/offloading/bug49334.cpp b/openmp/libomptarget/test/offloading/bug49334.cpp
index 4907d32ac9c0..0ba081555453 100644
--- a/openmp/libomptarget/test/offloading/bug49334.cpp
+++ b/openmp/libomptarget/test/offloading/bug49334.cpp
@@ -2,7 +2,7 @@
 
 // Currently hangs on amdgpu
 // UNSUPPORTED: amdgcn-amd-amdhsa
-// UNSUPPORTED: amdgcn-amd-amdhsa-newRTL
+
 // UNSUPPORTED: x86_64-pc-linux-gnu
 
 #include <cassert>

diff  --git a/openmp/libomptarget/test/offloading/bug50022.cpp b/openmp/libomptarget/test/offloading/bug50022.cpp
index ca1f0e1ec3e3..a520442c835c 100644
--- a/openmp/libomptarget/test/offloading/bug50022.cpp
+++ b/openmp/libomptarget/test/offloading/bug50022.cpp
@@ -1,7 +1,6 @@
 // RUN: %libomptarget-compilexx-and-run-generic
 
 // UNSUPPORTED: amdgcn-amd-amdhsa
-// UNSUPPORTED: amdgcn-amd-amdhsa-newRTL
 
 #include <cassert>
 #include <iostream>

diff  --git a/openmp/libomptarget/test/offloading/global_constructor.cpp b/openmp/libomptarget/test/offloading/global_constructor.cpp
index ae602df8c32e..d73fe1ad938f 100644
--- a/openmp/libomptarget/test/offloading/global_constructor.cpp
+++ b/openmp/libomptarget/test/offloading/global_constructor.cpp
@@ -2,7 +2,6 @@
 
 // Fails in DAGToDAG on an address space problem
 // UNSUPPORTED: amdgcn-amd-amdhsa
-// UNSUPPORTED: amdgcn-amd-amdhsa-newRTL
 
 #include <cmath>
 #include <cstdio>

diff  --git a/openmp/libomptarget/test/offloading/host_as_target.c b/openmp/libomptarget/test/offloading/host_as_target.c
index 1e7cdef03caa..c25a4809c244 100644
--- a/openmp/libomptarget/test/offloading/host_as_target.c
+++ b/openmp/libomptarget/test/offloading/host_as_target.c
@@ -9,7 +9,6 @@
 
 // amdgcn does not have printf definition
 // XFAIL: amdgcn-amd-amdhsa
-// XFAIL: amdgcn-amd-amdhsa-newRTL
 
 #include <stdio.h>
 #include <omp.h>

diff  --git a/openmp/libomptarget/test/unified_shared_memory/api.c b/openmp/libomptarget/test/unified_shared_memory/api.c
index fcb531808edf..7282491b2a18 100644
--- a/openmp/libomptarget/test/unified_shared_memory/api.c
+++ b/openmp/libomptarget/test/unified_shared_memory/api.c
@@ -4,7 +4,6 @@
 
 // Fails on amdgcn with error: GPU Memory Error
 // XFAIL: amdgcn-amd-amdhsa
-// XFAIL: amdgcn-amd-amdhsa-newRTL
 
 #include <stdio.h>
 #include <omp.h>

diff  --git a/openmp/libomptarget/test/unified_shared_memory/close_enter_exit.c b/openmp/libomptarget/test/unified_shared_memory/close_enter_exit.c
index 62555d2eb4d9..e159ed82c25c 100644
--- a/openmp/libomptarget/test/unified_shared_memory/close_enter_exit.c
+++ b/openmp/libomptarget/test/unified_shared_memory/close_enter_exit.c
@@ -5,7 +5,6 @@
 
 // Fails on amdgcn with error: GPU Memory Error
 // XFAIL: amdgcn-amd-amdhsa
-// XFAIL: amdgcn-amd-amdhsa-newRTL
 
 #include <omp.h>
 #include <stdio.h>

diff  --git a/openmp/libomptarget/test/unified_shared_memory/close_modifier.c b/openmp/libomptarget/test/unified_shared_memory/close_modifier.c
index 98f1322ff2cb..6667fd85ec53 100644
--- a/openmp/libomptarget/test/unified_shared_memory/close_modifier.c
+++ b/openmp/libomptarget/test/unified_shared_memory/close_modifier.c
@@ -5,7 +5,6 @@
 
 // amdgcn does not have printf definition
 // XFAIL: amdgcn-amd-amdhsa
-// XFAIL: amdgcn-amd-amdhsa-newRTL
 
 #include <omp.h>
 #include <stdio.h>

diff  --git a/openmp/libomptarget/test/unified_shared_memory/shared_update.c b/openmp/libomptarget/test/unified_shared_memory/shared_update.c
index 2b90cf362ea1..ab9b3e86f0a2 100644
--- a/openmp/libomptarget/test/unified_shared_memory/shared_update.c
+++ b/openmp/libomptarget/test/unified_shared_memory/shared_update.c
@@ -4,7 +4,6 @@
 
 // amdgcn does not have printf definition
 // XFAIL: amdgcn-amd-amdhsa
-// XFAIL: amdgcn-amd-amdhsa-newRTL
 
 #include <stdio.h>
 #include <omp.h>


        


More information about the Openmp-commits mailing list