Skip to content

Commit ba3b5fb

Browse files
Pierre-vhPierre-vh
committed
[AMDGPU] Introduce Code Object V6 (llvm#76954)
Introduce Code Object V6 in Clang, LLD, Flang and LLVM. This is the same as V5 except a new "generic version" flag can be present in EFLAGS. This is related to new generic targets that'll be added in a follow-up patch. It's also likely V6 will have new changes (possibly new metadata entries) added later. Docs change are part of the follow-up patch llvm#76955 Change-Id: I6c4311bf124de1455eb3e5eecb2df7f98deb71bb
1 parent 6387cf2 commit ba3b5fb

File tree

54 files changed

+503
-162
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

54 files changed

+503
-162
lines changed

‎clang/include/clang/Basic/DiagnosticDriverKinds.td‎

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -94,6 +94,10 @@ def err_drv_hipspv_no_hip_path : Error<
9494
"'--hip-path' must be specified when offloading to "
9595
"SPIR-V%select{| unless %1 is given}0.">;
9696

97+
// TODO: Remove when COV6 is fully supported by ROCm.
98+
def warn_drv_amdgpu_cov6: Warning<
99+
"code object v6 is still in development and not ready for production use yet;"
100+
" use at your own risk">;
97101
def err_drv_undetermined_gpu_arch : Error<
98102
"cannot determine %0 architecture: %1; consider passing it via "
99103
"'%2'">;

‎clang/include/clang/Driver/Options.td‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -4856,9 +4856,9 @@ defm amdgpu_ieee : BoolMOption<"amdgpu-ieee",
48564856
def mcode_object_version_EQ : Joined<["-"], "mcode-object-version=">, Group<m_Group>,
48574857
HelpText<"Specify code object ABI version. Defaults to 4. (AMDGPU only)">,
48584858
Visibility<[ClangOption, FlangOption, CC1Option, FC1Option]>,
4859-
Values<"none,4,5">,
4859+
Values<"none,4,5,6">,
48604860
NormalizedValuesScope<"llvm::CodeObjectVersionKind">,
4861-
NormalizedValues<["COV_None", "COV_4", "COV_5"]>,
4861+
NormalizedValues<["COV_None", "COV_4", "COV_5", "COV_6"]>,
48624862
MarshallingInfoEnum<TargetOpts<"CodeObjectVersion">, "COV_5">;
48634863

48644864
defm cumode : SimpleMFlag<"cumode",

‎clang/lib/CodeGen/CGBuiltin.cpp‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -17725,9 +17725,9 @@ Value *EmitAMDGPUImplicitArgPtr(CodeGenFunction &CGF) {
1772517725
// \p Index is 0, 1, and 2 for x, y, and z dimension, respectively.
1772617726
/// Emit code based on Code Object ABI version.
1772717727
/// COV_4 : Emit code to use dispatch ptr
17728-
/// COV_5 : Emit code to use implicitarg ptr
17728+
/// COV_5+ : Emit code to use implicitarg ptr
1772917729
/// COV_NONE : Emit code to load a global variable "__oclc_ABI_version"
17730-
/// and use its value for COV_4 or COV_5 approach. It is used for
17730+
/// and use its value for COV_4 or COV_5+ approach. It is used for
1773117731
/// compiling device libraries in an ABI-agnostic way.
1773217732
///
1773317733
/// Note: "__oclc_ABI_version" is supposed to be emitted and intialized by
@@ -17774,7 +17774,7 @@ Value *EmitAMDGPUWorkGroupSize(CodeGenFunction &CGF, unsigned Index) {
1777417774
Address(Result, CGF.Int16Ty, CharUnits::fromQuantity(2)));
1777517775
} else {
1777617776
Value *GEP = nullptr;
17777-
if (Cov == CodeObjectVersionKind::COV_5) {
17777+
if (Cov >= CodeObjectVersionKind::COV_5) {
1777817778
// Indexing the implicit kernarg segment.
1777917779
GEP = CGF.Builder.CreateConstGEP1_32(
1778017780
CGF.Int8Ty, EmitAMDGPUImplicitArgPtr(CGF), 12 + Index * 2);

‎clang/lib/Driver/ToolChains/CommonArgs.cpp‎

Lines changed: 7 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2755,7 +2755,7 @@ getAMDGPUCodeObjectArgument(const Driver &D, const llvm::opt::ArgList &Args) {
27552755
void tools::checkAMDGPUCodeObjectVersion(const Driver &D,
27562756
const llvm::opt::ArgList &Args) {
27572757
const unsigned MinCodeObjVer = 4;
2758-
const unsigned MaxCodeObjVer = 5;
2758+
const unsigned MaxCodeObjVer = 6;
27592759

27602760
if (auto *CodeObjArg = getAMDGPUCodeObjectArgument(D, Args)) {
27612761
if (CodeObjArg->getOption().getID() ==
@@ -2766,6 +2766,12 @@ void tools::checkAMDGPUCodeObjectVersion(const Driver &D,
27662766
if (Remnant || CodeObjVer < MinCodeObjVer || CodeObjVer > MaxCodeObjVer)
27672767
D.Diag(diag::err_drv_invalid_int_value)
27682768
<< CodeObjArg->getAsString(Args) << CodeObjArg->getValue();
2769+
2770+
// COV6 is only supported by LLVM at the time of writing this, and it's
2771+
// expected to take some time before all ROCm components fully
2772+
// support it. In the meantime, make sure users are aware of this.
2773+
if (CodeObjVer == 6)
2774+
D.Diag(diag::warn_drv_amdgpu_cov6);
27692775
}
27702776
}
27712777
}

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

Lines changed: 37 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4,6 +4,9 @@
44
// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm-bc \
55
// RUN: -mcode-object-version=5 -DUSER -x hip -o %t_5.bc %s
66

7+
// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm-bc \
8+
// RUN: -mcode-object-version=6 -DUSER -x hip -o %t_6.bc %s
9+
710
// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm-bc \
811
// RUN: -mcode-object-version=none -DDEVICELIB -x hip -o %t_0.bc %s
912

@@ -15,6 +18,10 @@
1518
// RUN: %t_5.bc -mlink-builtin-bitcode %t_0.bc -o - |\
1619
// RUN: FileCheck -check-prefix=LINKED5 %s
1720

21+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -emit-llvm -O3 \
22+
// RUN: %t_6.bc -mlink-builtin-bitcode %t_0.bc -o - |\
23+
// RUN: FileCheck -check-prefix=LINKED6 %s
24+
1825
#include "Inputs/cuda.h"
1926

2027
// LINKED4: @__oclc_ABI_version = weak_odr hidden local_unnamed_addr addrspace(4) constant i32 400
@@ -77,6 +84,36 @@
7784
// LINKED5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
7885
// LINKED5: "amdgpu_code_object_version", i32 500
7986

87+
// LINKED6: __oclc_ABI_version = weak_odr hidden local_unnamed_addr addrspace(4) constant i32 600
88+
// LINKED6-LABEL: bar
89+
// LINKED6-NOT: load i32, ptr addrspacecast (ptr addrspace(4) @__oclc_ABI_version to ptr), align {{.*}}
90+
// LINKED6-NOT: icmp sge i32 %{{.*}}, 500
91+
// LINKED6: call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
92+
// LINKED6: [[GEP_5_X:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 12
93+
// LINKED6: call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
94+
// LINKED6: [[GEP_4_X:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 4
95+
// LINKED6: select i1 true, ptr addrspace(4) [[GEP_5_X]], ptr addrspace(4) [[GEP_4_X]]
96+
// LINKED6: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
97+
98+
// LINKED6-NOT: load i32, ptr addrspacecast (ptr addrspace(4) @__oclc_ABI_version to ptr), align {{.*}}
99+
// LINKED6-NOT: icmp sge i32 %{{.*}}, 500
100+
// LINKED6: call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
101+
// LINKED6: [[GEP_5_Y:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 14
102+
// LINKED6: call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
103+
// LINKED6: [[GEP_4_Y:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 6
104+
// LINKED6: select i1 true, ptr addrspace(4) [[GEP_5_Y]], ptr addrspace(4) [[GEP_4_Y]]
105+
// LINKED6: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
106+
107+
// LINKED6-NOT: load i32, ptr addrspacecast (ptr addrspace(4) @__oclc_ABI_version to ptr), align {{.*}}
108+
// LINKED6-NOT: icmp sge i32 %{{.*}}, 500
109+
// LINKED6: call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
110+
// LINKED6: [[GEP_5_Z:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 16
111+
// LINKED6: call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
112+
// LINKED6: [[GEP_4_Z:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 8
113+
// LINKED6: select i1 true, ptr addrspace(4) [[GEP_5_Z]], ptr addrspace(4) [[GEP_4_Z]]
114+
// LINKED6: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
115+
// LINKED6: "amdgpu_code_object_version", i32 600
116+
80117
#ifdef DEVICELIB
81118
__device__ void bar(int *x, int *y, int *z)
82119
{

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

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -9,6 +9,9 @@
99
// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm \
1010
// RUN: -mcode-object-version=5 -o - %s | FileCheck -check-prefix=V5 %s
1111

12+
// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm \
13+
// RUN: -mcode-object-version=6 -o - %s | FileCheck -check-prefix=V6 %s
14+
1215
// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm \
1316
// RUN: -mcode-object-version=none -o - %s | FileCheck %s -check-prefix=NONE
1417

@@ -17,5 +20,6 @@
1720

1821
// V4: !{{.*}} = !{i32 1, !"amdgpu_code_object_version", i32 400}
1922
// V5: !{{.*}} = !{i32 1, !"amdgpu_code_object_version", i32 500}
23+
// V6: !{{.*}} = !{i32 1, !"amdgpu_code_object_version", i32 600}
2024
// NONE-NOT: !{{.*}} = !{i32 1, !"amdgpu_code_object_version",
2125
// INV: error: invalid value '4.1' in '-mcode-object-version=4.1'

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

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -7,6 +7,10 @@
77
// RUN: -fcuda-is-device -emit-llvm -o - -x hip %s \
88
// RUN: | FileCheck -check-prefix=COV5 %s
99

10+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
11+
// RUN: -fcuda-is-device -mcode-object-version=6 -emit-llvm -o - -x hip %s \
12+
// RUN: | FileCheck -check-prefix=COV5 %s
13+
1014
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
1115
// RUN: -fcuda-is-device -mcode-object-version=none -emit-llvm -o - -x hip %s \
1216
// RUN: | FileCheck -check-prefix=COVNONE %s

‎clang/test/Driver/Inputs/rocm/amdgcn/bitcode/oclc_abi_version_600.bc‎

Whitespace-only changes.

‎clang/test/Driver/hip-code-object-version.hip‎

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -23,6 +23,19 @@
2323
// V5: "-mllvm" "--amdhsa-code-object-version=5"
2424
// V5: "-targets=host-x86_64-unknown-linux,hipv4-amdgcn-amd-amdhsa--gfx906"
2525

26+
// Check bundle ID for code object version 6.
27+
28+
// RUN: not %clang -### --target=x86_64-linux-gnu \
29+
// RUN: -mcode-object-version=6 \
30+
// RUN: --offload-arch=gfx906 --rocm-path=%S/Inputs/rocm \
31+
// RUN: %s 2>&1 | FileCheck -check-prefix=V6 %s
32+
33+
// V6: warning: code object v6 is still in development and not ready for production use yet; use at your own risk
34+
// V6: "-mcode-object-version=6"
35+
// V6: "-mllvm" "--amdhsa-code-object-version=6"
36+
// V6: "-targets=host-x86_64-unknown-linux,hipv4-amdgcn-amd-amdhsa--gfx906"
37+
38+
2639
// Check bundle ID for code object version default
2740

2841
// RUN: %clang -### --target=x86_64-linux-gnu \

‎clang/test/Driver/hip-device-libs.hip‎

Lines changed: 16 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -194,6 +194,19 @@
194194
// RUN: --rocm-path=%S/Inputs/rocm %S/Inputs/hip_multiple_inputs/b.hip \
195195
// RUN: 2>&1 | FileCheck %s --check-prefixes=NOABI5
196196

197+
// Test -mcode-object-version=6
198+
// RUN: %clang -### --target=x86_64-linux-gnu --offload-arch=gfx900 \
199+
// RUN: -mcode-object-version=6 \
200+
// RUN: --rocm-path=%S/Inputs/rocm %S/Inputs/hip_multiple_inputs/b.hip \
201+
// RUN: 2>&1 | FileCheck %s --check-prefixes=ABI6
202+
203+
// Test -mcode-object-version=6 with old device library without abi_version_600.bc
204+
// RUN: not %clang -### --target=x86_64-linux-gnu --offload-arch=gfx900 \
205+
// RUN: -mcode-object-version=6 \
206+
// RUN: --hip-device-lib-path=%S/Inputs/rocm/amdgcn/bitcode-no-abi-ver \
207+
// RUN: --rocm-path=%S/Inputs/rocm %S/Inputs/hip_multiple_inputs/b.hip \
208+
// RUN: 2>&1 | FileCheck %s --check-prefixes=NOABI6
209+
197210
// ALL-NOT: error:
198211
// ALL: {{"[^"]*clang[^"]*"}}
199212

@@ -237,7 +250,10 @@
237250
// ABI4: "-mlink-builtin-bitcode" "{{.*}}oclc_abi_version_400.bc"
238251
// ABI5-NOT: error:
239252
// ABI5: "-mlink-builtin-bitcode" "{{.*}}oclc_abi_version_500.bc"
253+
// ABI6-NOT: error:
254+
// ABI6: "-mlink-builtin-bitcode" "{{.*}}oclc_abi_version_600.bc"
240255
// NOABI4-NOT: error:
241256
// NOABI4-NOT: "-mlink-builtin-bitcode" "{{.*}}oclc_abi_version_400.bc"
242257
// NOABI4-NOT: "-mlink-builtin-bitcode" "{{.*}}oclc_abi_version_500.bc"
243258
// NOABI5: error: cannot find ROCm device libraryfor ABI version 5; provide its path via '--rocm-path' or '--rocm-device-lib-path', or pass '-nogpulib' to build without ROCm device library
259+
// NOABI6: error: cannot find ROCm device libraryfor ABI version 6; provide its path via '--rocm-path' or '--rocm-device-lib-path', or pass '-nogpulib' to build without ROCm device library

0 commit comments

Comments
 (0)