Skip to content

Commit f1ac2af

Browse files
shiltianshiltian
authored
Reapply "[AMDGPU] Use COV6 by default (#118515)" (#130963)
This reverts commit 68bcba6.
1 parent c031579 commit f1ac2af

File tree

17 files changed

+29
-18
lines changed

17 files changed

+29
-18
lines changed

‎clang/docs/ReleaseNotes.rst

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -382,6 +382,8 @@ Target Specific Changes
382382
AMDGPU Support
383383
^^^^^^^^^^^^^^
384384

385+
- Bump the default code object version to 6. ROCm 6.3 is required to run any program compiled with COV6.
386+
385387
NVPTX Support
386388
^^^^^^^^^^^^^^
387389

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -5161,12 +5161,12 @@ defm amdgpu_ieee : BoolMOption<"amdgpu-ieee",
51615161
NegFlag<SetFalse, [], [ClangOption, CC1Option]>>;
51625162

51635163
def mcode_object_version_EQ : Joined<["-"], "mcode-object-version=">, Group<m_Group>,
5164-
HelpText<"Specify code object ABI version. Defaults to 5. (AMDGPU only)">,
5164+
HelpText<"Specify code object ABI version. Defaults to 6. (AMDGPU only)">,
51655165
Visibility<[ClangOption, FlangOption, CC1Option, FC1Option]>,
51665166
Values<"none,4,5,6">,
51675167
NormalizedValuesScope<"llvm::CodeObjectVersionKind">,
51685168
NormalizedValues<["COV_None", "COV_4", "COV_5", "COV_6"]>,
5169-
MarshallingInfoEnum<TargetOpts<"CodeObjectVersion">, "COV_5">;
5169+
MarshallingInfoEnum<TargetOpts<"CodeObjectVersion">, "COV_6">;
51705170

51715171
defm cumode : SimpleMFlag<"cumode",
51725172
"Specify CU wavefront", "Specify WGP wavefront",

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2760,7 +2760,7 @@ void tools::checkAMDGPUCodeObjectVersion(const Driver &D,
27602760

27612761
unsigned tools::getAMDGPUCodeObjectVersion(const Driver &D,
27622762
const llvm::opt::ArgList &Args) {
2763-
unsigned CodeObjVer = 5; // default
2763+
unsigned CodeObjVer = 6; // default
27642764
if (auto *CodeObjArg = getAMDGPUCodeObjectArgument(D, Args))
27652765
StringRef(CodeObjArg->getValue()).getAsInteger(0, CodeObjVer);
27662766
return CodeObjVer;

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -29,7 +29,7 @@ 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 500
32+
// CHECK: @__oclc_ABI_version = weak_odr hidden local_unnamed_addr addrspace(4) constant i32 600
3333
//.
3434
// CHECK-LABEL: define dso_local amdgpu_kernel void @foo(
3535
// CHECK-SAME: ) #[[ATTR0:[0-9]+]] {

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
11
// Create module flag for code object version.
22

33
// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm \
4-
// RUN: -o - %s | FileCheck %s -check-prefix=V5
4+
// RUN: -o - %s | FileCheck %s -check-prefix=V6
55

66
// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm \
77
// RUN: -mcode-object-version=4 -o - %s | FileCheck -check-prefix=V4 %s

‎clang/test/CodeGenCXX/dynamic-cast-address-space.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -13,7 +13,7 @@ B fail;
1313
// CHECK: @_ZTI1B = linkonce_odr addrspace(1) constant { ptr addrspace(1), ptr addrspace(1), ptr addrspace(1) } { ptr addrspace(1) getelementptr inbounds (ptr addrspace(1), ptr addrspace(1) @_ZTVN10__cxxabiv120__si_class_type_infoE, i64 2), ptr addrspace(1) @_ZTS1B, ptr addrspace(1) @_ZTI1A }, comdat, align 8
1414
// CHECK: @_ZTVN10__cxxabiv120__si_class_type_infoE = external addrspace(1) global [0 x ptr addrspace(1)]
1515
// CHECK: @_ZTS1B = linkonce_odr addrspace(1) constant [3 x i8] c"1B\00", comdat, align 1
16-
// CHECK: @__oclc_ABI_version = weak_odr hidden local_unnamed_addr addrspace(4) constant i32 500
16+
// CHECK: @__oclc_ABI_version = weak_odr hidden local_unnamed_addr addrspace(4) constant i32 600
1717
//.
1818
// WITH-NONZERO-DEFAULT-AS: @_ZTV1B = linkonce_odr unnamed_addr addrspace(1) constant { [3 x ptr addrspace(1)] } { [3 x ptr addrspace(1)] [ptr addrspace(1) null, ptr addrspace(1) @_ZTI1B, ptr addrspace(1) addrspacecast (ptr addrspace(4) @_ZN1A1fEv to ptr addrspace(1))] }, comdat, align 8
1919
// WITH-NONZERO-DEFAULT-AS: @fail = addrspace(1) global { ptr addrspace(1) } { ptr addrspace(1) getelementptr inbounds inrange(-16, 8) ({ [3 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTV1B, i32 0, i32 0, i32 2) }, align 8
@@ -118,11 +118,11 @@ const B& f(A *a) {
118118
// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR3]] = { nounwind }
119119
// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR4]] = { noreturn }
120120
//.
121-
// CHECK: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 500}
121+
// CHECK: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 600}
122122
// CHECK: [[META1:![0-9]+]] = !{i32 1, !"wchar_size", i32 4}
123123
// CHECK: [[META2:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"}
124124
//.
125-
// WITH-NONZERO-DEFAULT-AS: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 500}
125+
// WITH-NONZERO-DEFAULT-AS: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 600}
126126
// WITH-NONZERO-DEFAULT-AS: [[META1:![0-9]+]] = !{i32 1, !"wchar_size", i32 4}
127127
// WITH-NONZERO-DEFAULT-AS: [[META2:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"}
128128
//.

‎clang/test/CodeGenHIP/default-attributes.hip

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -8,7 +8,7 @@
88
//.
99
// OPTNONE: @__hip_cuid_ = addrspace(1) global i8 0
1010
// OPTNONE: @llvm.compiler.used = appending addrspace(1) global [1 x ptr] [ptr addrspacecast (ptr addrspace(1) @__hip_cuid_ to ptr)], section "llvm.metadata"
11-
// OPTNONE: @__oclc_ABI_version = weak_odr hidden local_unnamed_addr addrspace(4) constant i32 500
11+
// OPTNONE: @__oclc_ABI_version = weak_odr hidden local_unnamed_addr addrspace(4) constant i32 600
1212
//.
1313
__device__ void extern_func();
1414

@@ -39,7 +39,7 @@ __global__ void kernel() {
3939
// OPTNONE: attributes #[[ATTR2]] = { convergent mustprogress noinline norecurse nounwind optnone "amdgpu-flat-work-group-size"="1,1024" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" }
4040
// OPTNONE: attributes #[[ATTR3]] = { convergent nounwind }
4141
//.
42-
// OPTNONE: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 500}
42+
// OPTNONE: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 600}
4343
// OPTNONE: [[META1:![0-9]+]] = !{i32 1, !"amdgpu_printf_kind", !"hostcall"}
4444
// OPTNONE: [[META2:![0-9]+]] = !{i32 1, !"wchar_size", i32 4}
4545
//.

‎clang/test/CodeGenOpenCL/amdgpu-enqueue-kernel.cl

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -68,7 +68,7 @@ kernel void test_target_features_kernel(global int *i) {
6868
// CHECK: @__block_literal_global = internal addrspace(1) constant { i32, i32, ptr } { i32 16, i32 8, ptr @__test_target_features_kernel_block_invoke }, align 8 #0
6969
// CHECK: @__test_target_features_kernel_block_invoke_kernel.runtime.handle = internal addrspace(1) externally_initialized constant %block.runtime.handle.t.3 zeroinitializer, section ".amdgpu.kernel.runtime.handle"
7070
// CHECK: @llvm.used = appending addrspace(1) global [10 x ptr] [ptr @__test_block_invoke_kernel, ptr addrspacecast (ptr addrspace(1) @__test_block_invoke_kernel.runtime.handle to ptr), ptr @__test_block_invoke_2_kernel, ptr addrspacecast (ptr addrspace(1) @__test_block_invoke_2_kernel.runtime.handle to ptr), ptr @__test_block_invoke_3_kernel, ptr addrspacecast (ptr addrspace(1) @__test_block_invoke_3_kernel.runtime.handle to ptr), ptr @__test_block_invoke_4_kernel, ptr addrspacecast (ptr addrspace(1) @__test_block_invoke_4_kernel.runtime.handle to ptr), ptr @__test_target_features_kernel_block_invoke_kernel, ptr addrspacecast (ptr addrspace(1) @__test_target_features_kernel_block_invoke_kernel.runtime.handle to ptr)], section "llvm.metadata"
71-
// CHECK: @__oclc_ABI_version = weak_odr hidden local_unnamed_addr addrspace(4) constant i32 500
71+
// CHECK: @__oclc_ABI_version = weak_odr hidden local_unnamed_addr addrspace(4) constant i32 600
7272
//.
7373
// NOCPU: Function Attrs: convergent noinline norecurse nounwind optnone
7474
// NOCPU-LABEL: define {{[^@]+}}@callee
@@ -764,7 +764,7 @@ kernel void test_target_features_kernel(global int *i) {
764764
// GFX900: attributes #[[ATTR7]] = { nounwind }
765765
// GFX900: attributes #[[ATTR8]] = { convergent nounwind }
766766
//.
767-
// NOCPU: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 500}
767+
// NOCPU: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 600}
768768
// NOCPU: [[META1:![0-9]+]] = !{i32 1, !"wchar_size", i32 4}
769769
// NOCPU: [[META2:![0-9]+]] = !{i32 2, i32 0}
770770
// NOCPU: [[META3]] = !{i32 1, i32 0, i32 1, i32 0}
@@ -787,7 +787,7 @@ kernel void test_target_features_kernel(global int *i) {
787787
// NOCPU: [[META20]] = !{!"int*"}
788788
// NOCPU: [[META21]] = !{ptr addrspace(1) @__test_target_features_kernel_block_invoke_kernel.runtime.handle}
789789
//.
790-
// GFX900: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 500}
790+
// GFX900: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 600}
791791
// GFX900: [[META1:![0-9]+]] = !{i32 1, !"wchar_size", i32 4}
792792
// GFX900: [[META2:![0-9]+]] = !{i32 2, i32 0}
793793
// GFX900: [[TBAA3]] = !{[[META4:![0-9]+]], [[META4]], i64 0}

‎clang/test/Driver/Inputs/rocm-spack/llvm-amdgpu-4.0.0-ieagcs7inf7runpyfvepqkurasoglq4z/amdgcn/bitcode/oclc_abi_version_600.bc

Whitespace-only changes.

‎clang/test/Driver/Inputs/rocm_resource_dir/lib/amdgcn/bitcode/oclc_abi_version_600.bc

Whitespace-only changes.

‎clang/test/Driver/Inputs/rocm_resource_dir/lib64/amdgcn/bitcode/oclc_abi_version_600.bc

Whitespace-only changes.

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -157,7 +157,7 @@
157157
// Test default code object version.
158158
// RUN: %clang -### --target=x86_64-linux-gnu --offload-arch=gfx900 \
159159
// RUN: --rocm-path=%S/Inputs/rocm %S/Inputs/hip_multiple_inputs/b.hip \
160-
// RUN: 2>&1 | FileCheck %s --check-prefixes=ABI5
160+
// RUN: 2>&1 | FileCheck %s --check-prefixes=ABI6
161161

162162
// Test default code object version with old device library without abi_version_400.bc
163163
// RUN: %clang -### --target=x86_64-linux-gnu --offload-arch=gfx900 \

‎clang/test/OpenMP/amdgcn_target_global_constructor.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -29,7 +29,7 @@ S A;
2929
// CHECK: @A = addrspace(1) global %struct.S zeroinitializer, align 4
3030
// CHECK: @llvm.global_ctors = appending global [1 x { i32, ptr, ptr }] [{ i32, ptr, ptr } { i32 65535, ptr @_GLOBAL__sub_I_amdgcn_target_global_constructor.cpp, ptr null }]
3131
// CHECK: @llvm.global_dtors = appending global [1 x { i32, ptr, ptr }] [{ i32, ptr, ptr } { i32 65535, ptr @__dtor_A, ptr null }]
32-
// CHECK: @__oclc_ABI_version = weak_odr hidden local_unnamed_addr addrspace(4) constant i32 500
32+
// CHECK: @__oclc_ABI_version = weak_odr hidden local_unnamed_addr addrspace(4) constant i32 600
3333
//.
3434
// CHECK-LABEL: define {{[^@]+}}@__cxx_global_var_init
3535
// CHECK-SAME: () #[[ATTR0:[0-9]+]] {
@@ -104,7 +104,7 @@ S A;
104104
// CHECK: attributes #[[ATTR4]] = { convergent nounwind }
105105
//.
106106
// CHECK: [[META0:![0-9]+]] = !{i32 1, !"A", i32 0, i32 0}
107-
// CHECK: [[META1:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 500}
107+
// CHECK: [[META1:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 600}
108108
// CHECK: [[META2:![0-9]+]] = !{i32 1, !"wchar_size", i32 4}
109109
// CHECK: [[META3:![0-9]+]] = !{i32 7, !"openmp", i32 51}
110110
// CHECK: [[META4:![0-9]+]] = !{i32 7, !"openmp-device", i32 51}

‎libc/cmake/modules/prepare_libc_gpu_build.cmake

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -104,7 +104,7 @@ if(LIBC_TARGET_ARCHITECTURE_IS_AMDGPU)
104104
# The AMDGPU environment uses different code objects to encode the ABI for
105105
# kernel calls and intrinsic functions. We want to specify this manually to
106106
# conform to whatever the test suite was built to handle.
107-
set(LIBC_GPU_CODE_OBJECT_VERSION 5)
107+
set(LIBC_GPU_CODE_OBJECT_VERSION 6)
108108
endif()
109109

110110
if(LIBC_TARGET_ARCHITECTURE_IS_NVPTX)

‎llvm/docs/ReleaseNotes.md

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -90,6 +90,8 @@ Changes to the AMDGPU Backend
9090
[FWD_PROGRESS bit](https://llvm.org/docs/AMDGPUUsage.html#code-object-v3-kernel-descriptor)
9191
for all GFX ISAs greater or equal to 10, for the AMDHSA OS.
9292

93+
* Bump the default `.amdhsa_code_object_version` to 6. ROCm 6.3 is required to run any program compiled with COV6.
94+
9395
Changes to the ARM Backend
9496
--------------------------
9597

‎llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -34,7 +34,7 @@
3434

3535
static llvm::cl::opt<unsigned> DefaultAMDHSACodeObjectVersion(
3636
"amdhsa-code-object-version", llvm::cl::Hidden,
37-
llvm::cl::init(llvm::AMDGPU::AMDHSA_COV5),
37+
llvm::cl::init(llvm::AMDGPU::AMDHSA_COV6),
3838
llvm::cl::desc("Set default AMDHSA Code Object Version (module flag "
3939
"or asm directive still take priority if present)"));
4040

‎llvm/test/CodeGen/AMDGPU/default_amdhsa_code_object_version.ll

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,7 @@
1+
; RUN: llc -mtriple=amdgcn-amd-amdhsa %s -o - | FileCheck %s
2+
3+
; CHECK: .amdhsa_code_object_version 6
4+
5+
define amdgpu_kernel void @kernel() {
6+
ret void
7+
}

0 commit comments

Comments
 (0)