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> _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
