[AMDGPU] Rename COV module flag to amdhsa_code_object_version (#79905)

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.

Author: epilk

clang/lib/CodeGen/CodeGenModule.cpp

@@ -872,12 +872,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);
     }
 

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

@@ -52,7 +52,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
@@ -82,7 +82,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
 
 // LINKED6: __oclc_ABI_version = weak_odr hidden local_unnamed_addr addrspace(4) constant i32 600
 // LINKED6-LABEL: bar
@@ -112,7 +112,7 @@
 // LINKED6: [[GEP_4_Z:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 8
 // LINKED6: select i1 true, ptr addrspace(4) [[GEP_5_Z]], ptr addrspace(4) [[GEP_4_Z]]
 // LINKED6: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
-// LINKED6: "amdgpu_code_object_version", i32 600
+// LINKED6: "amdhsa_code_object_version", i32 600
 
 #ifdef DEVICELIB
 __device__ void bar(int *x, int *y, int *z)

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

@@ -18,8 +18,8 @@
 // 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}
-// V6: !{{.*}} = !{i32 1, !"amdgpu_code_object_version", i32 600}
-// 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}
+// V6: !{{.*}} = !{i32 1, !"amdhsa_code_object_version", i32 600}
+// NONE-NOT: !{{.*}} = !{i32 1, !"amdhsa_code_object_version",
 // INV: error: invalid value '4.1' in '-mcode-object-version=4.1'

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

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}

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

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

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

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}

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}

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}

llvm/test/CodeGen/AMDGPU/GlobalISel/implicit-kernarg-backend-usage-global-isel.ll

@@ -370,4 +370,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}

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}

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}

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}

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}

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}

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}

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}

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}

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}

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}

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}

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}

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}

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}

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}

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}

llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.kernarg.segment.ptr.ll

@@ -128,4 +128,4 @@ attributes #2 = { nounwind "amdgpu-implicitarg-num-bytes"="48" }
 attributes #3 = { nounwind "amdgpu-implicitarg-num-bytes"="38" }
 
 !llvm.module.flags = !{!0}
-!0 = !{i32 1, !"amdgpu_code_object_version", i32 400}
+!0 = !{i32 1, !"amdhsa_code_object_version", i32 400}

llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.queue.ptr.ll

@@ -17,4 +17,4 @@ declare noalias ptr addrspace(4) @llvm.amdgcn.queue.ptr() #0
 attributes #0 = { nounwind readnone }
 
 !llvm.module.flags = !{!0}
-!0 = !{i32 1, !"amdgpu_code_object_version", i32 400}
+!0 = !{i32 1, !"amdhsa_code_object_version", i32 400}

llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.workgroup.id.ll

@@ -104,4 +104,4 @@ attributes #0 = { nounwind readnone }
 attributes #1 = { nounwind }
 
 !llvm.module.flags = !{!0}
-!0 = !{i32 1, !"amdgpu_code_object_version", i32 400}
+!0 = !{i32 1, !"amdhsa_code_object_version", i32 400}

llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.workitem.id.ll

@@ -201,4 +201,4 @@ attributes #1 = { nounwind }
 !2 = !{i32 1, i32 1, i32 64}
 
 !llvm.module.flags = !{!99}
-!99 = !{i32 1, !"amdgpu_code_object_version", i32 CODE_OBJECT_VERSION}
+!99 = !{i32 1, !"amdhsa_code_object_version", i32 CODE_OBJECT_VERSION}

llvm/test/CodeGen/AMDGPU/GlobalISel/non-entry-alloca.ll

@@ -267,7 +267,7 @@ declare i32 @llvm.amdgcn.workitem.id.x() #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}
 ;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:
 ; ASSUME1024: {{.*}}
 ; DEFAULTSIZE: {{.*}}

llvm/test/CodeGen/AMDGPU/abi-attribute-hints-undefined-behavior.ll

@@ -402,4 +402,4 @@ declare void @llvm.debugtrap()
 attributes #0 = { "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-queue-ptr" "amdgpu-no-work-group-id-x" "amdgpu-no-work-group-id-y" "amdgpu-no-work-group-id-z" "amdgpu-no-work-item-id-x" "amdgpu-no-work-item-id-y" "amdgpu-no-work-item-id-z" }
 
 !llvm.module.flags = !{!0}
-!0 = !{i32 1, !"amdgpu_code_object_version", i32 500}
+!0 = !{i32 1, !"amdhsa_code_object_version", i32 500}

llvm/test/CodeGen/AMDGPU/addrspacecast-constantexpr.ll

@@ -227,7 +227,7 @@ attributes #0 = { argmemonly nounwind }
 attributes #1 = { nounwind }
 
 !llvm.module.flags = !{!0}
-!0 = !{i32 1, !"amdgpu_code_object_version", i32 500}
+!0 = !{i32 1, !"amdhsa_code_object_version", i32 500}
 ;.
 ; AKF_HSA: attributes #[[ATTR0:[0-9]+]] = { nocallback nofree nounwind willreturn memory(argmem: readwrite) }
 ; AKF_HSA: attributes #[[ATTR1]] = { nounwind }
@@ -237,7 +237,7 @@ attributes #1 = { nounwind }
 ; ATTRIBUTOR_HSA: attributes #[[ATTR2]] = { nounwind "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-heap-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-multigrid-sync-arg" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
 ; ATTRIBUTOR_HSA: attributes #[[ATTR3]] = { nounwind "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-heap-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-multigrid-sync-arg" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "amdgpu-waves-per-eu"="4,10" "uniform-work-group-size"="false" }
 ;.
-; AKF_HSA: [[META0:![0-9]+]] = !{i32 1, !"amdgpu_code_object_version", i32 500}
+; AKF_HSA: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 500}
 ;.
-; ATTRIBUTOR_HSA: [[META0:![0-9]+]] = !{i32 1, !"amdgpu_code_object_version", i32 500}
+; ATTRIBUTOR_HSA: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 500}
 ;.

llvm/test/CodeGen/AMDGPU/addrspacecast.gfx6.ll

@@ -208,4 +208,4 @@ define ptr addrspace(6) @addrspacecast_flat_null_to_constant32bit() {
 attributes #0 = { "amdgpu-32bit-address-high-bits"="0xffff8000" }
 
 !llvm.module.flags = !{!0}
-!0 = !{i32 1, !"amdgpu_code_object_version", i32 500}
+!0 = !{i32 1, !"amdhsa_code_object_version", i32 500}

llvm/test/CodeGen/AMDGPU/addrspacecast.ll

@@ -424,4 +424,4 @@ attributes #2 = { nounwind readnone }
 attributes #3 = { nounwind "amdgpu-32bit-address-high-bits"="0xffff8000" }
 
 !llvm.module.flags = !{!0}
-!0 = !{i32 1, !"amdgpu_code_object_version", i32 400}
+!0 = !{i32 1, !"amdhsa_code_object_version", i32 400}

llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-pow-codegen.ll

@@ -775,4 +775,4 @@ define double @test_pown_fast_f64_known_odd(double %x, i32 %y.arg) {
 }
 
 !llvm.module.flags = !{!0}
-!0 = !{i32 1, !"amdgpu_code_object_version", i32 500}
+!0 = !{i32 1, !"amdhsa_code_object_version", i32 500}

llvm/test/CodeGen/AMDGPU/amdgpu.private-memory.ll

@@ -530,7 +530,7 @@ attributes #0 = { nounwind "amdgpu-waves-per-eu"="1,2" "amdgpu-flat-work-group-s
 attributes #1 = { nounwind "amdgpu-flat-work-group-size"="1,256" }
 
 !llvm.module.flags = !{!99}
-!99 = !{i32 1, !"amdgpu_code_object_version", i32 400}
+!99 = !{i32 1, !"amdhsa_code_object_version", i32 400}
 
 ; HSAOPT: !1 = !{}
 ; HSAOPT: !2 = !{i32 0, i32 257}

llvm/test/CodeGen/AMDGPU/annotate-kernel-features-hsa-call.ll

@@ -1012,7 +1012,7 @@ attributes #6 = { "enqueued-block" }
 
 
 !llvm.module.flags = !{!0}
-!0 = !{i32 1, !"amdgpu_code_object_version", i32 500}
+!0 = !{i32 1, !"amdhsa_code_object_version", i32 500}
 ;.
 ; AKF_HSA: attributes #[[ATTR0:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
 ; AKF_HSA: attributes #[[ATTR1]] = { nounwind "target-cpu"="fiji" }
@@ -1055,7 +1055,7 @@ attributes #6 = { "enqueued-block" }
 ; ATTRIBUTOR_HSA: attributes #[[ATTR28]] = { nounwind }
 ; ATTRIBUTOR_HSA: attributes #[[ATTR29]] = { "enqueued-block" }
 ;.
-; AKF_HSA: [[META0:![0-9]+]] = !{i32 1, !"amdgpu_code_object_version", i32 500}
+; AKF_HSA: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 500}
 ;.
-; ATTRIBUTOR_HSA: [[META0:![0-9]+]] = !{i32 1, !"amdgpu_code_object_version", i32 500}
+; ATTRIBUTOR_HSA: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 500}
 ;.

llvm/test/CodeGen/AMDGPU/annotate-kernel-features-hsa.ll

@@ -635,7 +635,7 @@ attributes #0 = { nounwind readnone speculatable }
 attributes #1 = { nounwind }
 
 !llvm.module.flags = !{!0}
-!0 = !{i32 1, !"amdgpu_code_object_version", i32 500}
+!0 = !{i32 1, !"amdhsa_code_object_version", i32 500}
 
 ;.
 ; AKF_HSA: attributes #[[ATTR0:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
@@ -657,7 +657,7 @@ attributes #1 = { nounwind }
 ; ATTRIBUTOR_HSA: attributes #[[ATTR12]] = { nounwind "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-heap-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-multigrid-sync-arg" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
 ; ATTRIBUTOR_HSA: attributes #[[ATTR13]] = { nounwind "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-heap-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-multigrid-sync-arg" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "amdgpu-waves-per-eu"="4,10" "uniform-work-group-size"="false" }
 ;.
-; AKF_HSA: [[META0:![0-9]+]] = !{i32 1, !"amdgpu_code_object_version", i32 500}
+; AKF_HSA: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 500}
 ;.
-; ATTRIBUTOR_HSA: [[META0:![0-9]+]] = !{i32 1, !"amdgpu_code_object_version", i32 500}
+; ATTRIBUTOR_HSA: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 500}
 ;.

llvm/test/CodeGen/AMDGPU/attr-amdgpu-flat-work-group-size.ll

@@ -130,7 +130,7 @@ define amdgpu_kernel void @min_1024_max_1024() #3 {
 attributes #3 = {"amdgpu-flat-work-group-size"="1024,1024"}
 
 !llvm.module.flags = !{!0}
-!0 = !{i32 1, !"amdgpu_code_object_version", i32 400}
+!0 = !{i32 1, !"amdhsa_code_object_version", i32 400}
 
 ; HSAMD: amdhsa.kernels
 ; HSAMD:  .max_flat_workgroup_size: 64

llvm/test/CodeGen/AMDGPU/attributor-noopt.ll

@@ -36,4 +36,4 @@ define amdgpu_kernel void @foo() {
 }
 
 !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}

llvm/test/CodeGen/AMDGPU/blender-no-live-segment-at-def-implicit-def.ll

@@ -124,4 +124,4 @@ kernel_direct_lighting.exit:                      ; preds = %if.end294.i.i, %ent
 declare float @_Z3dotDv3_fS_(<3 x float>)
 
 !llvm.module.flags = !{!0}
-!0 = !{i32 1, !"amdgpu_code_object_version", i32 500}
+!0 = !{i32 1, !"amdhsa_code_object_version", i32 500}