Skip to content

Commit f501081

Browse files
kerbowakerbowa
committed
Diagnostic error and move LDSKernelId.
Use diagnostic error. Unify preload kernarg SGPR accounting and move LDSKernelId to UserSGPRInfo so that synthetic SGPRs can be properly tracked when calculating the number of available registers for preloads.
1 parent 87cc1db commit f501081

10 files changed

+134
-50
lines changed

‎llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp‎

Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -520,6 +520,17 @@ bool AMDGPUCallLowering::lowerFormalArgumentsKernel(
520520

521521
// TODO: Align down to dword alignment and extract bits for extending loads.
522522
for (auto &Arg : F.args()) {
523+
// Hidden arguments that are in the kernel signature must be preloded to
524+
// user SGPRs, or loaded via the implicit_arg ptr. Print a diagnostic
525+
// error if a hidden argument is in the argument list and is not
526+
// preloaded.
527+
if (Arg.hasAttribute("amdgpu-hidden-argument")) {
528+
DiagnosticInfoUnsupported NonPreloadHiddenArg(
529+
*Arg.getParent(),
530+
"Hidden argument in kernel signature was not preloaded");
531+
F.getContext().diagnose(NonPreloadHiddenArg);
532+
}
533+
523534
const bool IsByRef = Arg.hasByRefAttr();
524535
Type *ArgTy = IsByRef ? Arg.getParamByRefType() : Arg.getType();
525536
unsigned AllocSize = DL.getTypeAllocSize(ArgTy);

‎llvm/lib/Target/AMDGPU/AMDGPULowerKernelArguments.cpp‎

Lines changed: 6 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -144,19 +144,19 @@ class PreloadKernelArgInfo {
144144
// arguments.
145145
void setInitialFreeUserSGPRsCount() {
146146
GCNUserSGPRUsageInfo UserSGPRInfo(F, ST);
147-
NumFreeUserSGPRs =
148-
UserSGPRInfo.getNumFreeUserSGPRs() - 1 /* Synthetic SGPRs*/;
147+
NumFreeUserSGPRs = UserSGPRInfo.getNumFreeKernargPreloadSGPRs();
149148
}
150149

151150
bool tryAllocPreloadSGPRs(unsigned AllocSize, uint64_t ArgOffset,
152151
uint64_t LastExplicitArgOffset) {
153152
// Check if this argument may be loaded into the same register as the
154153
// previous argument.
155-
if (ArgOffset == LastExplicitArgOffset && !isAligned(Align(4), ArgOffset) &&
156-
AllocSize < 4)
154+
if (ArgOffset - LastExplicitArgOffset < 4 &&
155+
!isAligned(Align(4), ArgOffset))
157156
return true;
158157

159158
// Pad SGPRs for kernarg alignment.
159+
ArgOffset = alignDown(ArgOffset, 4);
160160
unsigned Padding = ArgOffset - LastExplicitArgOffset;
161161
unsigned PaddingSGPRs = alignTo(Padding, 4) / 4;
162162
unsigned NumPreloadSGPRs = alignTo(AllocSize, 4) / 4;
@@ -228,7 +228,8 @@ class PreloadKernelArgInfo {
228228
LastExplicitArgOffset))
229229
return true;
230230

231-
LastExplicitArgOffset = LoadOffset + LoadSize;
231+
LastExplicitArgOffset =
232+
ImplicitArgsBaseOffset + LoadOffset + LoadSize;
232233
return false;
233234
});
234235

‎llvm/lib/Target/AMDGPU/GCNSubtarget.cpp‎

Lines changed: 9 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -748,6 +748,10 @@ GCNUserSGPRUsageInfo::GCNUserSGPRUsageInfo(const Function &F,
748748
FlatScratchInit = true;
749749
}
750750

751+
if (!AMDGPU::isGraphics(CC) && !IsKernel &&
752+
!F.hasFnAttribute("amdgpu-no-lds-kernel-id"))
753+
LDSKernelId = true;
754+
751755
if (hasImplicitBufferPtr())
752756
NumUsedUserSGPRs += getNumUserSGPRForField(ImplicitBufferPtrID);
753757

@@ -771,6 +775,9 @@ GCNUserSGPRUsageInfo::GCNUserSGPRUsageInfo(const Function &F,
771775

772776
if (hasPrivateSegmentSize())
773777
NumUsedUserSGPRs += getNumUserSGPRForField(PrivateSegmentSizeID);
778+
779+
if (hasLDSKernelId())
780+
NumSyntheticSGPRs += getNumUserSGPRForField(LDSKernelIdID);
774781
}
775782

776783
void GCNUserSGPRUsageInfo::allocKernargPreloadSGPRs(unsigned NumSGPRs) {
@@ -779,6 +786,6 @@ void GCNUserSGPRUsageInfo::allocKernargPreloadSGPRs(unsigned NumSGPRs) {
779786
NumUsedUserSGPRs += NumSGPRs;
780787
}
781788

782-
unsigned GCNUserSGPRUsageInfo::getNumFreeUserSGPRs() {
783-
return AMDGPU::getMaxNumUserSGPRs(ST) - NumUsedUserSGPRs;
789+
unsigned GCNUserSGPRUsageInfo::getNumFreeKernargPreloadSGPRs() {
790+
return AMDGPU::getMaxNumUserSGPRs(ST) - (NumUsedUserSGPRs + NumSyntheticSGPRs);
784791
}

‎llvm/lib/Target/AMDGPU/GCNSubtarget.h‎

Lines changed: 12 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1676,11 +1676,13 @@ class GCNUserSGPRUsageInfo {
16761676

16771677
bool hasPrivateSegmentSize() const { return PrivateSegmentSize; }
16781678

1679+
bool hasLDSKernelId() const { return LDSKernelId; }
1680+
16791681
unsigned getNumKernargPreloadSGPRs() const { return NumKernargPreloadSGPRs; }
16801682

16811683
unsigned getNumUsedUserSGPRs() const { return NumUsedUserSGPRs; }
16821684

1683-
unsigned getNumFreeUserSGPRs();
1685+
unsigned getNumFreeKernargPreloadSGPRs();
16841686

16851687
void allocKernargPreloadSGPRs(unsigned NumSGPRs);
16861688

@@ -1692,11 +1694,12 @@ class GCNUserSGPRUsageInfo {
16921694
KernargSegmentPtrID = 4,
16931695
DispatchIdID = 5,
16941696
FlatScratchInitID = 6,
1695-
PrivateSegmentSizeID = 7
1697+
PrivateSegmentSizeID = 7,
1698+
LDSKernelIdID = 8
16961699
};
16971700

16981701
// Returns the size in number of SGPRs for preload user SGPR field.
1699-
static unsigned getNumUserSGPRForField(UserSGPRID ID) {
1702+
static constexpr unsigned getNumUserSGPRForField(UserSGPRID ID) {
17001703
switch (ID) {
17011704
case ImplicitBufferPtrID:
17021705
return 2;
@@ -1714,6 +1717,8 @@ class GCNUserSGPRUsageInfo {
17141717
return 2;
17151718
case PrivateSegmentSizeID:
17161719
return 1;
1720+
case LDSKernelIdID:
1721+
return 1;
17171722
}
17181723
llvm_unreachable("Unknown UserSGPRID.");
17191724
}
@@ -1742,9 +1747,13 @@ class GCNUserSGPRUsageInfo {
17421747

17431748
bool PrivateSegmentSize = false;
17441749

1750+
bool LDSKernelId = false;
1751+
17451752
unsigned NumKernargPreloadSGPRs = 0;
17461753

17471754
unsigned NumUsedUserSGPRs = 0;
1755+
1756+
unsigned NumSyntheticSGPRs = 0;
17481757
};
17491758

17501759
} // end namespace llvm

‎llvm/lib/Target/AMDGPU/SIISelLowering.cpp‎

Lines changed: 17 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -2418,7 +2418,7 @@ void SITargetLowering::allocateSpecialInputSGPRs(
24182418
if (Info.hasWorkGroupIDZ())
24192419
allocateSGPR32Input(CCInfo, ArgInfo.WorkGroupIDZ);
24202420

2421-
if (Info.hasLDSKernelId())
2421+
if (UserSGPRInfo.hasLDSKernelId())
24222422
allocateSGPR32Input(CCInfo, ArgInfo.LDSKernelId);
24232423
}
24242424

@@ -2541,8 +2541,8 @@ void SITargetLowering::allocatePreloadKernArgSGPRs(
25412541
unsigned Padding = ArgOffset - LastExplicitArgOffset;
25422542
unsigned PaddingSGPRs = alignTo(Padding, 4) / 4;
25432543
// Check for free user SGPRs for preloading.
2544-
if (PaddingSGPRs + NumAllocSGPRs + 1 /*Synthetic SGPRs*/ >
2545-
SGPRInfo.getNumFreeUserSGPRs()) {
2544+
if (PaddingSGPRs + NumAllocSGPRs >
2545+
SGPRInfo.getNumFreeKernargPreloadSGPRs()) {
25462546
InPreloadSequence = false;
25472547
break;
25482548
}
@@ -2570,7 +2570,8 @@ void SITargetLowering::allocateLDSKernelId(CCState &CCInfo, MachineFunction &MF,
25702570
const SIRegisterInfo &TRI,
25712571
SIMachineFunctionInfo &Info) const {
25722572
// Always allocate this last since it is a synthetic preload.
2573-
if (Info.hasLDSKernelId()) {
2573+
const GCNUserSGPRUsageInfo &UserSGPRInfo = Info.getUserSGPRInfo();
2574+
if (UserSGPRInfo.hasLDSKernelId()) {
25742575
Register Reg = Info.addLDSKernelId();
25752576
MF.addLiveIn(Reg, &AMDGPU::SGPR_32RegClass);
25762577
CCInfo.AllocateReg(Reg);
@@ -2820,7 +2821,7 @@ SDValue SITargetLowering::LowerFormalArguments(
28202821
const GCNUserSGPRUsageInfo &UserSGPRInfo = Info->getUserSGPRInfo();
28212822
assert(!UserSGPRInfo.hasDispatchPtr() &&
28222823
!UserSGPRInfo.hasKernargSegmentPtr() && !Info->hasWorkGroupInfo() &&
2823-
!Info->hasLDSKernelId() && !Info->hasWorkItemIDX() &&
2824+
!UserSGPRInfo.hasLDSKernelId() && !Info->hasWorkItemIDX() &&
28242825
!Info->hasWorkItemIDY() && !Info->hasWorkItemIDZ());
28252826
(void)UserSGPRInfo;
28262827
if (!Subtarget->enableFlatScratch())
@@ -3020,13 +3021,20 @@ SDValue SITargetLowering::LowerFormalArguments(
30203021
NewArg = DAG.getMergeValues({NewArg, Chain}, DL);
30213022
}
30223023
} else {
3023-
#ifndef NDEBUG
3024+
// Hidden arguments that are in the kernel signature must be preloded to
3025+
// user SGPRs, or loaded via the implicit_arg ptr. Print a diagnostic
3026+
// error if a hidden argument is in the argument list and is not
3027+
// preloaded.
30243028
if (Arg.isOrigArg()) {
30253029
Argument *OrigArg = Fn.getArg(Arg.getOrigArgIndex());
3026-
assert(!OrigArg->hasAttribute("amdgpu-hidden-argument") &&
3027-
"Hidden arguments should be preloaded");
3030+
if (OrigArg->hasAttribute("amdgpu-hidden-argument")) {
3031+
DiagnosticInfoUnsupported NonPreloadHiddenArg(
3032+
*OrigArg->getParent(),
3033+
"Hidden argument in kernel signature was not preloaded",
3034+
DL.getDebugLoc());
3035+
DAG.getContext()->diagnose(NonPreloadHiddenArg);
3036+
}
30283037
}
3029-
#endif // NDEBUG
30303038

30313039
NewArg =
30323040
lowerKernargMemParameter(DAG, VT, MemVT, DL, Chain, Offset,

‎llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp‎

Lines changed: 1 addition & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -38,7 +38,7 @@ SIMachineFunctionInfo::SIMachineFunctionInfo(const Function &F,
3838
const GCNSubtarget *STI)
3939
: AMDGPUMachineFunction(F, *STI), Mode(F, *STI), GWSResourcePSV(getTM(STI)),
4040
UserSGPRInfo(F, *STI), WorkGroupIDX(false), WorkGroupIDY(false),
41-
WorkGroupIDZ(false), WorkGroupInfo(false), LDSKernelId(false),
41+
WorkGroupIDZ(false), WorkGroupInfo(false),
4242
PrivateSegmentWaveByteOffset(false), WorkItemIDX(false),
4343
WorkItemIDY(false), WorkItemIDZ(false), ImplicitArgPtr(false),
4444
GITPtrHigh(0xffffffff), HighBitsOf32BitAddress(0) {
@@ -131,9 +131,6 @@ SIMachineFunctionInfo::SIMachineFunctionInfo(const Function &F,
131131
if (!F.hasFnAttribute("amdgpu-no-workitem-id-z") &&
132132
ST.getMaxWorkitemID(F, 2) != 0)
133133
WorkItemIDZ = true;
134-
135-
if (!IsKernel && !F.hasFnAttribute("amdgpu-no-lds-kernel-id"))
136-
LDSKernelId = true;
137134
}
138135

139136
if (isEntryFunction()) {

‎llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h‎

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -461,7 +461,6 @@ class SIMachineFunctionInfo final : public AMDGPUMachineFunction,
461461
bool WorkGroupIDY : 1;
462462
bool WorkGroupIDZ : 1;
463463
bool WorkGroupInfo : 1;
464-
bool LDSKernelId : 1;
465464
bool PrivateSegmentWaveByteOffset : 1;
466465

467466
bool WorkItemIDX : 1; // Always initialized.
@@ -822,8 +821,6 @@ class SIMachineFunctionInfo final : public AMDGPUMachineFunction,
822821
return ArgInfo.WorkGroupInfo.getRegister();
823822
}
824823

825-
bool hasLDSKernelId() const { return LDSKernelId; }
826-
827824
// Add special VGPR inputs
828825
void setWorkItemIDX(ArgDescriptor Arg) {
829826
ArgInfo.WorkItemIDX = Arg;

‎llvm/test/CodeGen/AMDGPU/invalid-hidden-kernarg-in-kernel-signature.ll‎

Lines changed: 17 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,17 @@
1+
; RUN: not llc -global-isel=1 -mtriple=amdgcn--amdhsa -mcpu=gfx940 -start-after=amdgpu-lower-kernel-arguments < %s 2>&1 | FileCheck -check-prefix=ERROR %s
2+
; RUN: not llc -global-isel=0 -mtriple=amdgcn--amdhsa -mcpu=gfx940 -start-after=amdgpu-lower-kernel-arguments < %s 2>&1 | FileCheck -check-prefix=ERROR %s
3+
4+
define amdgpu_kernel void @no_free_sgprs_block_count_x_no_preload_diag(ptr addrspace(1) inreg %out, i512 inreg, i32 inreg "amdgpu-hidden-argument" %_hidden_block_count_x) #0 {
5+
; ERROR: error: <unknown>:0:0: in function no_free_sgprs_block_count_x_no_preload_diag void (ptr addrspace(1), i512, i32): Hidden argument in kernel signature was not preloaded
6+
store i32 %_hidden_block_count_x, ptr addrspace(1) %out
7+
ret void
8+
}
9+
10+
define amdgpu_kernel void @preloadremainder_z_no_preload_diag(ptr addrspace(1) inreg %out, i256 inreg, i32 inreg "amdgpu-hidden-argument" %_hidden_block_count_x, i32 inreg "amdgpu-hidden-argument" %_hidden_block_count_y, i32 inreg "amdgpu-hidden-argument" %_hidden_block_count_z, i16 inreg "amdgpu-hidden-argument" %_hidden_group_size_x, i16 inreg "amdgpu-hidden-argument" %_hidden_group_size_y, i16 inreg "amdgpu-hidden-argument" %_hidden_group_size_z, i16 inreg "amdgpu-hidden-argument" %_hidden_remainder_x, i16 inreg "amdgpu-hidden-argument" %_hidden_remainder_y, i16 inreg "amdgpu-hidden-argument" %_hidden_remainder_z) #0 {
11+
; ERROR: error: <unknown>:0:0: in function preloadremainder_z_no_preload_diag void (ptr addrspace(1), i256, i32, i32, i32, i16, i16, i16, i16, i16, i16): Hidden argument in kernel signature was not preloaded
12+
%conv = zext i16 %_hidden_remainder_z to i32
13+
store i32 %conv, ptr addrspace(1) %out
14+
ret void
15+
}
16+
17+
attributes #0 = { "amdgpu-no-agpr" "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" }

‎llvm/test/CodeGen/AMDGPU/preload-implicit-kernargs.ll‎

Lines changed: 59 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -439,13 +439,13 @@ define amdgpu_kernel void @preload_workgroup_size_xyz(ptr addrspace(1) inreg %ou
439439
; GFX90a: s_trap 2 ; Kernarg preload header. Trap with incompatible firmware that doesn't support preloading kernel arguments.
440440
; GFX90a-NEXT: .fill 63, 4, 0xbf800000 ; s_nop 0
441441
; GFX90a-NEXT: ; %bb.0:
442-
; GFX90a-NEXT: v_mov_b32_e32 v3, 0
443-
; GFX90a-NEXT: global_load_ushort v2, v3, s[4:5] offset:24
444442
; GFX90a-NEXT: s_lshr_b32 s0, s11, 16
445443
; GFX90a-NEXT: s_and_b32 s1, s11, 0xffff
444+
; GFX90a-NEXT: s_and_b32 s2, s12, 0xffff
445+
; GFX90a-NEXT: v_mov_b32_e32 v3, 0
446446
; GFX90a-NEXT: v_mov_b32_e32 v0, s1
447447
; GFX90a-NEXT: v_mov_b32_e32 v1, s0
448-
; GFX90a-NEXT: s_waitcnt vmcnt(0)
448+
; GFX90a-NEXT: v_mov_b32_e32 v2, s2
449449
; GFX90a-NEXT: global_store_dwordx3 v3, v[0:2], s[6:7]
450450
; GFX90a-NEXT: s_endpgm
451451
%imp_arg_ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
@@ -554,28 +554,27 @@ define amdgpu_kernel void @preloadremainder_xyz(ptr addrspace(1) inreg %out) #0
554554
; GFX940: s_trap 2 ; Kernarg preload header. Trap with incompatible firmware that doesn't support preloading kernel arguments.
555555
; GFX940-NEXT: .fill 63, 4, 0xbf800000 ; s_nop 0
556556
; GFX940-NEXT: ; %bb.0:
557+
; GFX940-NEXT: s_lshr_b32 s0, s9, 16
558+
; GFX940-NEXT: s_lshr_b32 s1, s8, 16
559+
; GFX940-NEXT: s_and_b32 s4, s9, 0xffff
557560
; GFX940-NEXT: v_mov_b32_e32 v3, 0
558-
; GFX940-NEXT: global_load_ushort v2, v3, s[0:1] offset:30
559-
; GFX940-NEXT: s_lshr_b32 s0, s8, 16
560-
; GFX940-NEXT: s_and_b32 s1, s9, 0xffff
561-
; GFX940-NEXT: v_mov_b32_e32 v0, s0
562-
; GFX940-NEXT: v_mov_b32_e32 v1, s1
563-
; GFX940-NEXT: s_waitcnt vmcnt(0)
561+
; GFX940-NEXT: v_mov_b32_e32 v0, s1
562+
; GFX940-NEXT: v_mov_b32_e32 v1, s4
563+
; GFX940-NEXT: v_mov_b32_e32 v2, s0
564564
; GFX940-NEXT: global_store_dwordx3 v3, v[0:2], s[2:3] sc0 sc1
565565
; GFX940-NEXT: s_endpgm
566566
;
567567
; GFX90a-LABEL: preloadremainder_xyz:
568568
; GFX90a: s_trap 2 ; Kernarg preload header. Trap with incompatible firmware that doesn't support preloading kernel arguments.
569569
; GFX90a-NEXT: .fill 63, 4, 0xbf800000 ; s_nop 0
570570
; GFX90a-NEXT: ; %bb.0:
571+
; GFX90a-NEXT: s_lshr_b32 s0, s13, 16
572+
; GFX90a-NEXT: s_lshr_b32 s1, s12, 16
573+
; GFX90a-NEXT: s_and_b32 s2, s13, 0xffff
571574
; GFX90a-NEXT: v_mov_b32_e32 v3, 0
572-
; GFX90a-NEXT: global_load_dword v0, v3, s[4:5] offset:26
573-
; GFX90a-NEXT: global_load_ushort v2, v3, s[4:5] offset:30
574-
; GFX90a-NEXT: s_lshr_b32 s0, s12, 16
575-
; GFX90a-NEXT: s_waitcnt vmcnt(1)
576-
; GFX90a-NEXT: v_lshrrev_b32_e32 v1, 16, v0
577-
; GFX90a-NEXT: v_mov_b32_e32 v0, s0
578-
; GFX90a-NEXT: s_waitcnt vmcnt(0)
575+
; GFX90a-NEXT: v_mov_b32_e32 v0, s1
576+
; GFX90a-NEXT: v_mov_b32_e32 v1, s2
577+
; GFX90a-NEXT: v_mov_b32_e32 v2, s0
579578
; GFX90a-NEXT: global_store_dwordx3 v3, v[0:2], s[6:7]
580579
; GFX90a-NEXT: s_endpgm
581580
%imp_arg_ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
@@ -600,10 +599,8 @@ define amdgpu_kernel void @no_free_sgprs_preloadremainder_z(ptr addrspace(1) inr
600599
; GFX940: s_trap 2 ; Kernarg preload header. Trap with incompatible firmware that doesn't support preloading kernel arguments.
601600
; GFX940-NEXT: .fill 63, 4, 0xbf800000 ; s_nop 0
602601
; GFX940-NEXT: ; %bb.0:
603-
; GFX940-NEXT: s_load_dword s0, s[4:5], 0x1c
602+
; GFX940-NEXT: s_lshr_b32 s0, s15, 16
604603
; GFX940-NEXT: v_mov_b32_e32 v0, 0
605-
; GFX940-NEXT: s_waitcnt lgkmcnt(0)
606-
; GFX940-NEXT: s_lshr_b32 s0, s0, 16
607604
; GFX940-NEXT: v_mov_b32_e32 v1, s0
608605
; GFX940-NEXT: global_store_dword v0, v1, s[8:9] sc0 sc1
609606
; GFX940-NEXT: s_endpgm
@@ -627,7 +624,7 @@ define amdgpu_kernel void @no_free_sgprs_preloadremainder_z(ptr addrspace(1) inr
627624
ret void
628625
}
629626

630-
; Check for consistency between isel and earlier passes preload SGPR accounting.
627+
; Check for consistency between isel and earlier passes preload SGPR accounting with max preload SGPRs.
631628

632629
define amdgpu_kernel void @preload_block_max_user_sgprs(ptr addrspace(1) inreg %out, i192 inreg %t0, i32 inreg %t1) #0 {
633630
; GFX940-LABEL: preload_block_max_user_sgprs:
@@ -655,4 +652,46 @@ define amdgpu_kernel void @preload_block_max_user_sgprs(ptr addrspace(1) inreg %
655652
ret void
656653
}
657654

655+
define amdgpu_kernel void @preload_block_count_z_workgroup_size_z_remainder_z(ptr addrspace(1) inreg %out) #0 {
656+
; GFX940-LABEL: preload_block_count_z_workgroup_size_z_remainder_z:
657+
; GFX940: s_trap 2 ; Kernarg preload header. Trap with incompatible firmware that doesn't support preloading kernel arguments.
658+
; GFX940-NEXT: .fill 63, 4, 0xbf800000 ; s_nop 0
659+
; GFX940-NEXT: ; %bb.0:
660+
; GFX940-NEXT: s_lshr_b32 s0, s9, 16
661+
; GFX940-NEXT: s_and_b32 s1, s8, 0xffff
662+
; GFX940-NEXT: v_mov_b32_e32 v3, 0
663+
; GFX940-NEXT: v_mov_b32_e32 v0, s6
664+
; GFX940-NEXT: v_mov_b32_e32 v1, s1
665+
; GFX940-NEXT: v_mov_b32_e32 v2, s0
666+
; GFX940-NEXT: global_store_dwordx3 v3, v[0:2], s[2:3] sc0 sc1
667+
; GFX940-NEXT: s_endpgm
668+
;
669+
; GFX90a-LABEL: preload_block_count_z_workgroup_size_z_remainder_z:
670+
; GFX90a: s_trap 2 ; Kernarg preload header. Trap with incompatible firmware that doesn't support preloading kernel arguments.
671+
; GFX90a-NEXT: .fill 63, 4, 0xbf800000 ; s_nop 0
672+
; GFX90a-NEXT: ; %bb.0:
673+
; GFX90a-NEXT: s_lshr_b32 s0, s13, 16
674+
; GFX90a-NEXT: s_and_b32 s1, s12, 0xffff
675+
; GFX90a-NEXT: v_mov_b32_e32 v3, 0
676+
; GFX90a-NEXT: v_mov_b32_e32 v0, s10
677+
; GFX90a-NEXT: v_mov_b32_e32 v1, s1
678+
; GFX90a-NEXT: v_mov_b32_e32 v2, s0
679+
; GFX90a-NEXT: global_store_dwordx3 v3, v[0:2], s[6:7]
680+
; GFX90a-NEXT: s_endpgm
681+
%imp_arg_ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
682+
%gep0 = getelementptr i8, ptr addrspace(4) %imp_arg_ptr, i32 8
683+
%gep1 = getelementptr i8, ptr addrspace(4) %imp_arg_ptr, i32 16
684+
%gep2 = getelementptr i8, ptr addrspace(4) %imp_arg_ptr, i32 22
685+
%load0 = load i32, ptr addrspace(4) %gep0
686+
%load1 = load i16, ptr addrspace(4) %gep1
687+
%load2 = load i16, ptr addrspace(4) %gep2
688+
%conv1 = zext i16 %load1 to i32
689+
%conv2 = zext i16 %load2 to i32
690+
%ins.0 = insertelement <3 x i32> poison, i32 %load0, i32 0
691+
%ins.1 = insertelement <3 x i32> %ins.0, i32 %conv1, i32 1
692+
%ins.2 = insertelement <3 x i32> %ins.1, i32 %conv2, i32 2
693+
store <3 x i32> %ins.2, ptr addrspace(1) %out
694+
ret void
695+
}
696+
658697
attributes #0 = { "amdgpu-no-agpr" "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" }

‎llvm/test/CodeGen/AMDGPU/preload-kernargs-IR-lowering.ll‎

Lines changed: 2 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -187,17 +187,15 @@ define amdgpu_kernel void @test_preload_IR_lowering_kernel_8(ptr addrspace(1) %i
187187
; PRELOAD-8-LABEL: define {{[^@]+}}@test_preload_IR_lowering_kernel_8
188188
; PRELOAD-8-SAME: (ptr addrspace(1) inreg [[IN:%.*]], ptr addrspace(1) inreg [[IN1:%.*]], ptr addrspace(1) inreg [[IN2:%.*]], ptr addrspace(1) inreg [[IN3:%.*]], ptr addrspace(1) inreg [[OUT:%.*]], ptr addrspace(1) inreg [[OUT1:%.*]], ptr addrspace(1) inreg [[OUT2:%.*]], ptr addrspace(1) inreg [[OUT3:%.*]]) #[[ATTR0]] {
189189
; PRELOAD-8-NEXT: [[TEST_PRELOAD_IR_LOWERING_KERNEL_8_KERNARG_SEGMENT:%.*]] = call nonnull align 16 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.kernarg.segment.ptr()
190-
; PRELOAD-8-NEXT: [[OUT2_KERNARG_OFFSET:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[TEST_PRELOAD_IR_LOWERING_KERNEL_8_KERNARG_SEGMENT]], i64 48
191-
; PRELOAD-8-NEXT: [[OUT2_LOAD:%.*]] = load ptr addrspace(1), ptr addrspace(4) [[OUT2_KERNARG_OFFSET]], align 16, !invariant.load [[META0:![0-9]+]]
192190
; PRELOAD-8-NEXT: [[OUT3_KERNARG_OFFSET:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[TEST_PRELOAD_IR_LOWERING_KERNEL_8_KERNARG_SEGMENT]], i64 56
193-
; PRELOAD-8-NEXT: [[OUT3_LOAD:%.*]] = load ptr addrspace(1), ptr addrspace(4) [[OUT3_KERNARG_OFFSET]], align 8, !invariant.load [[META0]]
191+
; PRELOAD-8-NEXT: [[OUT3_LOAD:%.*]] = load ptr addrspace(1), ptr addrspace(4) [[OUT3_KERNARG_OFFSET]], align 8, !invariant.load [[META0:![0-9]+]]
194192
; PRELOAD-8-NEXT: [[LOAD:%.*]] = load i32, ptr addrspace(1) [[IN]], align 4
195193
; PRELOAD-8-NEXT: [[LOAD1:%.*]] = load i32, ptr addrspace(1) [[IN1]], align 4
196194
; PRELOAD-8-NEXT: [[LOAD2:%.*]] = load i32, ptr addrspace(1) [[IN2]], align 4
197195
; PRELOAD-8-NEXT: [[LOAD3:%.*]] = load i32, ptr addrspace(1) [[IN3]], align 4
198196
; PRELOAD-8-NEXT: store i32 [[LOAD]], ptr addrspace(1) [[OUT]], align 4
199197
; PRELOAD-8-NEXT: store i32 [[LOAD1]], ptr addrspace(1) [[OUT1]], align 4
200-
; PRELOAD-8-NEXT: store i32 [[LOAD2]], ptr addrspace(1) [[OUT2_LOAD]], align 4
198+
; PRELOAD-8-NEXT: store i32 [[LOAD2]], ptr addrspace(1) [[OUT2]], align 4
201199
; PRELOAD-8-NEXT: store i32 [[LOAD3]], ptr addrspace(1) [[OUT3_LOAD]], align 4
202200
; PRELOAD-8-NEXT: ret void
203201
;

0 commit comments

Comments
 (0)