shylie/llvm-project
1
0
Fork 0
Code Issues Pull Requests Actions 6 Packages Projects Releases Wiki Activity

[AMDGPU] Set TGID_EN_X/Y/Z when cluster ID intrinsics are used (#159120)

Browse Source 
Hardware initializes a single value in ttmp9 which is either the
workgroup ID X or cluster ID X. Most of this patch is a refactoring to
use a single `PreloadedValue` enumerator for this value, instead of two
enumerators `WORKGROUP_ID_X` and `CLUSTER_ID_X` referring to the same
value.

This makes it simpler to have a single attribute
`amdgpu-no-workgroup-id-x` indicating that this value is not used, which
in turns sets the TGID_EN_X bit appropriately to tell the hardware
whether to initialize it.

All of the above applies to Y and Z similarly.

Fixes: LWPSCGFX13-568

Co-authored-by: Jay Foad <jay.foad@amd.com>
This commit is contained in:
Shilei Tian 2025-09-16 15:37:01 -04:00 committed by GitHub
parent 158eeb344b
commit 8122ccdca9
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
31 changed files with 652 additions and 85 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

22
llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp
View File

@ -809,15 +809,15 @@ bool AMDGPUCallLowering::passSpecialInputs(MachineIRBuilder &MIRBuilder,
AMDGPUFunctionArgInfo::LDS_KERNEL_ID,
};
static constexpr StringLiteral ImplicitAttrNames[] = {
"amdgpu-no-dispatch-ptr",
"amdgpu-no-queue-ptr",
"amdgpu-no-implicitarg-ptr",
"amdgpu-no-dispatch-id",
"amdgpu-no-workgroup-id-x",
"amdgpu-no-workgroup-id-y",
"amdgpu-no-workgroup-id-z",
"amdgpu-no-lds-kernel-id",
static constexpr StringLiteral ImplicitAttrNames[][2] = {
{"amdgpu-no-dispatch-ptr", ""},
{"amdgpu-no-queue-ptr", ""},
{"amdgpu-no-implicitarg-ptr", ""},
{"amdgpu-no-dispatch-id", ""},
{"amdgpu-no-workgroup-id-x", "amdgpu-no-cluster-id-x"},
{"amdgpu-no-workgroup-id-y", "amdgpu-no-cluster-id-y"},
{"amdgpu-no-workgroup-id-z", "amdgpu-no-cluster-id-z"},
{"amdgpu-no-lds-kernel-id", ""},
};
MachineRegisterInfo &MRI = MF.getRegInfo();
@ -833,7 +833,9 @@ bool AMDGPUCallLowering::passSpecialInputs(MachineIRBuilder &MIRBuilder,
LLT ArgTy;
// If the callee does not use the attribute value, skip copying the value.
if (Info.CB->hasFnAttr(ImplicitAttrNames[I++]))
if (all_of(ImplicitAttrNames[I++], [&](StringRef AttrName) {
return AttrName.empty() || Info.CB->hasFnAttr(AttrName);
}))
continue;
std::tie(OutgoingArg, ArgRC, ArgTy) =

26
llvm/lib/Target/AMDGPU/SIISelLowering.cpp
View File

@ -3793,21 +3793,23 @@ void SITargetLowering::passSpecialInputs(
// in the same location as the input.
// clang-format off
static constexpr std::pair<AMDGPUFunctionArgInfo::PreloadedValue,
StringLiteral> ImplicitAttrs[] = {
{AMDGPUFunctionArgInfo::DISPATCH_PTR, "amdgpu-no-dispatch-ptr"},
{AMDGPUFunctionArgInfo::QUEUE_PTR, "amdgpu-no-queue-ptr" },
{AMDGPUFunctionArgInfo::IMPLICIT_ARG_PTR, "amdgpu-no-implicitarg-ptr"},
{AMDGPUFunctionArgInfo::DISPATCH_ID, "amdgpu-no-dispatch-id"},
{AMDGPUFunctionArgInfo::WORKGROUP_ID_X, "amdgpu-no-workgroup-id-x"},
{AMDGPUFunctionArgInfo::WORKGROUP_ID_Y,"amdgpu-no-workgroup-id-y"},
{AMDGPUFunctionArgInfo::WORKGROUP_ID_Z,"amdgpu-no-workgroup-id-z"},
{AMDGPUFunctionArgInfo::LDS_KERNEL_ID,"amdgpu-no-lds-kernel-id"},
};
std::array<StringLiteral, 2>> ImplicitAttrs[] = {
{AMDGPUFunctionArgInfo::DISPATCH_PTR, {"amdgpu-no-dispatch-ptr", ""}},
{AMDGPUFunctionArgInfo::QUEUE_PTR, {"amdgpu-no-queue-ptr", ""}},
{AMDGPUFunctionArgInfo::IMPLICIT_ARG_PTR, {"amdgpu-no-implicitarg-ptr", ""}},
{AMDGPUFunctionArgInfo::DISPATCH_ID, {"amdgpu-no-dispatch-id", ""}},
{AMDGPUFunctionArgInfo::WORKGROUP_ID_X, {"amdgpu-no-workgroup-id-x", "amdgpu-no-cluster-id-x"}},
{AMDGPUFunctionArgInfo::WORKGROUP_ID_Y, {"amdgpu-no-workgroup-id-y", "amdgpu-no-cluster-id-y"}},
{AMDGPUFunctionArgInfo::WORKGROUP_ID_Z, {"amdgpu-no-workgroup-id-z", "amdgpu-no-cluster-id-z"}},
{AMDGPUFunctionArgInfo::LDS_KERNEL_ID, {"amdgpu-no-lds-kernel-id", ""}},
};
// clang-format on
for (auto [InputID, Attr] : ImplicitAttrs) {
for (auto [InputID, Attrs] : ImplicitAttrs) {
// If the callee does not use the attribute value, skip copying the value.
if (CLI.CB->hasFnAttr(Attr))
if (all_of(Attrs, [&](StringRef Attr) {
return Attr.empty() || CLI.CB->hasFnAttr(Attr);
}))
continue;
const auto [OutgoingArg, ArgRC, ArgTy] =

9
llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp
View File

@ -132,13 +132,16 @@ SIMachineFunctionInfo::SIMachineFunctionInfo(const Function &F,
if (!AMDGPU::isGraphics(CC) ||
((CC == CallingConv::AMDGPU_CS || CC == CallingConv::AMDGPU_Gfx) &&
ST.hasArchitectedSGPRs())) {
if (IsKernel || !F.hasFnAttribute("amdgpu-no-workgroup-id-x"))
if (IsKernel || !F.hasFnAttribute("amdgpu-no-workgroup-id-x") ||
!F.hasFnAttribute("amdgpu-no-cluster-id-x"))
WorkGroupIDX = true;
if (!F.hasFnAttribute("amdgpu-no-workgroup-id-y"))
if (!F.hasFnAttribute("amdgpu-no-workgroup-id-y") ||
!F.hasFnAttribute("amdgpu-no-cluster-id-y"))
WorkGroupIDY = true;
if (!F.hasFnAttribute("amdgpu-no-workgroup-id-z"))
if (!F.hasFnAttribute("amdgpu-no-workgroup-id-z") ||
!F.hasFnAttribute("amdgpu-no-cluster-id-z"))
WorkGroupIDZ = true;
}

2
llvm/test/CodeGen/AMDGPU/GlobalISel/call-outgoing-stack-args.ll
View File

@ -430,5 +430,5 @@ define void @func_caller_byval(ptr addrspace(5) %argptr) {
declare void @llvm.memset.p5.i32(ptr addrspace(5) nocapture writeonly, i8, i32, i1 immarg) #1
attributes #0 = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "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" }
attributes #0 = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }
attributes #1 = { argmemonly nofree nounwind willreturn writeonly }

2
llvm/test/CodeGen/AMDGPU/GlobalISel/dereferenceable-declaration.ll
View File

@ -232,4 +232,4 @@ define i64 @load_deref_or_null_maxmimum_callsite_declaration_only() {
ret i64 %add
}
attributes #0 = { "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" }
attributes #0 = { "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-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }

8
llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-call-abi-attribute-hints.ll
View File

@ -88,7 +88,7 @@ define amdgpu_kernel void @kernel_call_no_workgroup_ids() {
; CHECK-NEXT: $sgpr30_sgpr31 = noconvergent G_SI_CALL [[GV]](p0), @extern, csr_amdgpu, implicit $sgpr0_sgpr1_sgpr2_sgpr3, implicit $sgpr4_sgpr5, implicit $sgpr6_sgpr7, implicit $sgpr8_sgpr9, implicit $sgpr10_sgpr11, implicit $sgpr15, implicit $vgpr31
; CHECK-NEXT: ADJCALLSTACKDOWN 0, 0, implicit-def $scc
; CHECK-NEXT: S_ENDPGM 0
call void @extern() "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z"
call void @extern() "amdgpu-no-workgroup-id-x" "amdgpu-no-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z"
ret void
}
@ -124,7 +124,7 @@ define amdgpu_kernel void @kernel_call_no_other_sgprs() {
; CHECK-NEXT: $sgpr30_sgpr31 = noconvergent G_SI_CALL [[GV]](p0), @extern, csr_amdgpu, implicit $sgpr0_sgpr1_sgpr2_sgpr3, implicit $sgpr8_sgpr9, implicit $sgpr15, implicit $vgpr31
; CHECK-NEXT: ADJCALLSTACKDOWN 0, 0, implicit-def $scc
; CHECK-NEXT: S_ENDPGM 0
call void @extern() "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z"
call void @extern() "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z"
ret void
}
@ -198,7 +198,7 @@ define void @func_call_no_workgroup_ids() {
; CHECK-NEXT: $sgpr30_sgpr31 = noconvergent G_SI_CALL [[GV]](p0), @extern, csr_amdgpu, implicit $sgpr0_sgpr1_sgpr2_sgpr3, implicit $sgpr4_sgpr5, implicit $sgpr6_sgpr7, implicit $sgpr8_sgpr9, implicit $sgpr10_sgpr11, implicit $sgpr15, implicit $vgpr31
; CHECK-NEXT: ADJCALLSTACKDOWN 0, 0, implicit-def $scc
; CHECK-NEXT: SI_RETURN
call void @extern() "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z"
call void @extern() "amdgpu-no-workgroup-id-x" "amdgpu-no-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z"
ret void
}
@ -223,7 +223,7 @@ define void @func_call_no_other_sgprs() {
; CHECK-NEXT: $sgpr30_sgpr31 = noconvergent G_SI_CALL [[GV]](p0), @extern, csr_amdgpu, implicit $sgpr0_sgpr1_sgpr2_sgpr3, implicit $sgpr8_sgpr9, implicit $sgpr15, implicit $vgpr31
; CHECK-NEXT: ADJCALLSTACKDOWN 0, 0, implicit-def $scc
; CHECK-NEXT: SI_RETURN
call void @extern() "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z"
call void @extern() "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z"
ret void
}

2
llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-sibling-call.ll
View File

@ -1486,5 +1486,5 @@ entry:
}
attributes #0 = { nounwind }
attributes #1 = { nounwind noinline "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" }
attributes #1 = { nounwind noinline "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-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }

2
llvm/test/CodeGen/AMDGPU/amdhsa-kernarg-preload-num-sgprs.ll
View File

@ -70,4 +70,4 @@ define amdgpu_kernel void @amdhsa_kernarg_preload_1_implicit_2(i32 inreg) #0 { r
define amdgpu_kernel void @amdhsa_kernarg_preload_0_implicit_2(i32) #0 { ret void }
attributes #0 = { "amdgpu-agpr-alloc"="0" "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" }
attributes #0 = { "amdgpu-agpr-alloc"="0" "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-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }

2
llvm/test/CodeGen/AMDGPU/call-argument-types.ll
View File

@ -7283,7 +7283,7 @@ declare hidden void @external_void_func_12xv3f32(<3 x float>, <3 x float>, <3 x
declare hidden void @external_void_func_8xv5f32(<5 x float>, <5 x float>, <5 x float>, <5 x float>,
<5 x float>, <5 x float>, <5 x float>, <5 x float>) #0
attributes #0 = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-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" }
attributes #0 = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }
attributes #1 = { nounwind readnone }
attributes #2 = { nounwind noinline }

2
llvm/test/CodeGen/AMDGPU/call-defs-mode-register.ll
View File

@ -54,4 +54,4 @@ define float @asm_changes_mode(float %x, float %y) #0 {
declare float @llvm.experimental.constrained.fadd.f32(float, float, metadata, metadata)
attributes #0 = { strictfp "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" }
attributes #0 = { strictfp "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-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }

2
llvm/test/CodeGen/AMDGPU/call-preserved-registers.ll
View File

@ -747,4 +747,4 @@ define amdgpu_kernel void @callee_saved_sgpr_vgpr_kernel() #2 {
attributes #0 = { nounwind }
attributes #1 = { nounwind readnone }
attributes #2 = { nounwind noinline }
attributes #3 = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-lds-kernel-id" "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" }
attributes #3 = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }

2
llvm/test/CodeGen/AMDGPU/call-reqd-group-size.ll
View File

@ -208,7 +208,7 @@ define amdgpu_kernel void @known_xyz_0(ptr addrspace(1) %out) !reqd_work_group_s
}
; CHECK: .amdhsa_system_vgpr_workitem_id 0
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-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" }
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-workgroup-id-x" "amdgpu-no-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z" }
!0 = !{i32 1, i32 64, i32 64}
!1 = !{i32 64, i32 1, i32 64}

2
llvm/test/CodeGen/AMDGPU/call-waitcnt.ll
View File

@ -154,5 +154,5 @@ declare hidden void @func(i32) #0
declare hidden i32 @func.return(i32) #0
declare void @got.func(i32) #0
attributes #0 = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-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" }
attributes #0 = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }

2
llvm/test/CodeGen/AMDGPU/elf-notes.ll
View File

@ -85,7 +85,7 @@ define amdgpu_kernel void @elf_notes() #0 {
ret void
}
attributes #0 = { "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-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" }
attributes #0 = { "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-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-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }
!llvm.module.flags = !{!0}
!0 = !{i32 1, !"amdhsa_code_object_version", i32 400}

6
llvm/test/CodeGen/AMDGPU/gfx11-user-sgpr-init16-bug.ll
View File

@ -200,8 +200,8 @@ declare align 4 ptr addrspace(4) @llvm.amdgcn.queue.ptr() #3
declare align 4 ptr addrspace(4) @llvm.amdgcn.kernarg.segment.ptr() #3
declare i64 @llvm.amdgcn.dispatch.id() #3
attributes #0 = { "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-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" }
attributes #1 = { "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-multigrid-sync-arg" "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" }
attributes #2 = { "amdgpu-no-lds-kernel-id" "amdgpu-no-multigrid-sync-arg" "amdgpu-no-workgroup-id-x" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }
attributes #0 = { "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-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-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }
attributes #1 = { "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-multigrid-sync-arg" "amdgpu-no-workgroup-id-x" "amdgpu-no-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }
attributes #2 = { "amdgpu-no-lds-kernel-id" "amdgpu-no-multigrid-sync-arg" "amdgpu-no-workgroup-id-x" "amdgpu-no-cluster-id-x" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }
attributes #3 = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }

4
llvm/test/CodeGen/AMDGPU/hsa.ll
View File

@ -43,7 +43,7 @@
; ELF: 00E0: 6E616D65 A673696D 706C65BB 2E707269
; ELF: 00F0: 76617465 5F736567 6D656E74 5F666978
; ELF: 0100: 65645F73 697A6500 AB2E7367 70725F63
; ELF: 0110: 6F756E74 0EB12E73 6770725F 7370696C
; ELF: 0110: 6F756E74 10B12E73 6770725F 7370696C
; ELF: 0120: 6C5F636F 756E7400 A72E7379 6D626F6C
; ELF: 0130: A973696D 706C652E 6B64AB2E 76677072
; ELF: 0140: 5F636F75 6E7403B1 2E766770 725F7370
@ -59,7 +59,7 @@
; ELF: 01E0: 73696D70 6C655F6E 6F5F6B65 726E6172
; ELF: 01F0: 6773BB2E 70726976 6174655F 7365676D
; ELF: 0200: 656E745F 66697865 645F7369 7A6500AB
; ELF: 0210: 2E736770 725F636F 756E740C B12E7367
; ELF: 0210: 2E736770 725F636F 756E740E B12E7367
; ELF: 0220: 70725F73 70696C6C 5F636F75 6E7400A7
; ELF: 0230: 2E73796D 626F6CB5 73696D70 6C655F6E
; ELF: 0240: 6F5F6B65 726E6172 67732E6B 64AB2E76

554
llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cluster.id.ll Normal file
View File

@ -0,0 +1,554 @@
; RUN: opt -mtriple=amdgcn-amd-amdhsa -passes=amdgpu-attributor %s -o %t.bc
; RUN: llc -mtriple=amdgcn -mcpu=gfx1250 %t.bc -o - | FileCheck --check-prefixes=CHECK-UNKNOWN %s
; RUN: llc -mtriple=amdgcn-unknown-mesa3d -mcpu=gfx1250 %t.bc -o - | FileCheck -check-prefixes=CHECK-MESA3D %s
; RUN: llc -global-isel -mtriple=amdgcn -mcpu=gfx1250 %t.bc -o - | FileCheck --check-prefixes=CHECK-G-UNKNOWN %s
; RUN: llc -global-isel -mtriple=amdgcn-unknown-mesa3d -mcpu=gfx1250 %t.bc -o - | FileCheck -check-prefixes=CHECK-G-MESA3D %s
declare i32 @llvm.amdgcn.cluster.id.x() #0
declare i32 @llvm.amdgcn.cluster.id.y() #0
declare i32 @llvm.amdgcn.cluster.id.z() #0
define amdgpu_kernel void @test_cluster_id_x(ptr addrspace(1) %out) {
; CHECK-UNKNOWN-LABEL: test_cluster_id_x:
; CHECK-UNKNOWN: ; %bb.0:
; CHECK-UNKNOWN-NEXT: s_load_b64 s[0:1], s[0:1], 0x24
; CHECK-UNKNOWN-NEXT: v_dual_mov_b32 v0, ttmp9 :: v_dual_mov_b32 v1, 0
; CHECK-UNKNOWN-NEXT: s_wait_kmcnt 0x0
; CHECK-UNKNOWN-NEXT: global_store_b32 v1, v0, s[0:1]
; CHECK-UNKNOWN-NEXT: s_endpgm
; CHECK-UNKNOWN: COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; CHECK-UNKNOWN: COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; CHECK-UNKNOWN: COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
;
; CHECK-MESA3D-LABEL: test_cluster_id_x:
; CHECK-MESA3D: .amd_kernel_code_t
; CHECK-MESA3D-NEXT: amd_code_version_major = 1
; CHECK-MESA3D-NEXT: amd_code_version_minor = 2
; CHECK-MESA3D-NEXT: amd_machine_kind = 1
; CHECK-MESA3D-NEXT: amd_machine_version_major = 12
; CHECK-MESA3D-NEXT: amd_machine_version_minor = 5
; CHECK-MESA3D-NEXT: amd_machine_version_stepping = 0
; CHECK-MESA3D-NEXT: kernel_code_entry_byte_offset = 256
; CHECK-MESA3D-NEXT: kernel_code_prefetch_byte_size = 0
; CHECK-MESA3D-NEXT: granulated_workitem_vgpr_count = 0
; CHECK-MESA3D-NEXT: granulated_wavefront_sgpr_count = 0
; CHECK-MESA3D-NEXT: priority = 0
; CHECK-MESA3D-NEXT: float_mode = 240
; CHECK-MESA3D-NEXT: priv = 0
; CHECK-MESA3D-NEXT: enable_dx10_clamp = 0
; CHECK-MESA3D-NEXT: debug_mode = 0
; CHECK-MESA3D-NEXT: enable_ieee_mode = 0
; CHECK-MESA3D-NEXT: enable_wgp_mode = 0
; CHECK-MESA3D-NEXT: enable_mem_ordered = 1
; CHECK-MESA3D-NEXT: enable_fwd_progress = 1
; CHECK-MESA3D-NEXT: enable_sgpr_private_segment_wave_byte_offset = 0
; CHECK-MESA3D-NEXT: user_sgpr_count = 2
; CHECK-MESA3D-NEXT: enable_trap_handler = 0
; CHECK-MESA3D-NEXT: enable_sgpr_workgroup_id_x = 1
; CHECK-MESA3D-NEXT: enable_sgpr_workgroup_id_y = 0
; CHECK-MESA3D-NEXT: enable_sgpr_workgroup_id_z = 0
; CHECK-MESA3D-NEXT: enable_sgpr_workgroup_info = 0
; CHECK-MESA3D-NEXT: enable_vgpr_workitem_id = 0
; CHECK-MESA3D-NEXT: enable_exception_msb = 0
; CHECK-MESA3D-NEXT: granulated_lds_size = 0
; CHECK-MESA3D-NEXT: enable_exception = 0
; CHECK-MESA3D-NEXT: enable_sgpr_private_segment_buffer = 0
; CHECK-MESA3D-NEXT: enable_sgpr_dispatch_ptr = 0
; CHECK-MESA3D-NEXT: enable_sgpr_queue_ptr = 0
; CHECK-MESA3D-NEXT: enable_sgpr_kernarg_segment_ptr = 1
; CHECK-MESA3D-NEXT: enable_sgpr_dispatch_id = 0
; CHECK-MESA3D-NEXT: enable_sgpr_flat_scratch_init = 0
; CHECK-MESA3D-NEXT: enable_sgpr_private_segment_size = 0
; CHECK-MESA3D-NEXT: enable_sgpr_grid_workgroup_count_x = 0
; CHECK-MESA3D-NEXT: enable_sgpr_grid_workgroup_count_y = 0
; CHECK-MESA3D-NEXT: enable_sgpr_grid_workgroup_count_z = 0
; CHECK-MESA3D-NEXT: enable_wavefront_size32 = 1
; CHECK-MESA3D-NEXT: enable_ordered_append_gds = 0
; CHECK-MESA3D-NEXT: private_element_size = 1
; CHECK-MESA3D-NEXT: is_ptr64 = 1
; CHECK-MESA3D-NEXT: is_dynamic_callstack = 0
; CHECK-MESA3D-NEXT: is_debug_enabled = 0
; CHECK-MESA3D-NEXT: is_xnack_enabled = 0
; CHECK-MESA3D-NEXT: workitem_private_segment_byte_size = 0
; CHECK-MESA3D-NEXT: workgroup_group_segment_byte_size = 0
; CHECK-MESA3D-NEXT: gds_segment_byte_size = 0
; CHECK-MESA3D-NEXT: kernarg_segment_byte_size = 8
; CHECK-MESA3D-NEXT: workgroup_fbarrier_count = 0
; CHECK-MESA3D-NEXT: wavefront_sgpr_count = 2
; CHECK-MESA3D-NEXT: workitem_vgpr_count = 2
; CHECK-MESA3D-NEXT: reserved_vgpr_first = 0
; CHECK-MESA3D-NEXT: reserved_vgpr_count = 0
; CHECK-MESA3D-NEXT: reserved_sgpr_first = 0
; CHECK-MESA3D-NEXT: reserved_sgpr_count = 0
; CHECK-MESA3D-NEXT: debug_wavefront_private_segment_offset_sgpr = 0
; CHECK-MESA3D-NEXT: debug_private_segment_buffer_sgpr = 0
; CHECK-MESA3D-NEXT: kernarg_segment_alignment = 4
; CHECK-MESA3D-NEXT: group_segment_alignment = 4
; CHECK-MESA3D-NEXT: private_segment_alignment = 4
; CHECK-MESA3D-NEXT: wavefront_size = 5
; CHECK-MESA3D-NEXT: call_convention = -1
; CHECK-MESA3D-NEXT: runtime_loader_kernel_symbol = 0
; CHECK-MESA3D-NEXT: .end_amd_kernel_code_t
; CHECK-MESA3D-NEXT: ; %bb.0:
; CHECK-MESA3D-NEXT: s_load_b64 s[0:1], s[0:1], 0x0
; CHECK-MESA3D-NEXT: v_dual_mov_b32 v0, ttmp9 :: v_dual_mov_b32 v1, 0
; CHECK-MESA3D-NEXT: s_wait_kmcnt 0x0
; CHECK-MESA3D-NEXT: global_store_b32 v1, v0, s[0:1]
; CHECK-MESA3D-NEXT: s_endpgm
;
; CHECK-G-UNKNOWN-LABEL: test_cluster_id_x:
; CHECK-G-UNKNOWN: ; %bb.0:
; CHECK-G-UNKNOWN-NEXT: s_load_b64 s[0:1], s[0:1], 0x24
; CHECK-G-UNKNOWN-NEXT: v_dual_mov_b32 v0, ttmp9 :: v_dual_mov_b32 v1, 0
; CHECK-G-UNKNOWN-NEXT: s_wait_kmcnt 0x0
; CHECK-G-UNKNOWN-NEXT: global_store_b32 v1, v0, s[0:1]
; CHECK-G-UNKNOWN-NEXT: s_endpgm
; CHECK-G-UNKNOWN: COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; CHECK-G-UNKNOWN: COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; CHECK-G-UNKNOWN: COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
;
; CHECK-G-MESA3D-LABEL: test_cluster_id_x:
; CHECK-G-MESA3D: .amd_kernel_code_t
; CHECK-G-MESA3D-NEXT: amd_code_version_major = 1
; CHECK-G-MESA3D-NEXT: amd_code_version_minor = 2
; CHECK-G-MESA3D-NEXT: amd_machine_kind = 1
; CHECK-G-MESA3D-NEXT: amd_machine_version_major = 12
; CHECK-G-MESA3D-NEXT: amd_machine_version_minor = 5
; CHECK-G-MESA3D-NEXT: amd_machine_version_stepping = 0
; CHECK-G-MESA3D-NEXT: kernel_code_entry_byte_offset = 256
; CHECK-G-MESA3D-NEXT: kernel_code_prefetch_byte_size = 0
; CHECK-G-MESA3D-NEXT: granulated_workitem_vgpr_count = 0
; CHECK-G-MESA3D-NEXT: granulated_wavefront_sgpr_count = 0
; CHECK-G-MESA3D-NEXT: priority = 0
; CHECK-G-MESA3D-NEXT: float_mode = 240
; CHECK-G-MESA3D-NEXT: priv = 0
; CHECK-G-MESA3D-NEXT: enable_dx10_clamp = 0
; CHECK-G-MESA3D-NEXT: debug_mode = 0
; CHECK-G-MESA3D-NEXT: enable_ieee_mode = 0
; CHECK-G-MESA3D-NEXT: enable_wgp_mode = 0
; CHECK-G-MESA3D-NEXT: enable_mem_ordered = 1
; CHECK-G-MESA3D-NEXT: enable_fwd_progress = 1
; CHECK-G-MESA3D-NEXT: enable_sgpr_private_segment_wave_byte_offset = 0
; CHECK-G-MESA3D-NEXT: user_sgpr_count = 2
; CHECK-G-MESA3D-NEXT: enable_trap_handler = 0
; CHECK-G-MESA3D-NEXT: enable_sgpr_workgroup_id_x = 1
; CHECK-G-MESA3D-NEXT: enable_sgpr_workgroup_id_y = 0
; CHECK-G-MESA3D-NEXT: enable_sgpr_workgroup_id_z = 0
; CHECK-G-MESA3D-NEXT: enable_sgpr_workgroup_info = 0
; CHECK-G-MESA3D-NEXT: enable_vgpr_workitem_id = 0
; CHECK-G-MESA3D-NEXT: enable_exception_msb = 0
; CHECK-G-MESA3D-NEXT: granulated_lds_size = 0
; CHECK-G-MESA3D-NEXT: enable_exception = 0
; CHECK-G-MESA3D-NEXT: enable_sgpr_private_segment_buffer = 0
; CHECK-G-MESA3D-NEXT: enable_sgpr_dispatch_ptr = 0
; CHECK-G-MESA3D-NEXT: enable_sgpr_queue_ptr = 0
; CHECK-G-MESA3D-NEXT: enable_sgpr_kernarg_segment_ptr = 1
; CHECK-G-MESA3D-NEXT: enable_sgpr_dispatch_id = 0
; CHECK-G-MESA3D-NEXT: enable_sgpr_flat_scratch_init = 0
; CHECK-G-MESA3D-NEXT: enable_sgpr_private_segment_size = 0
; CHECK-G-MESA3D-NEXT: enable_sgpr_grid_workgroup_count_x = 0
; CHECK-G-MESA3D-NEXT: enable_sgpr_grid_workgroup_count_y = 0
; CHECK-G-MESA3D-NEXT: enable_sgpr_grid_workgroup_count_z = 0
; CHECK-G-MESA3D-NEXT: enable_wavefront_size32 = 1
; CHECK-G-MESA3D-NEXT: enable_ordered_append_gds = 0
; CHECK-G-MESA3D-NEXT: private_element_size = 1
; CHECK-G-MESA3D-NEXT: is_ptr64 = 1
; CHECK-G-MESA3D-NEXT: is_dynamic_callstack = 0
; CHECK-G-MESA3D-NEXT: is_debug_enabled = 0
; CHECK-G-MESA3D-NEXT: is_xnack_enabled = 0
; CHECK-G-MESA3D-NEXT: workitem_private_segment_byte_size = 0
; CHECK-G-MESA3D-NEXT: workgroup_group_segment_byte_size = 0
; CHECK-G-MESA3D-NEXT: gds_segment_byte_size = 0
; CHECK-G-MESA3D-NEXT: kernarg_segment_byte_size = 8
; CHECK-G-MESA3D-NEXT: workgroup_fbarrier_count = 0
; CHECK-G-MESA3D-NEXT: wavefront_sgpr_count = 2
; CHECK-G-MESA3D-NEXT: workitem_vgpr_count = 2
; CHECK-G-MESA3D-NEXT: reserved_vgpr_first = 0
; CHECK-G-MESA3D-NEXT: reserved_vgpr_count = 0
; CHECK-G-MESA3D-NEXT: reserved_sgpr_first = 0
; CHECK-G-MESA3D-NEXT: reserved_sgpr_count = 0
; CHECK-G-MESA3D-NEXT: debug_wavefront_private_segment_offset_sgpr = 0
; CHECK-G-MESA3D-NEXT: debug_private_segment_buffer_sgpr = 0
; CHECK-G-MESA3D-NEXT: kernarg_segment_alignment = 4
; CHECK-G-MESA3D-NEXT: group_segment_alignment = 4
; CHECK-G-MESA3D-NEXT: private_segment_alignment = 4
; CHECK-G-MESA3D-NEXT: wavefront_size = 5
; CHECK-G-MESA3D-NEXT: call_convention = -1
; CHECK-G-MESA3D-NEXT: runtime_loader_kernel_symbol = 0
; CHECK-G-MESA3D-NEXT: .end_amd_kernel_code_t
; CHECK-G-MESA3D-NEXT: ; %bb.0:
; CHECK-G-MESA3D-NEXT: s_load_b64 s[0:1], s[0:1], 0x0
; CHECK-G-MESA3D-NEXT: v_dual_mov_b32 v0, ttmp9 :: v_dual_mov_b32 v1, 0
; CHECK-G-MESA3D-NEXT: s_wait_kmcnt 0x0
; CHECK-G-MESA3D-NEXT: global_store_b32 v1, v0, s[0:1]
; CHECK-G-MESA3D-NEXT: s_endpgm
%id = call i32 @llvm.amdgcn.cluster.id.x()
store i32 %id, ptr addrspace(1) %out
ret void
}
define amdgpu_kernel void @test_cluster_id_y(ptr addrspace(1) %out) #1 {
; CHECK-UNKNOWN-LABEL: test_cluster_id_y:
; CHECK-UNKNOWN: ; %bb.0:
; CHECK-UNKNOWN-NEXT: s_load_b64 s[0:1], s[0:1], 0x24
; CHECK-UNKNOWN-NEXT: v_dual_mov_b32 v0, ttmp7 :: v_dual_mov_b32 v1, 0
; CHECK-UNKNOWN-NEXT: s_wait_kmcnt 0x0
; CHECK-UNKNOWN-NEXT: global_store_b32 v1, v0, s[0:1]
; CHECK-UNKNOWN-NEXT: s_endpgm
; CHECK-UNKNOWN: COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; CHECK-UNKNOWN: COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; CHECK-UNKNOWN: COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
;
; CHECK-MESA3D-LABEL: test_cluster_id_y:
; CHECK-MESA3D: .amd_kernel_code_t
; CHECK-MESA3D-NEXT: amd_code_version_major = 1
; CHECK-MESA3D-NEXT: amd_code_version_minor = 2
; CHECK-MESA3D-NEXT: amd_machine_kind = 1
; CHECK-MESA3D-NEXT: amd_machine_version_major = 12
; CHECK-MESA3D-NEXT: amd_machine_version_minor = 5
; CHECK-MESA3D-NEXT: amd_machine_version_stepping = 0
; CHECK-MESA3D-NEXT: kernel_code_entry_byte_offset = 256
; CHECK-MESA3D-NEXT: kernel_code_prefetch_byte_size = 0
; CHECK-MESA3D-NEXT: granulated_workitem_vgpr_count = 0
; CHECK-MESA3D-NEXT: granulated_wavefront_sgpr_count = 0
; CHECK-MESA3D-NEXT: priority = 0
; CHECK-MESA3D-NEXT: float_mode = 240
; CHECK-MESA3D-NEXT: priv = 0
; CHECK-MESA3D-NEXT: enable_dx10_clamp = 0
; CHECK-MESA3D-NEXT: debug_mode = 0
; CHECK-MESA3D-NEXT: enable_ieee_mode = 0
; CHECK-MESA3D-NEXT: enable_wgp_mode = 0
; CHECK-MESA3D-NEXT: enable_mem_ordered = 1
; CHECK-MESA3D-NEXT: enable_fwd_progress = 1
; CHECK-MESA3D-NEXT: enable_sgpr_private_segment_wave_byte_offset = 0
; CHECK-MESA3D-NEXT: user_sgpr_count = 2
; CHECK-MESA3D-NEXT: enable_trap_handler = 0
; CHECK-MESA3D-NEXT: enable_sgpr_workgroup_id_x = 1
; CHECK-MESA3D-NEXT: enable_sgpr_workgroup_id_y = 1
; CHECK-MESA3D-NEXT: enable_sgpr_workgroup_id_z = 0
; CHECK-MESA3D-NEXT: enable_sgpr_workgroup_info = 0
; CHECK-MESA3D-NEXT: enable_vgpr_workitem_id = 0
; CHECK-MESA3D-NEXT: enable_exception_msb = 0
; CHECK-MESA3D-NEXT: granulated_lds_size = 0
; CHECK-MESA3D-NEXT: enable_exception = 0
; CHECK-MESA3D-NEXT: enable_sgpr_private_segment_buffer = 0
; CHECK-MESA3D-NEXT: enable_sgpr_dispatch_ptr = 0
; CHECK-MESA3D-NEXT: enable_sgpr_queue_ptr = 0
; CHECK-MESA3D-NEXT: enable_sgpr_kernarg_segment_ptr = 1
; CHECK-MESA3D-NEXT: enable_sgpr_dispatch_id = 0
; CHECK-MESA3D-NEXT: enable_sgpr_flat_scratch_init = 0
; CHECK-MESA3D-NEXT: enable_sgpr_private_segment_size = 0
; CHECK-MESA3D-NEXT: enable_sgpr_grid_workgroup_count_x = 0
; CHECK-MESA3D-NEXT: enable_sgpr_grid_workgroup_count_y = 0
; CHECK-MESA3D-NEXT: enable_sgpr_grid_workgroup_count_z = 0
; CHECK-MESA3D-NEXT: enable_wavefront_size32 = 1
; CHECK-MESA3D-NEXT: enable_ordered_append_gds = 0
; CHECK-MESA3D-NEXT: private_element_size = 1
; CHECK-MESA3D-NEXT: is_ptr64 = 1
; CHECK-MESA3D-NEXT: is_dynamic_callstack = 0
; CHECK-MESA3D-NEXT: is_debug_enabled = 0
; CHECK-MESA3D-NEXT: is_xnack_enabled = 0
; CHECK-MESA3D-NEXT: workitem_private_segment_byte_size = 0
; CHECK-MESA3D-NEXT: workgroup_group_segment_byte_size = 0
; CHECK-MESA3D-NEXT: gds_segment_byte_size = 0
; CHECK-MESA3D-NEXT: kernarg_segment_byte_size = 8
; CHECK-MESA3D-NEXT: workgroup_fbarrier_count = 0
; CHECK-MESA3D-NEXT: wavefront_sgpr_count = 2
; CHECK-MESA3D-NEXT: workitem_vgpr_count = 2
; CHECK-MESA3D-NEXT: reserved_vgpr_first = 0
; CHECK-MESA3D-NEXT: reserved_vgpr_count = 0
; CHECK-MESA3D-NEXT: reserved_sgpr_first = 0
; CHECK-MESA3D-NEXT: reserved_sgpr_count = 0
; CHECK-MESA3D-NEXT: debug_wavefront_private_segment_offset_sgpr = 0
; CHECK-MESA3D-NEXT: debug_private_segment_buffer_sgpr = 0
; CHECK-MESA3D-NEXT: kernarg_segment_alignment = 4
; CHECK-MESA3D-NEXT: group_segment_alignment = 4
; CHECK-MESA3D-NEXT: private_segment_alignment = 4
; CHECK-MESA3D-NEXT: wavefront_size = 5
; CHECK-MESA3D-NEXT: call_convention = -1
; CHECK-MESA3D-NEXT: runtime_loader_kernel_symbol = 0
; CHECK-MESA3D-NEXT: .end_amd_kernel_code_t
; CHECK-MESA3D-NEXT: ; %bb.0:
; CHECK-MESA3D-NEXT: s_load_b64 s[0:1], s[0:1], 0x0
; CHECK-MESA3D-NEXT: v_dual_mov_b32 v0, ttmp7 :: v_dual_mov_b32 v1, 0
; CHECK-MESA3D-NEXT: s_wait_kmcnt 0x0
; CHECK-MESA3D-NEXT: global_store_b32 v1, v0, s[0:1]
; CHECK-MESA3D-NEXT: s_endpgm
;
; CHECK-G-UNKNOWN-LABEL: test_cluster_id_y:
; CHECK-G-UNKNOWN: ; %bb.0:
; CHECK-G-UNKNOWN-NEXT: s_load_b64 s[0:1], s[0:1], 0x24
; CHECK-G-UNKNOWN-NEXT: v_dual_mov_b32 v0, ttmp7 :: v_dual_mov_b32 v1, 0
; CHECK-G-UNKNOWN-NEXT: s_wait_kmcnt 0x0
; CHECK-G-UNKNOWN-NEXT: global_store_b32 v1, v0, s[0:1]
; CHECK-G-UNKNOWN-NEXT: s_endpgm
; CHECK-G-UNKNOWN: COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; CHECK-G-UNKNOWN: COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; CHECK-G-UNKNOWN: COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
;
; CHECK-G-MESA3D-LABEL: test_cluster_id_y:
; CHECK-G-MESA3D: .amd_kernel_code_t
; CHECK-G-MESA3D-NEXT: amd_code_version_major = 1
; CHECK-G-MESA3D-NEXT: amd_code_version_minor = 2
; CHECK-G-MESA3D-NEXT: amd_machine_kind = 1
; CHECK-G-MESA3D-NEXT: amd_machine_version_major = 12
; CHECK-G-MESA3D-NEXT: amd_machine_version_minor = 5
; CHECK-G-MESA3D-NEXT: amd_machine_version_stepping = 0
; CHECK-G-MESA3D-NEXT: kernel_code_entry_byte_offset = 256
; CHECK-G-MESA3D-NEXT: kernel_code_prefetch_byte_size = 0
; CHECK-G-MESA3D-NEXT: granulated_workitem_vgpr_count = 0
; CHECK-G-MESA3D-NEXT: granulated_wavefront_sgpr_count = 0
; CHECK-G-MESA3D-NEXT: priority = 0
; CHECK-G-MESA3D-NEXT: float_mode = 240
; CHECK-G-MESA3D-NEXT: priv = 0
; CHECK-G-MESA3D-NEXT: enable_dx10_clamp = 0
; CHECK-G-MESA3D-NEXT: debug_mode = 0
; CHECK-G-MESA3D-NEXT: enable_ieee_mode = 0
; CHECK-G-MESA3D-NEXT: enable_wgp_mode = 0
; CHECK-G-MESA3D-NEXT: enable_mem_ordered = 1
; CHECK-G-MESA3D-NEXT: enable_fwd_progress = 1
; CHECK-G-MESA3D-NEXT: enable_sgpr_private_segment_wave_byte_offset = 0
; CHECK-G-MESA3D-NEXT: user_sgpr_count = 2
; CHECK-G-MESA3D-NEXT: enable_trap_handler = 0
; CHECK-G-MESA3D-NEXT: enable_sgpr_workgroup_id_x = 1
; CHECK-G-MESA3D-NEXT: enable_sgpr_workgroup_id_y = 1
; CHECK-G-MESA3D-NEXT: enable_sgpr_workgroup_id_z = 0
; CHECK-G-MESA3D-NEXT: enable_sgpr_workgroup_info = 0
; CHECK-G-MESA3D-NEXT: enable_vgpr_workitem_id = 0
; CHECK-G-MESA3D-NEXT: enable_exception_msb = 0
; CHECK-G-MESA3D-NEXT: granulated_lds_size = 0
; CHECK-G-MESA3D-NEXT: enable_exception = 0
; CHECK-G-MESA3D-NEXT: enable_sgpr_private_segment_buffer = 0
; CHECK-G-MESA3D-NEXT: enable_sgpr_dispatch_ptr = 0
; CHECK-G-MESA3D-NEXT: enable_sgpr_queue_ptr = 0
; CHECK-G-MESA3D-NEXT: enable_sgpr_kernarg_segment_ptr = 1
; CHECK-G-MESA3D-NEXT: enable_sgpr_dispatch_id = 0
; CHECK-G-MESA3D-NEXT: enable_sgpr_flat_scratch_init = 0
; CHECK-G-MESA3D-NEXT: enable_sgpr_private_segment_size = 0
; CHECK-G-MESA3D-NEXT: enable_sgpr_grid_workgroup_count_x = 0
; CHECK-G-MESA3D-NEXT: enable_sgpr_grid_workgroup_count_y = 0
; CHECK-G-MESA3D-NEXT: enable_sgpr_grid_workgroup_count_z = 0
; CHECK-G-MESA3D-NEXT: enable_wavefront_size32 = 1
; CHECK-G-MESA3D-NEXT: enable_ordered_append_gds = 0
; CHECK-G-MESA3D-NEXT: private_element_size = 1
; CHECK-G-MESA3D-NEXT: is_ptr64 = 1
; CHECK-G-MESA3D-NEXT: is_dynamic_callstack = 0
; CHECK-G-MESA3D-NEXT: is_debug_enabled = 0
; CHECK-G-MESA3D-NEXT: is_xnack_enabled = 0
; CHECK-G-MESA3D-NEXT: workitem_private_segment_byte_size = 0
; CHECK-G-MESA3D-NEXT: workgroup_group_segment_byte_size = 0
; CHECK-G-MESA3D-NEXT: gds_segment_byte_size = 0
; CHECK-G-MESA3D-NEXT: kernarg_segment_byte_size = 8
; CHECK-G-MESA3D-NEXT: workgroup_fbarrier_count = 0
; CHECK-G-MESA3D-NEXT: wavefront_sgpr_count = 2
; CHECK-G-MESA3D-NEXT: workitem_vgpr_count = 2
; CHECK-G-MESA3D-NEXT: reserved_vgpr_first = 0
; CHECK-G-MESA3D-NEXT: reserved_vgpr_count = 0
; CHECK-G-MESA3D-NEXT: reserved_sgpr_first = 0
; CHECK-G-MESA3D-NEXT: reserved_sgpr_count = 0
; CHECK-G-MESA3D-NEXT: debug_wavefront_private_segment_offset_sgpr = 0
; CHECK-G-MESA3D-NEXT: debug_private_segment_buffer_sgpr = 0
; CHECK-G-MESA3D-NEXT: kernarg_segment_alignment = 4
; CHECK-G-MESA3D-NEXT: group_segment_alignment = 4
; CHECK-G-MESA3D-NEXT: private_segment_alignment = 4
; CHECK-G-MESA3D-NEXT: wavefront_size = 5
; CHECK-G-MESA3D-NEXT: call_convention = -1
; CHECK-G-MESA3D-NEXT: runtime_loader_kernel_symbol = 0
; CHECK-G-MESA3D-NEXT: .end_amd_kernel_code_t
; CHECK-G-MESA3D-NEXT: ; %bb.0:
; CHECK-G-MESA3D-NEXT: s_load_b64 s[0:1], s[0:1], 0x0
; CHECK-G-MESA3D-NEXT: v_dual_mov_b32 v0, ttmp7 :: v_dual_mov_b32 v1, 0
; CHECK-G-MESA3D-NEXT: s_wait_kmcnt 0x0
; CHECK-G-MESA3D-NEXT: global_store_b32 v1, v0, s[0:1]
; CHECK-G-MESA3D-NEXT: s_endpgm
%id = call i32 @llvm.amdgcn.cluster.id.y()
store i32 %id, ptr addrspace(1) %out
ret void
}
define amdgpu_kernel void @test_cluster_id_z(ptr addrspace(1) %out) #1 {
; CHECK-UNKNOWN-LABEL: test_cluster_id_z:
; CHECK-UNKNOWN: ; %bb.0:
; CHECK-UNKNOWN-NEXT: s_load_b64 s[0:1], s[0:1], 0x24
; CHECK-UNKNOWN-NEXT: s_lshr_b32 s2, ttmp7, 16
; CHECK-UNKNOWN-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
; CHECK-UNKNOWN-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s2
; CHECK-UNKNOWN-NEXT: s_wait_kmcnt 0x0
; CHECK-UNKNOWN-NEXT: global_store_b32 v0, v1, s[0:1]
; CHECK-UNKNOWN-NEXT: s_endpgm
; CHECK-UNKNOWN: COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; CHECK-UNKNOWN: COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; CHECK-UNKNOWN: COMPUTE_PGM_RSRC2:TGID_Z_EN: 1
;
; CHECK-MESA3D-LABEL: test_cluster_id_z:
; CHECK-MESA3D: .amd_kernel_code_t
; CHECK-MESA3D-NEXT: amd_code_version_major = 1
; CHECK-MESA3D-NEXT: amd_code_version_minor = 2
; CHECK-MESA3D-NEXT: amd_machine_kind = 1
; CHECK-MESA3D-NEXT: amd_machine_version_major = 12
; CHECK-MESA3D-NEXT: amd_machine_version_minor = 5
; CHECK-MESA3D-NEXT: amd_machine_version_stepping = 0
; CHECK-MESA3D-NEXT: kernel_code_entry_byte_offset = 256
; CHECK-MESA3D-NEXT: kernel_code_prefetch_byte_size = 0
; CHECK-MESA3D-NEXT: granulated_workitem_vgpr_count = 0
; CHECK-MESA3D-NEXT: granulated_wavefront_sgpr_count = 0
; CHECK-MESA3D-NEXT: priority = 0
; CHECK-MESA3D-NEXT: float_mode = 240
; CHECK-MESA3D-NEXT: priv = 0
; CHECK-MESA3D-NEXT: enable_dx10_clamp = 0
; CHECK-MESA3D-NEXT: debug_mode = 0
; CHECK-MESA3D-NEXT: enable_ieee_mode = 0
; CHECK-MESA3D-NEXT: enable_wgp_mode = 0
; CHECK-MESA3D-NEXT: enable_mem_ordered = 1
; CHECK-MESA3D-NEXT: enable_fwd_progress = 1
; CHECK-MESA3D-NEXT: enable_sgpr_private_segment_wave_byte_offset = 0
; CHECK-MESA3D-NEXT: user_sgpr_count = 2
; CHECK-MESA3D-NEXT: enable_trap_handler = 0
; CHECK-MESA3D-NEXT: enable_sgpr_workgroup_id_x = 1
; CHECK-MESA3D-NEXT: enable_sgpr_workgroup_id_y = 0
; CHECK-MESA3D-NEXT: enable_sgpr_workgroup_id_z = 1
; CHECK-MESA3D-NEXT: enable_sgpr_workgroup_info = 0
; CHECK-MESA3D-NEXT: enable_vgpr_workitem_id = 0
; CHECK-MESA3D-NEXT: enable_exception_msb = 0
; CHECK-MESA3D-NEXT: granulated_lds_size = 0
; CHECK-MESA3D-NEXT: enable_exception = 0
; CHECK-MESA3D-NEXT: enable_sgpr_private_segment_buffer = 0
; CHECK-MESA3D-NEXT: enable_sgpr_dispatch_ptr = 0
; CHECK-MESA3D-NEXT: enable_sgpr_queue_ptr = 0
; CHECK-MESA3D-NEXT: enable_sgpr_kernarg_segment_ptr = 1
; CHECK-MESA3D-NEXT: enable_sgpr_dispatch_id = 0
; CHECK-MESA3D-NEXT: enable_sgpr_flat_scratch_init = 0
; CHECK-MESA3D-NEXT: enable_sgpr_private_segment_size = 0
; CHECK-MESA3D-NEXT: enable_sgpr_grid_workgroup_count_x = 0
; CHECK-MESA3D-NEXT: enable_sgpr_grid_workgroup_count_y = 0
; CHECK-MESA3D-NEXT: enable_sgpr_grid_workgroup_count_z = 0
; CHECK-MESA3D-NEXT: enable_wavefront_size32 = 1
; CHECK-MESA3D-NEXT: enable_ordered_append_gds = 0
; CHECK-MESA3D-NEXT: private_element_size = 1
; CHECK-MESA3D-NEXT: is_ptr64 = 1
; CHECK-MESA3D-NEXT: is_dynamic_callstack = 0
; CHECK-MESA3D-NEXT: is_debug_enabled = 0
; CHECK-MESA3D-NEXT: is_xnack_enabled = 0
; CHECK-MESA3D-NEXT: workitem_private_segment_byte_size = 0
; CHECK-MESA3D-NEXT: workgroup_group_segment_byte_size = 0
; CHECK-MESA3D-NEXT: gds_segment_byte_size = 0
; CHECK-MESA3D-NEXT: kernarg_segment_byte_size = 8
; CHECK-MESA3D-NEXT: workgroup_fbarrier_count = 0
; CHECK-MESA3D-NEXT: wavefront_sgpr_count = 3
; CHECK-MESA3D-NEXT: workitem_vgpr_count = 2
; CHECK-MESA3D-NEXT: reserved_vgpr_first = 0
; CHECK-MESA3D-NEXT: reserved_vgpr_count = 0
; CHECK-MESA3D-NEXT: reserved_sgpr_first = 0
; CHECK-MESA3D-NEXT: reserved_sgpr_count = 0
; CHECK-MESA3D-NEXT: debug_wavefront_private_segment_offset_sgpr = 0
; CHECK-MESA3D-NEXT: debug_private_segment_buffer_sgpr = 0
; CHECK-MESA3D-NEXT: kernarg_segment_alignment = 4
; CHECK-MESA3D-NEXT: group_segment_alignment = 4
; CHECK-MESA3D-NEXT: private_segment_alignment = 4
; CHECK-MESA3D-NEXT: wavefront_size = 5
; CHECK-MESA3D-NEXT: call_convention = -1
; CHECK-MESA3D-NEXT: runtime_loader_kernel_symbol = 0
; CHECK-MESA3D-NEXT: .end_amd_kernel_code_t
; CHECK-MESA3D-NEXT: ; %bb.0:
; CHECK-MESA3D-NEXT: s_load_b64 s[0:1], s[0:1], 0x0
; CHECK-MESA3D-NEXT: s_lshr_b32 s2, ttmp7, 16
; CHECK-MESA3D-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
; CHECK-MESA3D-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s2
; CHECK-MESA3D-NEXT: s_wait_kmcnt 0x0
; CHECK-MESA3D-NEXT: global_store_b32 v0, v1, s[0:1]
; CHECK-MESA3D-NEXT: s_endpgm
;
; CHECK-G-UNKNOWN-LABEL: test_cluster_id_z:
; CHECK-G-UNKNOWN: ; %bb.0:
; CHECK-G-UNKNOWN-NEXT: s_load_b64 s[0:1], s[0:1], 0x24
; CHECK-G-UNKNOWN-NEXT: s_lshr_b32 s2, ttmp7, 16
; CHECK-G-UNKNOWN-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
; CHECK-G-UNKNOWN-NEXT: v_dual_mov_b32 v1, 0 :: v_dual_mov_b32 v0, s2
; CHECK-G-UNKNOWN-NEXT: s_wait_kmcnt 0x0
; CHECK-G-UNKNOWN-NEXT: global_store_b32 v1, v0, s[0:1]
; CHECK-G-UNKNOWN-NEXT: s_endpgm
; CHECK-G-UNKNOWN: COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; CHECK-G-UNKNOWN: COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; CHECK-G-UNKNOWN: COMPUTE_PGM_RSRC2:TGID_Z_EN: 1
;
; CHECK-G-MESA3D-LABEL: test_cluster_id_z:
; CHECK-G-MESA3D: .amd_kernel_code_t
; CHECK-G-MESA3D-NEXT: amd_code_version_major = 1
; CHECK-G-MESA3D-NEXT: amd_code_version_minor = 2
; CHECK-G-MESA3D-NEXT: amd_machine_kind = 1
; CHECK-G-MESA3D-NEXT: amd_machine_version_major = 12
; CHECK-G-MESA3D-NEXT: amd_machine_version_minor = 5
; CHECK-G-MESA3D-NEXT: amd_machine_version_stepping = 0
; CHECK-G-MESA3D-NEXT: kernel_code_entry_byte_offset = 256
; CHECK-G-MESA3D-NEXT: kernel_code_prefetch_byte_size = 0
; CHECK-G-MESA3D-NEXT: granulated_workitem_vgpr_count = 0
; CHECK-G-MESA3D-NEXT: granulated_wavefront_sgpr_count = 0
; CHECK-G-MESA3D-NEXT: priority = 0
; CHECK-G-MESA3D-NEXT: float_mode = 240
; CHECK-G-MESA3D-NEXT: priv = 0
; CHECK-G-MESA3D-NEXT: enable_dx10_clamp = 0
; CHECK-G-MESA3D-NEXT: debug_mode = 0
; CHECK-G-MESA3D-NEXT: enable_ieee_mode = 0
; CHECK-G-MESA3D-NEXT: enable_wgp_mode = 0
; CHECK-G-MESA3D-NEXT: enable_mem_ordered = 1
; CHECK-G-MESA3D-NEXT: enable_fwd_progress = 1
; CHECK-G-MESA3D-NEXT: enable_sgpr_private_segment_wave_byte_offset = 0
; CHECK-G-MESA3D-NEXT: user_sgpr_count = 2
; CHECK-G-MESA3D-NEXT: enable_trap_handler = 0
; CHECK-G-MESA3D-NEXT: enable_sgpr_workgroup_id_x = 1
; CHECK-G-MESA3D-NEXT: enable_sgpr_workgroup_id_y = 0
; CHECK-G-MESA3D-NEXT: enable_sgpr_workgroup_id_z = 1
; CHECK-G-MESA3D-NEXT: enable_sgpr_workgroup_info = 0
; CHECK-G-MESA3D-NEXT: enable_vgpr_workitem_id = 0
; CHECK-G-MESA3D-NEXT: enable_exception_msb = 0
; CHECK-G-MESA3D-NEXT: granulated_lds_size = 0
; CHECK-G-MESA3D-NEXT: enable_exception = 0
; CHECK-G-MESA3D-NEXT: enable_sgpr_private_segment_buffer = 0
; CHECK-G-MESA3D-NEXT: enable_sgpr_dispatch_ptr = 0
; CHECK-G-MESA3D-NEXT: enable_sgpr_queue_ptr = 0
; CHECK-G-MESA3D-NEXT: enable_sgpr_kernarg_segment_ptr = 1
; CHECK-G-MESA3D-NEXT: enable_sgpr_dispatch_id = 0
; CHECK-G-MESA3D-NEXT: enable_sgpr_flat_scratch_init = 0
; CHECK-G-MESA3D-NEXT: enable_sgpr_private_segment_size = 0
; CHECK-G-MESA3D-NEXT: enable_sgpr_grid_workgroup_count_x = 0
; CHECK-G-MESA3D-NEXT: enable_sgpr_grid_workgroup_count_y = 0
; CHECK-G-MESA3D-NEXT: enable_sgpr_grid_workgroup_count_z = 0
; CHECK-G-MESA3D-NEXT: enable_wavefront_size32 = 1
; CHECK-G-MESA3D-NEXT: enable_ordered_append_gds = 0
; CHECK-G-MESA3D-NEXT: private_element_size = 1
; CHECK-G-MESA3D-NEXT: is_ptr64 = 1
; CHECK-G-MESA3D-NEXT: is_dynamic_callstack = 0
; CHECK-G-MESA3D-NEXT: is_debug_enabled = 0
; CHECK-G-MESA3D-NEXT: is_xnack_enabled = 0
; CHECK-G-MESA3D-NEXT: workitem_private_segment_byte_size = 0
; CHECK-G-MESA3D-NEXT: workgroup_group_segment_byte_size = 0
; CHECK-G-MESA3D-NEXT: gds_segment_byte_size = 0
; CHECK-G-MESA3D-NEXT: kernarg_segment_byte_size = 8
; CHECK-G-MESA3D-NEXT: workgroup_fbarrier_count = 0
; CHECK-G-MESA3D-NEXT: wavefront_sgpr_count = 3
; CHECK-G-MESA3D-NEXT: workitem_vgpr_count = 2
; CHECK-G-MESA3D-NEXT: reserved_vgpr_first = 0
; CHECK-G-MESA3D-NEXT: reserved_vgpr_count = 0
; CHECK-G-MESA3D-NEXT: reserved_sgpr_first = 0
; CHECK-G-MESA3D-NEXT: reserved_sgpr_count = 0
; CHECK-G-MESA3D-NEXT: debug_wavefront_private_segment_offset_sgpr = 0
; CHECK-G-MESA3D-NEXT: debug_private_segment_buffer_sgpr = 0
; CHECK-G-MESA3D-NEXT: kernarg_segment_alignment = 4
; CHECK-G-MESA3D-NEXT: group_segment_alignment = 4
; CHECK-G-MESA3D-NEXT: private_segment_alignment = 4
; CHECK-G-MESA3D-NEXT: wavefront_size = 5
; CHECK-G-MESA3D-NEXT: call_convention = -1
; CHECK-G-MESA3D-NEXT: runtime_loader_kernel_symbol = 0
; CHECK-G-MESA3D-NEXT: .end_amd_kernel_code_t
; CHECK-G-MESA3D-NEXT: ; %bb.0:
; CHECK-G-MESA3D-NEXT: s_load_b64 s[0:1], s[0:1], 0x0
; CHECK-G-MESA3D-NEXT: s_lshr_b32 s2, ttmp7, 16
; CHECK-G-MESA3D-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
; CHECK-G-MESA3D-NEXT: v_dual_mov_b32 v1, 0 :: v_dual_mov_b32 v0, s2
; CHECK-G-MESA3D-NEXT: s_wait_kmcnt 0x0
; CHECK-G-MESA3D-NEXT: global_store_b32 v1, v0, s[0:1]
; CHECK-G-MESA3D-NEXT: s_endpgm
%id = call i32 @llvm.amdgcn.cluster.id.z()
store i32 %id, ptr addrspace(1) %out
ret void
}

2
llvm/test/CodeGen/AMDGPU/mul24-pass-ordering.ll
View File

@ -265,6 +265,6 @@ declare float @llvm.fmuladd.f32(float, float, float) #1
attributes #0 = { nounwind willreturn "denormal-fp-math-f32"="preserve-sign,preserve-sign" }
attributes #1 = { nounwind readnone speculatable }
attributes #2 = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "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" "denormal-fp-math-f32"="preserve-sign,preserve-sign" }
attributes #2 = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" "denormal-fp-math-f32"="preserve-sign,preserve-sign" }
!0 = !{float 2.500000e+00}

2
llvm/test/CodeGen/AMDGPU/non-entry-alloca.ll
View File

@ -406,7 +406,7 @@ bb.1:
declare i32 @llvm.amdgcn.workitem.id.x() #0
attributes #0 = { nounwind readnone speculatable }
attributes #1 = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "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" }
attributes #1 = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
!llvm.module.flags = !{!0}
!0 = !{i32 1, !"amdhsa_code_object_version", i32 CODE_OBJECT_VERSION}

2
llvm/test/CodeGen/AMDGPU/partial-regcopy-and-spill-missed-at-regalloc.ll
View File

@ -104,4 +104,4 @@ define amdgpu_kernel void @partial_copy(<4 x i32> %arg) #0 {
declare <4 x i32> @llvm.amdgcn.mfma.i32.4x4x4i8(i32, i32, <4 x i32>, i32, i32, i32)
attributes #0 = { nounwind "amdgpu-num-vgpr"="5" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-lds-kernel-id" "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" }
attributes #0 = { nounwind "amdgpu-num-vgpr"="5" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }

2
llvm/test/CodeGen/AMDGPU/promote-constOffset-to-imm.ll
View File

@ -2688,4 +2688,4 @@ end:
}
attributes #0 = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "target-cpu"="fiji" "uniform-work-group-size"="false" }
attributes #0 = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "target-cpu"="fiji" "uniform-work-group-size"="false" }

2
llvm/test/CodeGen/AMDGPU/sibling-call.ll
View File

@ -1072,4 +1072,4 @@ entry:
}
attributes #0 = { nounwind }
attributes #1 = { nounwind noinline "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" }
attributes #1 = { nounwind noinline "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-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }

62
llvm/test/CodeGen/AMDGPU/spill-agpr.ll
View File

@ -363,11 +363,6 @@ use:
define amdgpu_kernel void @max_6regs_used_8a(ptr addrspace(1) %arg) #4 {
; GFX908-LABEL: max_6regs_used_8a:
; GFX908: ; %bb.0:
; GFX908-NEXT: s_mov_b32 s4, SCRATCH_RSRC_DWORD0
; GFX908-NEXT: s_mov_b32 s5, SCRATCH_RSRC_DWORD1
; GFX908-NEXT: s_mov_b32 s6, -1
; GFX908-NEXT: s_mov_b32 s7, 0xe00000
; GFX908-NEXT: s_add_u32 s4, s4, s3
; GFX908-NEXT: s_load_dwordx2 s[2:3], s[0:1], 0x24
; GFX908-NEXT: ;;#ASMSTART
; GFX908-NEXT: ; def v1
@ -376,22 +371,27 @@ define amdgpu_kernel void @max_6regs_used_8a(ptr addrspace(1) %arg) #4 {
; GFX908-NEXT: ;;#ASMSTART
; GFX908-NEXT: ; def a[0:3]
; GFX908-NEXT: ;;#ASMEND
; GFX908-NEXT: s_addc_u32 s5, s5, 0
; GFX908-NEXT: s_mov_b32 s8, SCRATCH_RSRC_DWORD0
; GFX908-NEXT: v_accvgpr_write_b32 a4, v1
; GFX908-NEXT: s_waitcnt lgkmcnt(0)
; GFX908-NEXT: global_load_dwordx4 v[0:3], v4, s[2:3]
; GFX908-NEXT: s_mov_b32 s9, SCRATCH_RSRC_DWORD1
; GFX908-NEXT: s_mov_b32 s10, -1
; GFX908-NEXT: s_mov_b32 s11, 0xe00000
; GFX908-NEXT: s_add_u32 s8, s8, s5
; GFX908-NEXT: s_addc_u32 s9, s9, 0
; GFX908-NEXT: v_accvgpr_read_b32 v5, a0 ; Reload Reuse
; GFX908-NEXT: s_nop 1
; GFX908-NEXT: buffer_store_dword v5, off, s[4:7], 0 ; 4-byte Folded Spill
; GFX908-NEXT: buffer_store_dword v5, off, s[8:11], 0 ; 4-byte Folded Spill
; GFX908-NEXT: v_accvgpr_read_b32 v5, a1 ; Reload Reuse
; GFX908-NEXT: s_nop 1
; GFX908-NEXT: buffer_store_dword v5, off, s[4:7], 0 offset:4 ; 4-byte Folded Spill
; GFX908-NEXT: buffer_store_dword v5, off, s[8:11], 0 offset:4 ; 4-byte Folded Spill
; GFX908-NEXT: v_accvgpr_read_b32 v5, a2 ; Reload Reuse
; GFX908-NEXT: s_nop 1
; GFX908-NEXT: buffer_store_dword v5, off, s[4:7], 0 offset:8 ; 4-byte Folded Spill
; GFX908-NEXT: buffer_store_dword v5, off, s[8:11], 0 offset:8 ; 4-byte Folded Spill
; GFX908-NEXT: v_accvgpr_read_b32 v5, a3 ; Reload Reuse
; GFX908-NEXT: s_nop 1
; GFX908-NEXT: buffer_store_dword v5, off, s[4:7], 0 offset:12 ; 4-byte Folded Spill
; GFX908-NEXT: buffer_store_dword v5, off, s[8:11], 0 offset:12 ; 4-byte Folded Spill
; GFX908-NEXT: s_waitcnt vmcnt(4)
; GFX908-NEXT: v_accvgpr_write_b32 a0, v0
; GFX908-NEXT: v_accvgpr_write_b32 a1, v1
@ -407,11 +407,11 @@ define amdgpu_kernel void @max_6regs_used_8a(ptr addrspace(1) %arg) #4 {
; GFX908-NEXT: v_accvgpr_read_b32 v3, a3
; GFX908-NEXT: s_nop 1
; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[2:3]
; GFX908-NEXT: buffer_load_dword v0, off, s[4:7], 0 ; 4-byte Folded Reload
; GFX908-NEXT: buffer_load_dword v0, off, s[8:11], 0 ; 4-byte Folded Reload
; GFX908-NEXT: s_nop 0
; GFX908-NEXT: buffer_load_dword v1, off, s[4:7], 0 offset:4 ; 4-byte Folded Reload
; GFX908-NEXT: buffer_load_dword v2, off, s[4:7], 0 offset:8 ; 4-byte Folded Reload
; GFX908-NEXT: buffer_load_dword v3, off, s[4:7], 0 offset:12 ; 4-byte Folded Reload
; GFX908-NEXT: buffer_load_dword v1, off, s[8:11], 0 offset:4 ; 4-byte Folded Reload
; GFX908-NEXT: buffer_load_dword v2, off, s[8:11], 0 offset:8 ; 4-byte Folded Reload
; GFX908-NEXT: buffer_load_dword v3, off, s[8:11], 0 offset:12 ; 4-byte Folded Reload
; GFX908-NEXT: s_waitcnt vmcnt(0)
; GFX908-NEXT: global_store_dwordx4 v[0:1], v[0:3], off
; GFX908-NEXT: s_waitcnt vmcnt(0)
@ -423,24 +423,24 @@ define amdgpu_kernel void @max_6regs_used_8a(ptr addrspace(1) %arg) #4 {
;
; GFX90A-LABEL: max_6regs_used_8a:
; GFX90A: ; %bb.0:
; GFX90A-NEXT: s_mov_b32 s4, SCRATCH_RSRC_DWORD0
; GFX90A-NEXT: s_mov_b32 s5, SCRATCH_RSRC_DWORD1
; GFX90A-NEXT: s_mov_b32 s6, -1
; GFX90A-NEXT: s_mov_b32 s7, 0xe00000
; GFX90A-NEXT: s_add_u32 s4, s4, s3
; GFX90A-NEXT: s_mov_b32 s8, SCRATCH_RSRC_DWORD0
; GFX90A-NEXT: s_mov_b32 s9, SCRATCH_RSRC_DWORD1
; GFX90A-NEXT: s_mov_b32 s10, -1
; GFX90A-NEXT: s_load_dwordx2 s[2:3], s[0:1], 0x24
; GFX90A-NEXT: s_addc_u32 s5, s5, 0
; GFX90A-NEXT: s_mov_b32 s11, 0xe00000
; GFX90A-NEXT: s_add_u32 s8, s8, s5
; GFX90A-NEXT: s_addc_u32 s9, s9, 0
; GFX90A-NEXT: ;;#ASMSTART
; GFX90A-NEXT: ; def v1
; GFX90A-NEXT: ;;#ASMEND
; GFX90A-NEXT: ;;#ASMSTART
; GFX90A-NEXT: ; def a[0:3]
; GFX90A-NEXT: ;;#ASMEND
; GFX90A-NEXT: buffer_store_dword a0, off, s[4:7], 0 ; 4-byte Folded Spill
; GFX90A-NEXT: buffer_store_dword a0, off, s[8:11], 0 ; 4-byte Folded Spill
; GFX90A-NEXT: s_nop 0
; GFX90A-NEXT: buffer_store_dword a1, off, s[4:7], 0 offset:4 ; 4-byte Folded Spill
; GFX90A-NEXT: buffer_store_dword a2, off, s[4:7], 0 offset:8 ; 4-byte Folded Spill
; GFX90A-NEXT: buffer_store_dword a3, off, s[4:7], 0 offset:12 ; 4-byte Folded Spill
; GFX90A-NEXT: buffer_store_dword a1, off, s[8:11], 0 offset:4 ; 4-byte Folded Spill
; GFX90A-NEXT: buffer_store_dword a2, off, s[8:11], 0 offset:8 ; 4-byte Folded Spill
; GFX90A-NEXT: buffer_store_dword a3, off, s[8:11], 0 offset:12 ; 4-byte Folded Spill
; GFX90A-NEXT: v_lshlrev_b32_e32 v0, 4, v0
; GFX90A-NEXT: s_waitcnt lgkmcnt(0)
; GFX90A-NEXT: global_load_dwordx4 a[0:3], v0, s[2:3]
@ -450,10 +450,10 @@ define amdgpu_kernel void @max_6regs_used_8a(ptr addrspace(1) %arg) #4 {
; GFX90A-NEXT: v_mfma_f32_4x4x1f32 a[0:3], v2, v2, a[0:3]
; GFX90A-NEXT: s_nop 4
; GFX90A-NEXT: global_store_dwordx4 v0, a[0:3], s[2:3]
; GFX90A-NEXT: buffer_load_dword v2, off, s[4:7], 0 ; 4-byte Folded Reload
; GFX90A-NEXT: buffer_load_dword v3, off, s[4:7], 0 offset:4 ; 4-byte Folded Reload
; GFX90A-NEXT: buffer_load_dword v4, off, s[4:7], 0 offset:8 ; 4-byte Folded Reload
; GFX90A-NEXT: buffer_load_dword v5, off, s[4:7], 0 offset:12 ; 4-byte Folded Reload
; GFX90A-NEXT: buffer_load_dword v2, off, s[8:11], 0 ; 4-byte Folded Reload
; GFX90A-NEXT: buffer_load_dword v3, off, s[8:11], 0 offset:4 ; 4-byte Folded Reload
; GFX90A-NEXT: buffer_load_dword v4, off, s[8:11], 0 offset:8 ; 4-byte Folded Reload
; GFX90A-NEXT: buffer_load_dword v5, off, s[8:11], 0 offset:12 ; 4-byte Folded Reload
; GFX90A-NEXT: s_waitcnt vmcnt(0)
; GFX90A-NEXT: global_store_dwordx4 v[0:1], v[2:5], off
; GFX90A-NEXT: s_waitcnt vmcnt(0)
@ -479,6 +479,12 @@ declare <16 x float> @llvm.amdgcn.mfma.f32.16x16x1f32(float, float, <16 x float>
declare <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float, float, <4 x float>, i32, i32, i32)
declare <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float, float, <32 x float>, i32, i32, i32)
attributes #1 = { nounwind "amdgpu-num-vgpr"="10" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }
attributes #2 = { nounwind "amdgpu-num-vgpr"="12" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }
attributes #3 = { nounwind "amdgpu-num-vgpr"="32" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }
attributes #4 = { nounwind "amdgpu-num-vgpr"="6" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }
attributes #1 = { nounwind "amdgpu-num-vgpr"="10" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }
attributes #2 = { nounwind "amdgpu-num-vgpr"="12" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }
attributes #3 = { nounwind "amdgpu-num-vgpr"="32" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }

2
llvm/test/CodeGen/AMDGPU/stack-pointer-offset-relative-frameindex.ll
View File

@ -157,4 +157,4 @@ shader_eval_surface.exit: ; preds = %entry
declare hidden i32 @svm_eval_nodes(ptr addrspace(5), ptr addrspace(5), ptr addrspace(5), i32, i32) local_unnamed_addr #0
attributes #0 = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "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" }
attributes #0 = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }

2
llvm/test/CodeGen/AMDGPU/tid-kd-xnack-any.ll
View File

@ -23,7 +23,7 @@ entry:
ret void
}
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-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }
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-workgroup-id-x" "amdgpu-no-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }
!llvm.module.flags = !{!0}
!0 = !{i32 1, !"amdhsa_code_object_version", i32 400}

2
llvm/test/CodeGen/AMDGPU/tid-kd-xnack-off.ll
View File

@ -23,7 +23,7 @@ entry:
ret void
}
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-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }
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-workgroup-id-x" "amdgpu-no-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }
!llvm.module.flags = !{!0}
!0 = !{i32 1, !"amdhsa_code_object_version", i32 400}

2
llvm/test/CodeGen/AMDGPU/tid-kd-xnack-on.ll
View File

@ -23,7 +23,7 @@ entry:
ret void
}
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-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }
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-workgroup-id-x" "amdgpu-no-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }
!llvm.module.flags = !{!0}
!0 = !{i32 1, !"amdhsa_code_object_version", i32 400}

2
llvm/test/CodeGen/MIR/AMDGPU/long-branch-reg-all-sgpr-used.ll
View File

@ -540,7 +540,7 @@
; Function Attrs: nocallback nofree nosync nounwind speculatable willreturn memory(none)
declare align 4 ptr addrspace(4) @llvm.amdgcn.kernarg.segment.ptr() #2
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-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" }
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-workgroup-id-x" "amdgpu-no-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }
attributes #1 = { nounwind }
attributes #2 = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
!0 = !{}

2
llvm/test/CodeGen/MIR/AMDGPU/machine-function-info-after-pei.ll
View File

@ -63,4 +63,4 @@ define amdgpu_kernel void @scavenge_fi(ptr addrspace(1) %out, i32 %in) #0 {
ret void
}
attributes #0 = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-lds-kernel-id" "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" }
attributes #0 = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }

2
llvm/test/CodeGen/MIR/AMDGPU/machine-function-info-long-branch-reg-debug.ll
View File

@ -99,7 +99,7 @@
; Function Attrs: convergent nocallback nofree nounwind willreturn
declare void @llvm.amdgcn.end.cf.i64(i64) #2
attributes #0 = { "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-flat-scratch-init" "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" "uniform-work-group-size"="false" }
attributes #0 = { "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-flat-scratch-init" "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-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
attributes #1 = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
attributes #2 = { convergent nocallback nofree nounwind willreturn }
attributes #3 = { convergent nocallback nofree nounwind willreturn memory(none) }

2
llvm/test/CodeGen/MIR/AMDGPU/machine-function-info-long-branch-reg.ll
View File

@ -73,5 +73,5 @@ bb4:
ret void
}
attributes #0 = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-flat-scratch-init" "amdgpu-no-implicitarg-ptr" "amdgpu-no-lds-kernel-id" "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" }
attributes #0 = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-flat-scratch-init" "amdgpu-no-implicitarg-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }
attributes #1 = { nounwind readnone }