replica
/
hanchenye-llvm-project
2
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

Revert "[libomptarget] Build DeviceRTL for amdgpu" 

- more tests failing on CI than failed locally when writing this patch

This reverts commit 33427fdb7b.
This commit is contained in:
Jon Chesterfield 2021-10-28 01:01:53 +01:00
parent cf37a94c1e
commit 6c7b203d1d
22 changed files with 15 additions and 91 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
clang/lib/Driver/ToolChains/AMDGPUOpenMP.cpp
View File

@ -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;

3
openmp/libomptarget/DeviceRTL/CMakeLists.txt
View File

@ -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()

4
openmp/libomptarget/DeviceRTL/src/Configuration.cpp
View File

@ -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));

78
openmp/libomptarget/DeviceRTL/src/Synchronization.cpp
View File

@ -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
///}

1
openmp/libomptarget/plugins/amdgpu/CMakeLists.txt
View File

@ -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)

1
openmp/libomptarget/test/mapping/data_member_ref.cpp
View File

@ -2,7 +2,6 @@
// amdgcn does not have printf definition
// XFAIL: amdgcn-amd-amdhsa
// XFAIL: amdgcn-amd-amdhsa-newRTL
#include <stdio.h>

1
openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers.cpp
View File

@ -2,7 +2,6 @@
// amdgcn does not have printf definition
// XFAIL: amdgcn-amd-amdhsa
// XFAIL: amdgcn-amd-amdhsa-newRTL
#include <cstdio>
#include <cstdlib>

1
openmp/libomptarget/test/mapping/declare_mapper_nested_mappers.cpp
View File

@ -2,7 +2,6 @@
// amdgcn does not have printf definition
// XFAIL: amdgcn-amd-amdhsa
// XFAIL: amdgcn-amd-amdhsa-newRTL
#include <cstdio>
#include <cstdlib>

1
openmp/libomptarget/test/mapping/delete_inf_refcount.c
View File

@ -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>

1
openmp/libomptarget/test/mapping/lambda_by_value.cpp
View File

@ -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>

1
openmp/libomptarget/test/mapping/ompx_hold/struct.c
View File

@ -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>

1
openmp/libomptarget/test/mapping/ptr_and_obj_motion.c
View File

@ -2,7 +2,6 @@
// amdgcn does not have printf definition
// XFAIL: amdgcn-amd-amdhsa
// XFAIL: amdgcn-amd-amdhsa-newRTL
#include <stdio.h>

1
openmp/libomptarget/test/mapping/reduction_implicit_map.cpp
View File

@ -2,7 +2,6 @@
// amdgcn does not have printf definition
// UNSUPPORTED: amdgcn-amd-amdhsa
// UNSUPPORTED: amdgcn-amd-amdhsa-newRTL
#include <stdio.h>

1
openmp/libomptarget/test/offloading/bug49021.cpp
View File

@ -2,7 +2,6 @@
// Wrong results on amdgcn
// UNSUPPORTED: amdgcn-amd-amdhsa
// UNSUPPORTED: amdgcn-amd-amdhsa-newRTL
#include <iostream>

2
openmp/libomptarget/test/offloading/bug49334.cpp
View File

@ -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>

1
openmp/libomptarget/test/offloading/bug50022.cpp
View File

@ -1,7 +1,6 @@
// RUN: %libomptarget-compilexx-and-run-generic
// UNSUPPORTED: amdgcn-amd-amdhsa
// UNSUPPORTED: amdgcn-amd-amdhsa-newRTL
#include <cassert>
#include <iostream>

1
openmp/libomptarget/test/offloading/global_constructor.cpp
View File

@ -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>

1
openmp/libomptarget/test/offloading/host_as_target.c
View File

@ -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>

1
openmp/libomptarget/test/unified_shared_memory/api.c
View File

@ -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>

1
openmp/libomptarget/test/unified_shared_memory/close_enter_exit.c
View File

@ -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>

1
openmp/libomptarget/test/unified_shared_memory/close_modifier.c
View File

@ -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>

1
openmp/libomptarget/test/unified_shared_memory/shared_update.c
View File

@ -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>