[AMDGPU] Rename COV module flag to amdhsa_code_object_version #79905
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
|
@tstellar there was some uncertainty about whether to leave the spelling as-is since mesa used to following HSA code object versions (but doesn't anymore). Would you have any objections to this change? |
llvmbot
added
clang
|
@llvm/pr-subscribers-backend-amdgpu @llvm/pr-subscribers-clang Author: Emma Pilkington (epilk) ChangesThe previous name 'amdgpu_code_object_version', was misleading since this is really a property of the HSA OS. The new spelling also matches the asm directive I added in bc82cfb. See comment thread here: #76267 (comment) Patch is 107.69 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/79905.diff 182 Files Affected:
diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp
index 6ec54cc01c923dc..4fbf63a27de9f66 100644
--- a/clang/lib/CodeGen/CodeGenModule.cpp
+++ b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -873,12 +873,12 @@ void CodeGenModule::Release() {
EmitMainVoidAlias();
if (getTriple().isAMDGPU()) {
- // Emit amdgpu_code_object_version module flag, which is code object version
+ // Emit amdhsa_code_object_version module flag, which is code object version
// times 100.
if (getTarget().getTargetOpts().CodeObjectVersion !=
llvm::CodeObjectVersionKind::COV_None) {
getModule().addModuleFlag(llvm::Module::Error,
- "amdgpu_code_object_version",
+ "amdhsa_code_object_version",
getTarget().getTargetOpts().CodeObjectVersion);
}
diff --git a/clang/test/CodeGenCUDA/amdgpu-code-object-version-linking.cu b/clang/test/CodeGenCUDA/amdgpu-code-object-version-linking.cu
index 663687ae227f234..2f01cf6522946b5 100644
--- a/clang/test/CodeGenCUDA/amdgpu-code-object-version-linking.cu
+++ b/clang/test/CodeGenCUDA/amdgpu-code-object-version-linking.cu
@@ -45,7 +45,7 @@
// LINKED4: [[GEP_4_Z:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 8
// LINKED4: select i1 false, ptr addrspace(4) [[GEP_5_Z]], ptr addrspace(4) [[GEP_4_Z]]
// LINKED4: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
-// LINKED4: "amdgpu_code_object_version", i32 400
+// LINKED4: "amdhsa_code_object_version", i32 400
// LINKED5: __oclc_ABI_version = weak_odr hidden local_unnamed_addr addrspace(4) constant i32 500
// LINKED5-LABEL: bar
@@ -75,7 +75,7 @@
// LINKED5: [[GEP_4_Z:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 8
// LINKED5: select i1 true, ptr addrspace(4) [[GEP_5_Z]], ptr addrspace(4) [[GEP_4_Z]]
// LINKED5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
-// LINKED5: "amdgpu_code_object_version", i32 500
+// LINKED5: "amdhsa_code_object_version", i32 500
#ifdef DEVICELIB
__device__ void bar(int *x, int *y, int *z)
diff --git a/clang/test/CodeGenCUDA/amdgpu-code-object-version.cu b/clang/test/CodeGenCUDA/amdgpu-code-object-version.cu
index 3cb6632fc0b63d3..3cefb10a110baa5 100644
--- a/clang/test/CodeGenCUDA/amdgpu-code-object-version.cu
+++ b/clang/test/CodeGenCUDA/amdgpu-code-object-version.cu
@@ -15,7 +15,7 @@
// RUN: not %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm \
// RUN: -mcode-object-version=4.1 -o - %s 2>&1| FileCheck %s -check-prefix=INV
-// V4: !{{.*}} = !{i32 1, !"amdgpu_code_object_version", i32 400}
-// V5: !{{.*}} = !{i32 1, !"amdgpu_code_object_version", i32 500}
-// NONE-NOT: !{{.*}} = !{i32 1, !"amdgpu_code_object_version",
+// V4: !{{.*}} = !{i32 1, !"amdhsa_code_object_version", i32 400}
+// V5: !{{.*}} = !{i32 1, !"amdhsa_code_object_version", i32 500}
+// NONE-NOT: !{{.*}} = !{i32 1, !"amdhsa_code_object_version",
// INV: error: invalid value '4.1' in '-mcode-object-version=4.1'
diff --git a/clang/test/CodeGenHIP/default-attributes.hip b/clang/test/CodeGenHIP/default-attributes.hip
index 9c9ea521271b99b..63572bfd242b98d 100644
--- a/clang/test/CodeGenHIP/default-attributes.hip
+++ b/clang/test/CodeGenHIP/default-attributes.hip
@@ -46,11 +46,11 @@ __global__ void kernel() {
// OPT: attributes #0 = { mustprogress nofree norecurse nosync nounwind willreturn memory(none) "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
// OPT: attributes #1 = { mustprogress nofree norecurse nosync nounwind willreturn memory(none) "amdgpu-flat-work-group-size"="1,1024" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" }
//.
-// OPTNONE: !0 = !{i32 1, !"amdgpu_code_object_version", i32 500}
+// OPTNONE: !0 = !{i32 1, !"amdhsa_code_object_version", i32 500}
// OPTNONE: !1 = !{i32 1, !"amdgpu_printf_kind", !"hostcall"}
// OPTNONE: !2 = !{i32 1, !"wchar_size", i32 4}
//.
-// OPT: !0 = !{i32 1, !"amdgpu_code_object_version", i32 500}
+// OPT: !0 = !{i32 1, !"amdhsa_code_object_version", i32 500}
// OPT: !1 = !{i32 1, !"amdgpu_printf_kind", !"hostcall"}
// OPT: !2 = !{i32 1, !"wchar_size", i32 4}
//.
diff --git a/clang/test/CodeGenOpenCL/amdgpu-enqueue-kernel.cl b/clang/test/CodeGenOpenCL/amdgpu-enqueue-kernel.cl
index 2cf1286e2b54e8e..a5c9f69bc2de066 100644
--- a/clang/test/CodeGenOpenCL/amdgpu-enqueue-kernel.cl
+++ b/clang/test/CodeGenOpenCL/amdgpu-enqueue-kernel.cl
@@ -703,7 +703,7 @@ kernel void test_target_features_kernel(global int *i) {
// GFX900: attributes #8 = { nounwind }
// GFX900: attributes #9 = { convergent nounwind }
//.
-// NOCPU: !0 = !{i32 1, !"amdgpu_code_object_version", i32 500}
+// NOCPU: !0 = !{i32 1, !"amdhsa_code_object_version", i32 500}
// NOCPU: !1 = !{i32 1, !"wchar_size", i32 4}
// NOCPU: !2 = !{i32 2, i32 0}
// NOCPU: !3 = !{i32 1, i32 0, i32 1, i32 0}
@@ -721,7 +721,7 @@ kernel void test_target_features_kernel(global int *i) {
// NOCPU: !15 = !{i32 1}
// NOCPU: !16 = !{!"int*"}
//.
-// GFX900: !0 = !{i32 1, !"amdgpu_code_object_version", i32 500}
+// GFX900: !0 = !{i32 1, !"amdhsa_code_object_version", i32 500}
// GFX900: !1 = !{i32 1, !"wchar_size", i32 4}
// GFX900: !2 = !{i32 2, i32 0}
// GFX900: !3 = !{!4, !4, i64 0}
diff --git a/lld/test/ELF/lto/amdgcn-oses.ll b/lld/test/ELF/lto/amdgcn-oses.ll
index 0fd0ce4b94775d1..7a74d0317f2b9ea 100644
--- a/lld/test/ELF/lto/amdgcn-oses.ll
+++ b/lld/test/ELF/lto/amdgcn-oses.ll
@@ -28,7 +28,7 @@ target triple = "amdgcn-amd-amdhsa"
target datalayout = "e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-S32-A5"
!llvm.module.flags = !{!0}
-!0 = !{i32 1, !"amdgpu_code_object_version", i32 500}
+!0 = !{i32 1, !"amdhsa_code_object_version", i32 500}
define void @_start() {
ret void
diff --git a/llvm/lib/IR/AutoUpgrade.cpp b/llvm/lib/IR/AutoUpgrade.cpp
index b90bbe71ac1896b..be0abb4b71dae2c 100644
--- a/llvm/lib/IR/AutoUpgrade.cpp
+++ b/llvm/lib/IR/AutoUpgrade.cpp
@@ -5072,6 +5072,15 @@ bool llvm::UpgradeModuleFlags(Module &M) {
Changed = true;
}
}
+
+ if (ID->getString() == "amdgpu_code_object_version") {
+ Metadata *Ops[3] = {
+ Op->getOperand(0),
+ MDString::get(M.getContext(), "amdhsa_code_object_version"),
+ Op->getOperand(2)};
+ ModFlags->setOperand(I, MDNode::get(M.getContext(), Ops));
+ Changed = true;
+ }
}
// "Objective-C Class Properties" is recently added for Objective-C. We
diff --git a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
index 106fdb19f27895b..59470c1d8e0aa04 100644
--- a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
@@ -164,7 +164,7 @@ bool isHsaAbi(const MCSubtargetInfo &STI) {
unsigned getAMDHSACodeObjectVersion(const Module &M) {
if (auto Ver = mdconst::extract_or_null<ConstantInt>(
- M.getModuleFlag("amdgpu_code_object_version"))) {
+ M.getModuleFlag("amdhsa_code_object_version"))) {
return (unsigned)Ver->getZExtValue() / 100;
}
diff --git a/llvm/test/Bitcode/upgrade-module-flag.ll b/llvm/test/Bitcode/upgrade-module-flag.ll
index 1004fd88d18560b..6a523be93fb1d81 100644
--- a/llvm/test/Bitcode/upgrade-module-flag.ll
+++ b/llvm/test/Bitcode/upgrade-module-flag.ll
@@ -1,15 +1,17 @@
; RUN: llvm-as < %s | llvm-dis | FileCheck %s
; RUN: verify-uselistorder < %s
-!llvm.module.flags = !{!0, !1, !2, !3}
+!llvm.module.flags = !{!0, !1, !2, !3, !4}
!0 = !{i32 1, !"PIC Level", i32 1}
!1 = !{i32 1, !"PIE Level", i32 1}
!2 = !{i32 1, !"Objective-C Image Info Version", i32 0}
!3 = !{i32 1, !"Objective-C Image Info Section", !"__DATA, __objc_imageinfo, regular, no_dead_strip"}
+!4 = !{i32 1, !"amdgpu_code_object_version", i32 500}
; CHECK: !0 = !{i32 8, !"PIC Level", i32 1}
; CHECK: !1 = !{i32 7, !"PIE Level", i32 1}
; CHECK: !2 = !{i32 1, !"Objective-C Image Info Version", i32 0}
; CHECK: !3 = !{i32 1, !"Objective-C Image Info Section", !"__DATA,__objc_imageinfo,regular,no_dead_strip"}
-; CHECK: !4 = !{i32 4, !"Objective-C Class Properties", i32 0}
+; CHECK: !4 = !{i32 1, !"amdhsa_code_object_version", i32 500}
+; CHECK: !5 = !{i32 4, !"Objective-C Class Properties", i32 0}
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/crash-stack-address-O0.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/crash-stack-address-O0.ll
index 9580326d7b78faf..2622aaf70cc8d6c 100644
--- a/llvm/test/CodeGen/AMDGPU/GlobalISel/crash-stack-address-O0.ll
+++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/crash-stack-address-O0.ll
@@ -25,4 +25,4 @@ entry:
}
!llvm.module.flags = !{!0}
-!0 = !{i32 1, !"amdgpu_code_object_version", i32 500}
+!0 = !{i32 1, !"amdhsa_code_object_version", i32 500}
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/dropped_debug_info_assert.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/dropped_debug_info_assert.ll
index 44c4910bac7ea1e..56bd7ddde6f5279 100644
--- a/llvm/test/CodeGen/AMDGPU/GlobalISel/dropped_debug_info_assert.ll
+++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/dropped_debug_info_assert.ll
@@ -89,4 +89,4 @@ attributes #0 = { nofree nosync nounwind readnone willreturn }
!7 = distinct !DISubprogram(name: "call_debug_loc", scope: !1, file: !1, line: 8, type: !8, scopeLine: 9, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0, retainedNodes: !9)
!8 = !DISubroutineType(types: !9)
!9 = !{}
-!10 = !{i32 1, !"amdgpu_code_object_version", i32 500}
+!10 = !{i32 1, !"amdhsa_code_object_version", i32 500}
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/implicit-kernarg-backend-usage-global-isel.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/implicit-kernarg-backend-usage-global-isel.ll
index 4bdbe6604782a8b..086f495683b7b80 100644
--- a/llvm/test/CodeGen/AMDGPU/GlobalISel/implicit-kernarg-backend-usage-global-isel.ll
+++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/implicit-kernarg-backend-usage-global-isel.ll
@@ -368,4 +368,4 @@ declare void @llvm.trap()
declare void @llvm.debugtrap()
!llvm.module.flags = !{!0}
-!0 = !{i32 1, !"amdgpu_code_object_version", i32 CODE_OBJECT_VERSION}
+!0 = !{i32 1, !"amdhsa_code_object_version", i32 CODE_OBJECT_VERSION}
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-assert-align.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-assert-align.ll
index 0576d9781e3df46..3150f8cac12846c 100644
--- a/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-assert-align.ll
+++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-assert-align.ll
@@ -211,4 +211,4 @@ entry:
}
!llvm.module.flags = !{!0}
-!0 = !{i32 1, !"amdgpu_code_object_version", i32 500}
+!0 = !{i32 1, !"amdhsa_code_object_version", i32 500}
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-atomicrmw.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-atomicrmw.ll
index 1e1c632ee96fc41..fa49b26847e548c 100644
--- a/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-atomicrmw.ll
+++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-atomicrmw.ll
@@ -50,4 +50,4 @@ define float @test_atomicrmw_fsub(ptr addrspace(3) %addr) {
}
!llvm.module.flags = !{!0}
-!0 = !{i32 1, !"amdgpu_code_object_version", i32 500}
+!0 = !{i32 1, !"amdhsa_code_object_version", i32 500}
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-call-abi-attribute-hints.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-call-abi-attribute-hints.ll
index ca889de30c65000..ca33eae148819f7 100644
--- a/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-call-abi-attribute-hints.ll
+++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-call-abi-attribute-hints.ll
@@ -226,4 +226,4 @@ define void @func_call_no_other_sgprs() {
}
!llvm.module.flags = !{!0}
-!0 = !{i32 1, !"amdgpu_code_object_version", i32 500}
+!0 = !{i32 1, !"amdhsa_code_object_version", i32 500}
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-call-implicit-args.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-call-implicit-args.ll
index 213598b2e812666..a5f59b15c11b846 100644
--- a/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-call-implicit-args.ll
+++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-call-implicit-args.ll
@@ -1228,4 +1228,4 @@ attributes #1 = { nounwind readnone speculatable willreturn }
!3 = !{i32 32, i32 2, i32 1}
!4 = !{i32 1, i32 32, i32 2}
!5 = !{i32 32, i32 1, i32 2}
-!6 = !{i32 1, !"amdgpu_code_object_version", i32 500}
+!6 = !{i32 1, !"amdhsa_code_object_version", i32 500}
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-call-return-values.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-call-return-values.ll
index 8b0a006e29c0002..37f2118572d84e2 100644
--- a/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-call-return-values.ll
+++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-call-return-values.ll
@@ -2969,4 +2969,4 @@ attributes #1 = { nounwind readnone }
attributes #2 = { nounwind noinline }
!llvm.module.flags = !{!0}
-!0 = !{i32 1, !"amdgpu_code_object_version", i32 500}
+!0 = !{i32 1, !"amdhsa_code_object_version", i32 500}
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-call-sret.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-call-sret.ll
index ed68d82997f54bb..854f3463b64d820 100644
--- a/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-call-sret.ll
+++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-call-sret.ll
@@ -90,4 +90,4 @@ define amdgpu_kernel void @test_call_external_void_func_sret_struct_i8_i32_byval
}
!llvm.module.flags = !{!0}
-!0 = !{i32 1, !"amdgpu_code_object_version", i32 500}
+!0 = !{i32 1, !"amdhsa_code_object_version", i32 500}
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-call.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-call.ll
index cb0efc19169dc2a..392b0ae6823e449 100644
--- a/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-call.ll
+++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-call.ll
@@ -6059,4 +6059,4 @@ attributes #1 = { nounwind readnone }
attributes #2 = { nounwind noinline }
!llvm.module.flags = !{!0}
-!0 = !{i32 1, !"amdgpu_code_object_version", i32 500}
+!0 = !{i32 1, !"amdhsa_code_object_version", i32 500}
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-constant-fold-vector-op.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-constant-fold-vector-op.ll
index ab407079abc6677..ce0e2e40e5d19d6 100644
--- a/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-constant-fold-vector-op.ll
+++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-constant-fold-vector-op.ll
@@ -24,4 +24,4 @@ entry:
}
!llvm.module.flags = !{!0}
-!0 = !{i32 1, !"amdgpu_code_object_version", i32 500}
+!0 = !{i32 1, !"amdhsa_code_object_version", i32 500}
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-function-args.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-function-args.ll
index 2f0156d67bdfe58..0b21c2112f05b86 100644
--- a/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-function-args.ll
+++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-function-args.ll
@@ -3236,4 +3236,4 @@ define void @void_func_v2p3_inreg(<2 x ptr addrspace(3)> inreg %arg0) #0 {
attributes #0 = { nounwind }
!llvm.module.flags = !{!0}
-!0 = !{i32 1, !"amdgpu_code_object_version", i32 400}
+!0 = !{i32 1, !"amdhsa_code_object_version", i32 400}
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-indirect-call.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-indirect-call.ll
index 0fb13bf9c8a2b02..0c918def3dc5c03 100644
--- a/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-indirect-call.ll
+++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-indirect-call.ll
@@ -74,4 +74,4 @@ define amdgpu_gfx void @test_gfx_indirect_call_sgpr_ptr(ptr %fptr) {
}
!llvm.module.flags = !{!0}
-!0 = !{i32 1, !"amdgpu_code_object_version", i32 500}
+!0 = !{i32 1, !"amdhsa_code_object_version", i32 500}
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-inline-asm.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-inline-asm.ll
index ceff84ea1812235..326df0750cfbd62 100644
--- a/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-inline-asm.ll
+++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-inline-asm.ll
@@ -333,4 +333,4 @@ define amdgpu_kernel void @asm_constraint_n_n() {
!llvm.module.flags = !{!1}
!0 = !{i32 70}
-!1 = !{i32 1, !"amdgpu_code_object_version", i32 500}
+!1 = !{i32 1, !"amdhsa_code_object_version", i32 500}
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-sibling-call.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-sibling-call.ll
index b45b307f890d0d8..02bf7725015151c 100644
--- a/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-sibling-call.ll
+++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-sibling-call.ll
@@ -1516,4 +1516,4 @@ attributes #0 = { nounwind }
attributes #1 = { nounwind noinline }
!llvm.module.flags = !{!0}
-!0 = !{i32 1, !"amdgpu_code_object_version", i32 500}
+!0 = !{i32 1, !"amdhsa_code_object_version", i32 500}
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-tail-call.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-tail-call.ll
index 1ead1e443dfe13b..81d2f36ac8746d9 100644
--- a/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-tail-call.ll
+++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-tail-call.ll
@@ -44,4 +44,4 @@ define void @tail_call_void_func_void() {
}
!llvm.module.flags = !{!0}
-!0 = !{i32 1, !"amdgpu_code_object_version", i32 500}
+!0 = !{i32 1, !"amdhsa_code_object_version", i32 500}
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.dispatch.id.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.dispatch.id.ll
index fc22d5ec66f0739..0948b03a0b545a6 100644
--- a/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.dispatch.id.ll
+++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.dispatch.id.ll
@@ -17,4 +17,4 @@ attributes #0 = { nounwind }
attributes #1 = { nounwind readnone }
!llvm.module.flags = !{!0}
-!0 = !{i32 1, !"amdgpu_code_object_version", i32 400}
+!0 = !{i32 1, !"amdhsa_code_object_version", i32 400}
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.dispatch.ptr.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.dispatch.ptr.ll
index 26ab4408a5f571a..d165fb577efc2ca 100644
--- a/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.dispatch.ptr.ll
+++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.dispatch.ptr.ll
@@ -17,4 +17,4 @@ declare noalias ptr addrspace(4) @llvm.amdgcn.dispatch.ptr() #0
attributes #0 = { readnone }
!llvm.module.flags = !{!0}
-!0 = !{i32 1, !"amdgpu_code_object_version", i32 400}
+!0 = !{i32 1, !"amdhsa_code_object_version", i32 400}
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.is.private.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.is.private.ll
index a0f54bfee2dca54..303dc46e2c88452 100644
--- a/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.is.private.ll
+++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.is.private.ll
@@ -151,4 +151,4 @@ declare i1 @llvm.amdgcn.is.private(ptr nocapture) #0
attributes #0 = { nounwind readnone speculatable }
!llvm.module.flags = !{!0}
-!0 = !{i32 1, !"amdgpu_code_object_version", i32 500}
+!0 = !{i32 1, !"amdhsa_code_object_version", i32 500}
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.is.shared.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.is.shared.ll
index 4b3a4758910315a..63702d2587574bb 100644
--- a/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.is.shared.ll
+++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.is.shared.ll
@@ -151,4 +151,4 @@ declare i1 @llvm.amdgcn.is.shared(ptr nocapture) #0
attributes #0 = { nounwind readnone speculatable }
!llvm.module.flags = !{!0}
-!0 = !{i32 1, !"amdgpu_code_object_version", i32 500}
+!0 = !{i32 1, !"amdhsa_code_object_version", i32 500}
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.kernarg.segment.ptr.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.kernarg.segment.ptr.ll
index 9d0ede60c7cf172..7fc9842824b01db 100644
--- a/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.kernarg.segment.ptr.ll
+++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.kernarg.segment.ptr.ll
@@ -128,4 +128,4 @@ attributes #2 = { nounwind "amdgpu-implicitarg-num-...
[truncated]
|
arsenm
approved these changes
Jan 30, 2024
|
FYI @DominikAdamski, in case anything needs to change in the flang driver. |
saiislam
reviewed
Jan 30, 2024
| @@ -25,4 +25,4 @@ entry: | |||
| } | |||
|
|
|||
| !llvm.module.flags = !{!0} | |||
| !0 = !{i32 1, !"amdgpu_code_object_version", i32 500} | |||
| !0 = !{i32 1, !"amdhsa_code_object_version", i32 500} | |||
There was a problem hiding this comment.
Can we remove explicit mention of COV5 module flag in all these test files given that the current default is COV5?
Or, may be a separate patch would be better for that.
There was a problem hiding this comment.
Sure, I'll make a PR for that too after this lands.
The previous name 'amdgpu_code_object_version', was misleading since this is really a property of the HSA OS. The new spelling also matches the asm directive I added in bc82cfb.
antiagainst
added a commit
to antiagainst/triton
that referenced
this pull request
Mar 30, 2024
triton-lang#3493 moved the AMD backend to use code object v4. But that reveals some issues that we still need to investigate and resolve. For now switch back to v4. Though we need to be consistent by attaching the LLVM module attribute `amdhsa_code_object_version`, given that's what the AMDGPU LLVM backend is trying to query: https://github.com/llvm/llvm-project/blob/6f44bb77/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp#L166-L173 If missing it will fallback to use v5, which was happening before triton-lang#3493. And FWIW, the magic string used to be `amdgpu_code_object_version`: llvm/llvm-project#79905. The last LLVM bump triton-lang#3463 picked up the renaming patch.
antiagainst
added a commit
to antiagainst/triton
that referenced
this pull request
Mar 30, 2024
triton-lang#3493 moved the AMD backend to use code object v4. But that reveals some issues that we still need to investigate and resolve. For now switch back to v4. Though we need to be consistent by attaching the LLVM module attribute `amdhsa_code_object_version`, given that's what the AMDGPU LLVM backend is trying to query: https://github.com/llvm/llvm-project/blob/6f44bb77/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp#L166-L173 If missing it will fallback to use v5, which was happening before triton-lang#3493. And FWIW, the magic string used to be `amdgpu_code_object_version`: llvm/llvm-project#79905. The last LLVM bump triton-lang#3463 picked up the renaming patch.
ThomasRaoux
pushed a commit
to triton-lang/triton
that referenced
this pull request
Mar 30, 2024
#3493 moved the AMD backend to use code object v5. But that reveals some issues that we still need to investigate and resolve. For now switch back to v4. Though we need to be consistent by attaching the LLVM module attribute `amdhsa_code_object_version`, given that's what the AMDGPU LLVM backend is trying to query: https://github.com/llvm/llvm-project/blob/6f44bb77/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp#L166-L173 If missing it will fallback to use v5, which was happening before #3493. And FWIW, the magic string used to be `amdgpu_code_object_version`: llvm/llvm-project#79905. The last LLVM bump #3463 picked up the renaming patch.
ptillet
pushed a commit
to triton-lang/triton
that referenced
this pull request
Apr 1, 2024
#3493 moved the AMD backend to use code object v5. But that reveals some issues that we still need to investigate and resolve. For now switch back to v4. Though we need to be consistent by attaching the LLVM module attribute `amdhsa_code_object_version`, given that's what the AMDGPU LLVM backend is trying to query: https://github.com/llvm/llvm-project/blob/6f44bb77/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp#L166-L173 If missing it will fallback to use v5, which was happening before #3493. And FWIW, the magic string used to be `amdgpu_code_object_version`: llvm/llvm-project#79905. The last LLVM bump #3463 picked up the renaming patch.
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
The previous name 'amdgpu_code_object_version', was misleading since this is really a property of the HSA OS. The new spelling also matches the asm directive I added in bc82cfb.
See comment thread here: #76267 (comment)