Skip to content

Commit f39d106

Browse files
jhuber6jhuber6
committed
[Clang][AMDGPU] Remove special handling for COV4 libraries
Summary: When we were first porting to COV5, this lead to some ABI issues due to a change in how we looked up the work group size. Bitcode libraries relied on the builtins to emit code, but this was changed between versions. This prevented the bitcode libraries, like OpenMP or libc, from being used for both COV4 and COV5. The solution was to have this 'none' functionality which effectively emitted code that branched off of a global to resolve to either version. This isn't a great solution because it forced every TU to have this variable in it. The patch in llvm#131033 removed support for COV4 from OpenMP, which was the only consumer of this functionality. Other users like HIP and OpenCL did not use this because they linked the ROCm Device Library directly which has its own handling (The name was borrowed from it after all). So, now that we don't need to worry about backward compatibility with COV4, we can remove this special handling. Users can still emit COV4 code, this simply removes the special handling used to make the OpenMP device runtime bitcode version agnostic.
1 parent 25bf4e2 commit f39d106

File tree

14 files changed

+29
-277
lines changed

14 files changed

+29
-277
lines changed

‎clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp‎

Lines changed: 8 additions & 47 deletions
Original file line numberDiff line numberDiff line change
@@ -62,62 +62,23 @@ Value *EmitAMDGPUImplicitArgPtr(CodeGenFunction &CGF) {
6262
/// Emit code based on Code Object ABI version.
6363
/// COV_4 : Emit code to use dispatch ptr
6464
/// COV_5+ : Emit code to use implicitarg ptr
65-
/// COV_NONE : Emit code to load a global variable "__oclc_ABI_version"
66-
/// and use its value for COV_4 or COV_5+ approach. It is used for
67-
/// compiling device libraries in an ABI-agnostic way.
68-
///
69-
/// Note: "__oclc_ABI_version" is supposed to be emitted and intialized by
70-
/// clang during compilation of user code.
7165
Value *EmitAMDGPUWorkGroupSize(CodeGenFunction &CGF, unsigned Index) {
7266
llvm::LoadInst *LD;
7367

7468
auto Cov = CGF.getTarget().getTargetOpts().CodeObjectVersion;
7569

76-
if (Cov == CodeObjectVersionKind::COV_None) {
77-
StringRef Name = "__oclc_ABI_version";
78-
auto *ABIVersionC = CGF.CGM.getModule().getNamedGlobal(Name);
79-
if (!ABIVersionC)
80-
ABIVersionC = new llvm::GlobalVariable(
81-
CGF.CGM.getModule(), CGF.Int32Ty, false,
82-
llvm::GlobalValue::ExternalLinkage, nullptr, Name, nullptr,
83-
llvm::GlobalVariable::NotThreadLocal,
84-
CGF.CGM.getContext().getTargetAddressSpace(LangAS::opencl_constant));
85-
86-
// This load will be eliminated by the IPSCCP because it is constant
87-
// weak_odr without externally_initialized. Either changing it to weak or
88-
// adding externally_initialized will keep the load.
89-
Value *ABIVersion = CGF.Builder.CreateAlignedLoad(CGF.Int32Ty, ABIVersionC,
90-
CGF.CGM.getIntAlign());
91-
92-
Value *IsCOV5 = CGF.Builder.CreateICmpSGE(
93-
ABIVersion,
94-
llvm::ConstantInt::get(CGF.Int32Ty, CodeObjectVersionKind::COV_5));
95-
70+
Value *GEP = nullptr;
71+
if (Cov >= CodeObjectVersionKind::COV_5) {
9672
// Indexing the implicit kernarg segment.
97-
Value *ImplicitGEP = CGF.Builder.CreateConstGEP1_32(
73+
GEP = CGF.Builder.CreateConstGEP1_32(
9874
CGF.Int8Ty, EmitAMDGPUImplicitArgPtr(CGF), 12 + Index * 2);
99-
100-
// Indexing the HSA kernel_dispatch_packet struct.
101-
Value *DispatchGEP = CGF.Builder.CreateConstGEP1_32(
102-
CGF.Int8Ty, EmitAMDGPUDispatchPtr(CGF), 4 + Index * 2);
103-
104-
auto Result = CGF.Builder.CreateSelect(IsCOV5, ImplicitGEP, DispatchGEP);
105-
LD = CGF.Builder.CreateLoad(
106-
Address(Result, CGF.Int16Ty, CharUnits::fromQuantity(2)));
10775
} else {
108-
Value *GEP = nullptr;
109-
if (Cov >= CodeObjectVersionKind::COV_5) {
110-
// Indexing the implicit kernarg segment.
111-
GEP = CGF.Builder.CreateConstGEP1_32(
112-
CGF.Int8Ty, EmitAMDGPUImplicitArgPtr(CGF), 12 + Index * 2);
113-
} else {
114-
// Indexing the HSA kernel_dispatch_packet struct.
115-
GEP = CGF.Builder.CreateConstGEP1_32(
116-
CGF.Int8Ty, EmitAMDGPUDispatchPtr(CGF), 4 + Index * 2);
117-
}
118-
LD = CGF.Builder.CreateLoad(
119-
Address(GEP, CGF.Int16Ty, CharUnits::fromQuantity(2)));
76+
// Indexing the HSA kernel_dispatch_packet struct.
77+
GEP = CGF.Builder.CreateConstGEP1_32(CGF.Int8Ty, EmitAMDGPUDispatchPtr(CGF),
78+
4 + Index * 2);
12079
}
80+
LD = CGF.Builder.CreateLoad(
81+
Address(GEP, CGF.Int16Ty, CharUnits::fromQuantity(2)));
12182

12283
llvm::MDBuilder MDHelper(CGF.getLLVMContext());
12384
llvm::MDNode *RNode = MDHelper.createRange(APInt(16, 1),

‎clang/lib/CodeGen/Targets/AMDGPU.cpp‎

Lines changed: 0 additions & 36 deletions
Original file line numberDiff line numberDiff line change
@@ -305,8 +305,6 @@ class AMDGPUTargetCodeGenInfo : public TargetCodeGenInfo {
305305
void setFunctionDeclAttributes(const FunctionDecl *FD, llvm::Function *F,
306306
CodeGenModule &CGM) const;
307307

308-
void emitTargetGlobals(CodeGen::CodeGenModule &CGM) const override;
309-
310308
void setTargetAttributes(const Decl *D, llvm::GlobalValue *GV,
311309
CodeGen::CodeGenModule &M) const override;
312310
unsigned getOpenCLKernelCallingConv() const override;
@@ -414,40 +412,6 @@ void AMDGPUTargetCodeGenInfo::setFunctionDeclAttributes(
414412
}
415413
}
416414

417-
/// Emits control constants used to change per-architecture behaviour in the
418-
/// AMDGPU ROCm device libraries.
419-
void AMDGPUTargetCodeGenInfo::emitTargetGlobals(
420-
CodeGen::CodeGenModule &CGM) const {
421-
StringRef Name = "__oclc_ABI_version";
422-
llvm::GlobalVariable *OriginalGV = CGM.getModule().getNamedGlobal(Name);
423-
if (OriginalGV && !llvm::GlobalVariable::isExternalLinkage(OriginalGV->getLinkage()))
424-
return;
425-
426-
if (CGM.getTarget().getTargetOpts().CodeObjectVersion ==
427-
llvm::CodeObjectVersionKind::COV_None)
428-
return;
429-
430-
auto *Type = llvm::IntegerType::getIntNTy(CGM.getModule().getContext(), 32);
431-
llvm::Constant *COV = llvm::ConstantInt::get(
432-
Type, CGM.getTarget().getTargetOpts().CodeObjectVersion);
433-
434-
// It needs to be constant weak_odr without externally_initialized so that
435-
// the load instuction can be eliminated by the IPSCCP.
436-
auto *GV = new llvm::GlobalVariable(
437-
CGM.getModule(), Type, true, llvm::GlobalValue::WeakODRLinkage, COV, Name,
438-
nullptr, llvm::GlobalValue::ThreadLocalMode::NotThreadLocal,
439-
CGM.getContext().getTargetAddressSpace(LangAS::opencl_constant));
440-
GV->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::Local);
441-
GV->setVisibility(llvm::GlobalValue::VisibilityTypes::HiddenVisibility);
442-
443-
// Replace any external references to this variable with the new global.
444-
if (OriginalGV) {
445-
OriginalGV->replaceAllUsesWith(GV);
446-
GV->takeName(OriginalGV);
447-
OriginalGV->eraseFromParent();
448-
}
449-
}
450-
451415
void AMDGPUTargetCodeGenInfo::setTargetAttributes(
452416
const Decl *D, llvm::GlobalValue *GV, CodeGen::CodeGenModule &M) const {
453417
if (requiresAMDGPUProtectedVisibility(D, GV)) {

‎clang/test/CodeGen/amdgpu-abi-version.c‎

Lines changed: 13 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -1,23 +1,24 @@
11
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals --version 3
22
// RUN: %clang_cc1 -cc1 -triple amdgcn-amd-amdhsa -emit-llvm -mcode-object-version=none %s -o - | FileCheck %s
33

4-
//.
5-
// CHECK: @__oclc_ABI_version = external addrspace(4) global i32
6-
//.
74
// CHECK-LABEL: define dso_local i32 @foo(
85
// CHECK-SAME: ) #[[ATTR0:[0-9]+]] {
96
// CHECK-NEXT: entry:
107
// CHECK-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
118
// CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
12-
// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(4) @__oclc_ABI_version, align 4
13-
// CHECK-NEXT: [[TMP1:%.*]] = icmp sge i32 [[TMP0]], 500
14-
// CHECK-NEXT: [[TMP2:%.*]] = call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
15-
// CHECK-NEXT: [[TMP3:%.*]] = getelementptr i8, ptr addrspace(4) [[TMP2]], i32 12
16-
// CHECK-NEXT: [[TMP4:%.*]] = call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
17-
// CHECK-NEXT: [[TMP5:%.*]] = getelementptr i8, ptr addrspace(4) [[TMP4]], i32 4
18-
// CHECK-NEXT: [[TMP6:%.*]] = select i1 [[TMP1]], ptr addrspace(4) [[TMP3]], ptr addrspace(4) [[TMP5]]
19-
// CHECK-NEXT: [[TMP7:%.*]] = load i16, ptr addrspace(4) [[TMP6]], align 2, !range [[RNG2:![0-9]+]], !invariant.load !3, !noundef !3
20-
// CHECK-NEXT: [[CONV:%.*]] = zext i16 [[TMP7]] to i32
9+
// CHECK-NEXT: [[TMP0:%.*]] = call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
10+
// CHECK-NEXT: [[TMP1:%.*]] = getelementptr i8, ptr addrspace(4) [[TMP0]], i32 4
11+
// CHECK-NEXT: [[TMP2:%.*]] = load i16, ptr addrspace(4) [[TMP1]], align 2, !range [[RNG2:![0-9]+]], !invariant.load [[META3:![0-9]+]], !noundef [[META3]]
12+
// CHECK-NEXT: [[CONV:%.*]] = zext i16 [[TMP2]] to i32
2113
// CHECK-NEXT: ret i32 [[CONV]]
2214
//
2315
int foo() { return __builtin_amdgcn_workgroup_size_x(); }
16+
//.
17+
// CHECK: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
18+
// CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
19+
//.
20+
// CHECK: [[META0:![0-9]+]] = !{i32 1, !"wchar_size", i32 4}
21+
// CHECK: [[META1:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"}
22+
// CHECK: [[RNG2]] = !{i16 1, i16 1025}
23+
// CHECK: [[META3]] = !{}
24+
//.

‎clang/test/CodeGen/amdgpu-address-spaces.cpp‎

Lines changed: 7 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -29,7 +29,6 @@ int [[clang::address_space(999)]] bbb = 1234;
2929
// CHECK: @u = addrspace(5) global i32 undef, align 4
3030
// CHECK: @aaa = addrspace(6) global i32 1000, align 4
3131
// CHECK: @bbb = addrspace(999) global i32 1234, align 4
32-
// CHECK: @__oclc_ABI_version = weak_odr hidden local_unnamed_addr addrspace(4) constant i32 600
3332
//.
3433
// CHECK-LABEL: define dso_local amdgpu_kernel void @foo(
3534
// CHECK-SAME: ) #[[ATTR0:[0-9]+]] {
@@ -60,3 +59,10 @@ extern "C" [[clang::amdgpu_kernel]] void foo() {
6059
aaa = 0;
6160
bbb = 0;
6261
}
62+
//.
63+
// CHECK: attributes #[[ATTR0]] = { convergent mustprogress noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
64+
//.
65+
// CHECK: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 600}
66+
// CHECK: [[META1:![0-9]+]] = !{i32 1, !"wchar_size", i32 4}
67+
// CHECK: [[META2:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"}
68+
//.

‎clang/test/CodeGenCUDA/amdgpu-code-object-version-linking.cu‎

Lines changed: 0 additions & 133 deletions
This file was deleted.

‎clang/test/CodeGenCUDA/amdgpu-workgroup-size.cu‎

Lines changed: 0 additions & 34 deletions
Original file line numberDiff line numberDiff line change
@@ -2,7 +2,6 @@
22
// RUN: -fcuda-is-device -mcode-object-version=4 -emit-llvm -o - -x hip %s \
33
// RUN: | FileCheck -check-prefix=PRECOV5 %s
44

5-
65
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
76
// RUN: -fcuda-is-device -emit-llvm -o - -x hip %s \
87
// RUN: | FileCheck -check-prefix=COV5 %s
@@ -11,10 +10,6 @@
1110
// RUN: -fcuda-is-device -mcode-object-version=6 -emit-llvm -o - -x hip %s \
1211
// RUN: | FileCheck -check-prefix=COV5 %s
1312

14-
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
15-
// RUN: -fcuda-is-device -mcode-object-version=none -emit-llvm -o - -x hip %s \
16-
// RUN: | FileCheck -check-prefix=COVNONE %s
17-
1813
#include "Inputs/cuda.h"
1914

2015
// PRECOV5-LABEL: test_get_workgroup_size
@@ -35,35 +30,6 @@
3530
// COV5: getelementptr i8, ptr addrspace(4) %{{.*}}, i32 16
3631
// COV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
3732

38-
39-
// COVNONE-LABEL: test_get_workgroup_size
40-
// COVNONE: load i32, ptr addrspace(4) @__oclc_ABI_version
41-
// COVNONE: [[ABI5_X:%.*]] = icmp sge i32 %{{.*}}, 500
42-
// COVNONE: call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
43-
// COVNONE: [[GEP_5_X:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 12
44-
// COVNONE: call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
45-
// COVNONE: [[GEP_4_X:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 4
46-
// COVNONE: select i1 [[ABI5_X]], ptr addrspace(4) [[GEP_5_X]], ptr addrspace(4) [[GEP_4_X]]
47-
// COVNONE: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
48-
49-
// COVNONE: load i32, ptr addrspace(4) @__oclc_ABI_version
50-
// COVNONE: [[ABI5_Y:%.*]] = icmp sge i32 %{{.*}}, 500
51-
// COVNONE: call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
52-
// COVNONE: [[GEP_5_Y:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 14
53-
// COVNONE: call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
54-
// COVNONE: [[GEP_4_Y:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 6
55-
// COVNONE: select i1 [[ABI5_Y]], ptr addrspace(4) [[GEP_5_Y]], ptr addrspace(4) [[GEP_4_Y]]
56-
// COVNONE: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
57-
58-
// COVNONE: load i32, ptr addrspace(4) @__oclc_ABI_version
59-
// COVNONE: [[ABI5_Z:%.*]] = icmp sge i32 %{{.*}}, 500
60-
// COVNONE: call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
61-
// COVNONE: [[GEP_5_Z:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 16
62-
// COVNONE: call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
63-
// COVNONE: [[GEP_4_Z:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 8
64-
// COVNONE: select i1 [[ABI5_Z]], ptr addrspace(4) [[GEP_5_Z]], ptr addrspace(4) [[GEP_4_Z]]
65-
// COVNONE: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
66-
6733
__device__ void test_get_workgroup_size(int d, int *out)
6834
{
6935
switch (d) {

0 commit comments

Comments
 (0)