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 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
