[clang] [llvm] [lld] [AMDGPU] Rename COV module flag to amdhsa_code_object_version (PR #79905)
via cfe-commits
cfe-commits at lists.llvm.org
Mon Jan 29 13:56:58 PST 2024
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clang-codegen
@llvm/pr-subscribers-lld
@llvm/pr-subscribers-clang
Author: Emma Pilkington (epilk)
<details>
<summary>Changes</summary>
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: https://github.com/llvm/llvm-project/pull/76267#discussion_r1445578615
---
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:
- (modified) clang/lib/CodeGen/CodeGenModule.cpp (+2-2)
- (modified) clang/test/CodeGenCUDA/amdgpu-code-object-version-linking.cu (+2-2)
- (modified) clang/test/CodeGenCUDA/amdgpu-code-object-version.cu (+3-3)
- (modified) clang/test/CodeGenHIP/default-attributes.hip (+2-2)
- (modified) clang/test/CodeGenOpenCL/amdgpu-enqueue-kernel.cl (+2-2)
- (modified) lld/test/ELF/lto/amdgcn-oses.ll (+1-1)
- (modified) llvm/lib/IR/AutoUpgrade.cpp (+9)
- (modified) llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp (+1-1)
- (modified) llvm/test/Bitcode/upgrade-module-flag.ll (+4-2)
- (modified) llvm/test/CodeGen/AMDGPU/GlobalISel/crash-stack-address-O0.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/GlobalISel/dropped_debug_info_assert.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/GlobalISel/implicit-kernarg-backend-usage-global-isel.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-assert-align.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-atomicrmw.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-call-abi-attribute-hints.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-call-implicit-args.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-call-return-values.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-call-sret.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-call.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-constant-fold-vector-op.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-function-args.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-indirect-call.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-inline-asm.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-sibling-call.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-tail-call.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.dispatch.id.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.dispatch.ptr.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.is.private.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.is.shared.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.kernarg.segment.ptr.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.queue.ptr.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.workgroup.id.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.workitem.id.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/GlobalISel/non-entry-alloca.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/abi-attribute-hints-undefined-behavior.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/addrspacecast-constantexpr.ll (+3-3)
- (modified) llvm/test/CodeGen/AMDGPU/addrspacecast.gfx6.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/addrspacecast.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-pow-codegen.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/amdgpu.private-memory.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/annotate-kernel-features-hsa-call.ll (+3-3)
- (modified) llvm/test/CodeGen/AMDGPU/annotate-kernel-features-hsa.ll (+3-3)
- (modified) llvm/test/CodeGen/AMDGPU/attr-amdgpu-flat-work-group-size.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/attributor-noopt.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/blender-no-live-segment-at-def-implicit-def.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/branch-folding-implicit-def-subreg.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/call-alias-register-usage-agpr.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/call-alias-register-usage0.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/call-alias-register-usage1.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/call-alias-register-usage2.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/call-alias-register-usage3.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/call-args-inreg.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/call-argument-types.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/call-graph-register-usage.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/call-waitcnt.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/callee-special-input-sgprs-fixed-abi.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/callee-special-input-vgprs.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/cc-update.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/cf-loop-on-constant.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/codegen-internal-only-func.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/collapse-endcf.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/control-flow-fastregalloc.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/cross-block-use-is-not-abi-copy.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/cvt_f32_ubyte.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/dagcombine-lshr-and-cmp.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/ds_read2.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/dwarf-multi-register-use-crash.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/elf-notes.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/enable-scratch-only-dynamic-stack.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/flat-for-global-subtarget-feature.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/flat-scratch-init.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/flat-scratch-reg.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/fneg-fabs.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/gfx11-user-sgpr-init16-bug.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/gfx902-without-xnack.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/global_atomics_scan_fadd.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/global_atomics_scan_fmax.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/global_atomics_scan_fmin.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/global_atomics_scan_fsub.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/hsa-default-device.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/hsa-fp-mode.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/hsa-func.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/hsa-metadata-dynlds-func-hidden-args-v5.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/hsa-metadata-dynlds-funcarg-hidden-args-v5.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/hsa-metadata-dynlds-kernarg-hidden-args-v5.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/hsa-metadata-enqueue-kernel-.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/hsa-metadata-from-llvm-ctor-dtor-list.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/hsa-metadata-from-llvm-ir-full.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/hsa-metadata-heap-v5.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/hsa-metadata-hidden-args-v4.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/hsa-metadata-hidden-args-v5.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/hsa-metadata-hostcall-present-asan.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/hsa-metadata-hostcall-v4.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/hsa-metadata-hostcall-v5.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/hsa-metadata-images.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/hsa-metadata-invalid-ocl-version-1.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/hsa-metadata-invalid-ocl-version-3.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/hsa-metadata-kernel-code-props.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/hsa-metadata-multigrid-sync-arg-v5.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/hsa-metadata-queue-ptr-v5.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/hsa-metadata-queueptr-v5.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/hsa-metadata-resource-usage-function-ordering.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/hsa-metadata-uniform-workgroup-size-v5.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/hsa-metadata-workgroup-processor-mode-v5.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/hsa-note-no-func.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/hsa.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/implicit-arg-v5-opt.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/implicit-kernarg-backend-usage.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/implicit-kernel-argument-alignment.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/implicitarg-offset-attributes.ll (+3-3)
- (modified) llvm/test/CodeGen/AMDGPU/indirect-addressing-term.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/indirect-call.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/insert-delay-alu-bug.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/kernarg-size.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/kernel-argument-dag-lowering.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/kernel-vgpr-spill-mubuf-with-voffset.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/large-alloca-compute.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/lds-alignment.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/lds-frame-extern.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/lds-global-non-entry-func.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/lds-size.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.dispatch.id.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.dispatch.ptr.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.implicitarg.ptr.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.is.private.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.is.shared.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.kernarg.segment.ptr.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.queue.ptr.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.workgroup.id.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.workitem.id.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/llvm.dbg.value.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/lower-kernargs.ll (+3-3)
- (modified) llvm/test/CodeGen/AMDGPU/lower-module-lds-via-hybrid.ll (+2-2)
- (modified) llvm/test/CodeGen/AMDGPU/lower-module-lds-via-table.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/machine-sink-temporal-divergence-swdev407790.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/module-lds-false-sharing.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/need-fp-from-vgpr-spills.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/no-hsa-graphics-shaders.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/non-entry-alloca.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/nop-data.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/partial-sgpr-to-vgpr-spills.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/preserve-wwm-copy-dst-reg.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/promote-alloca-calling-conv.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/promote-alloca-no-opts.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/promote-alloca-padding-size-estimate.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/recursion.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/remove-no-kernel-id-attribute.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/reqd-work-group-size.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/resource-optimization-remarks.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/resource-usage-dead-function.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/sgpr-spill-no-vgprs.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/sgpr-spill-update-only-slot-indexes.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/simple-indirect-call.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/sopk-no-literal.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/spill-m0.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/stack-realign-kernel.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/stacksave_stackrestore.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/tid-kd-xnack-any.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/tid-kd-xnack-off.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/tid-kd-xnack-on.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-all-any.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-all-not-supported.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-all-off.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-all-on.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-any-off-1.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-any-off-2.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-any-on-1.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-any-on-2.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-invalid-any-off-on.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/tid-one-func-xnack-any.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/tid-one-func-xnack-not-supported.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/tid-one-func-xnack-off.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/tid-one-func-xnack-on.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/trap-abis.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/trap.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/tuple-allocation-failure.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/unstructured-cfg-def-use-issue.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/unsupported-code-object-version.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/vgpr-spill-placement-issue61083.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/vgpr_constant_to_sgpr.ll (+1-1)
- (modified) llvm/test/CodeGen/AMDGPU/wwm-reserved.ll (+1-1)
- (modified) llvm/test/tools/llvm-reduce/reduce-module-flags.ll (+1-1)
``````````diff
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]
``````````
</details>
https://github.com/llvm/llvm-project/pull/79905
More information about the cfe-commits
mailing list