[llvm] [AMDGPU][GlobalISel] Support lowering preloaded kernel arguments (PR #205049)
Keshav Vinayak Jha via llvm-commits
llvm-commits at lists.llvm.org
Tue Jun 23 10:09:47 PDT 2026
https://github.com/keshavvinayak01 updated https://github.com/llvm/llvm-project/pull/205049
>From ae654b6aa6ce13d6dd7dc0946d655543aad4b844 Mon Sep 17 00:00:00 2001
From: Keshav Vinayak Jha <keshavvinayakjha at gmail.com>
Date: Mon, 22 Jun 2026 13:23:45 +0530
Subject: [PATCH 1/2] Handle implicit hidden argument preloading for kernels
Signed-off-by: Keshav Vinayak Jha <keshavvinayakjha at gmail.com>
---
llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp | 346 +++++++++++++++---
llvm/lib/Target/AMDGPU/AMDGPUCallLowering.h | 5 +-
.../AMDGPU/GlobalISel/preload-kernargs.ll | 100 +++++
.../CodeGen/AMDGPU/llvm.amdgcn.cvt.sat.pk.ll | 16 +-
4 files changed, 406 insertions(+), 61 deletions(-)
create mode 100644 llvm/test/CodeGen/AMDGPU/GlobalISel/preload-kernargs.ll
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp b/llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp
index ca541c234ba0d..b3331d76df7f0 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp
@@ -20,6 +20,7 @@
#include "llvm/CodeGen/Analysis.h"
#include "llvm/CodeGen/FunctionLoweringInfo.h"
#include "llvm/CodeGen/GlobalISel/MachineIRBuilder.h"
+#include "llvm/CodeGen/GlobalISel/Utils.h"
#include "llvm/CodeGen/MachineFrameInfo.h"
#include "llvm/CodeGen/PseudoSourceValueManager.h"
#include "llvm/IR/IntrinsicsAMDGPU.h"
@@ -296,6 +297,198 @@ struct AMDGPUOutgoingArgHandler : public AMDGPUOutgoingValueHandler {
assignValueToAddress(ValVReg, Addr, MemTy, MPO, VA);
}
};
+
+// Return the virtual register and type for a preloaded physical SGPR live-in.
+static std::pair<Register, LLT> getPreloadedLiveIn(MachineIRBuilder &B,
+ MCRegister PhysReg,
+ const SIRegisterInfo &TRI) {
+ MachineFunction &MF = B.getMF();
+ const TargetInstrInfo *TII = MF.getSubtarget().getInstrInfo();
+ const TargetRegisterClass *RC = TRI.getPhysRegBaseClass(PhysReg);
+ assert(RC && "expected a register class for preload SGPR");
+ LLT RegTy = LLT::scalar(TRI.getRegSizeInBits(*RC).getFixedValue());
+ return {getFunctionLiveInPhysReg(MF, *TII, PhysReg, *RC, B.getDL(), RegTy),
+ RegTy};
+}
+
+// Convert raw SGPR preload bits to the type expected by the formal argument.
+static Register adjustPreloadedArgType(MachineIRBuilder &B, Register Src,
+ LLT SrcTy, LLT DstTy) {
+ if (SrcTy == DstTy)
+ return Src;
+
+ LLT IntDstTy = DstTy.isPointer() ? LLT::scalar(DstTy.getSizeInBits()) : DstTy;
+ if (SrcTy.getSizeInBits() != IntDstTy.getSizeInBits())
+ Src = B.buildAnyExtOrTrunc(IntDstTy, Src).getReg(0);
+ else if (SrcTy != IntDstTy)
+ Src = B.buildBitcast(IntDstTy, Src).getReg(0);
+
+ if (DstTy.isPointer())
+ return B.buildIntToPtr(DstTy, Src).getReg(0);
+
+ return Src;
+}
+
+// Materialize a preloaded kernarg from its assigned SGPRs into DstReg.
+static void lowerPreloadedKernArg(MachineIRBuilder &B, Register DstReg,
+ LLT DstTy, uint64_t Offset, Align Alignment,
+ ArrayRef<MCRegister> PreloadRegs,
+ const SIRegisterInfo &TRI) {
+ assert(!PreloadRegs.empty());
+
+ // Kernarg preloads are assigned in 32-bit SGPR units.
+ const LLT S32 = LLT::scalar(32);
+
+ Register Value;
+ LLT ValueTy;
+ const bool IsPackedSubDword = DstTy.getSizeInBits() < 32 && Alignment < 4;
+
+ if (PreloadRegs.size() == 1) {
+ auto [LiveIn, LiveInTy] = getPreloadedLiveIn(B, PreloadRegs[0], TRI);
+ if (IsPackedSubDword) {
+ // Extract sub-dword preloads from their containing 32-bit SGPR word.
+ Register Raw = B.buildCopy(S32, LiveIn).getReg(0);
+ uint64_t OffsetDiff = Offset - alignDown(Offset, 4);
+ Register ShiftAmt = B.buildConstant(S32, OffsetDiff * 8).getReg(0);
+ Value = B.buildLShr(S32, Raw, ShiftAmt).getReg(0);
+ ValueTy = S32;
+ } else {
+ // A single preload entry may be a 32-bit SGPR or a wider SGPR tuple.
+ ValueTy = LiveInTy;
+ Value = B.buildCopy(ValueTy, LiveIn).getReg(0);
+ }
+ } else {
+ assert(!IsPackedSubDword && "packed sub-dword preload should use one SGPR");
+ // Reconstruct wider preloads from separate 32-bit SGPR entries.
+ SmallVector<Register, 4> Regs;
+ Regs.reserve(PreloadRegs.size());
+ for (MCRegister Reg : PreloadRegs) {
+ Register LiveIn = getPreloadedLiveIn(B, Reg, TRI).first;
+ Regs.push_back(B.buildCopy(S32, LiveIn).getReg(0));
+ }
+
+ ValueTy = LLT::scalar(PreloadRegs.size() * 32);
+ Value = B.buildMergeLikeInstr(ValueTy, Regs).getReg(0);
+ }
+
+ Register Adjusted = adjustPreloadedArgType(B, Value, ValueTy, DstTy);
+ B.buildCopy(DstReg, Adjusted);
+}
+
+// CCValAssign records an MVT, so round odd-sized split argument parts to the
+// simple memory type used for kernarg layout.
+static MVT getLocVTForSplitArg(const TargetLowering &TLI, const DataLayout &DL,
+ Type *Ty, LLVMContext &Ctx) {
+ EVT LocVT = TLI.getValueType(DL, Ty, true);
+
+ if (LocVT.isVector() && LocVT.getVectorNumElements() == 1)
+ LocVT = LocVT.getScalarType();
+
+ if (LocVT.isVector() && !LocVT.isPow2VectorType())
+ LocVT = LocVT.getPow2VectorType(Ctx);
+ else if (!LocVT.isSimple() && !LocVT.isVector())
+ LocVT = LocVT.getRoundIntegerType(Ctx);
+
+ assert(LocVT.isSimple());
+ return LocVT.getSimpleVT();
+}
+
+// Assign SGPRs for the contiguous sequence of kernel argument parts marked for
+// kernarg preload.
+static void allocatePreloadKernArgSGPRs(
+ CCState &CCInfo, ArrayRef<CallLowering::ArgInfo> SplitArgs,
+ ArrayRef<CCValAssign> ArgLocs, MachineFunction &MF,
+ const GCNSubtarget &Subtarget, const SIRegisterInfo &TRI,
+ SIMachineFunctionInfo &Info) {
+ assert(SplitArgs.size() == ArgLocs.size());
+
+ const Function &F = MF.getFunction();
+ const DataLayout &DL = F.getDataLayout();
+ unsigned LastExplicitArgOffset = Subtarget.getExplicitKernelArgOffset();
+ GCNUserSGPRUsageInfo &SGPRInfo = Info.getUserSGPRInfo();
+ bool InPreloadSequence = true;
+ unsigned InIdx = 0;
+ bool AlignedForImplicitArgs = false;
+ unsigned ImplicitArgOffset = 0;
+
+ for (const Argument &Arg : F.args()) {
+ // Preload assignment follows the original argument order. Each original
+ // argument may already have been split into one or more lowered parts.
+ const bool IsByRef = Arg.hasByRefAttr();
+ Type *ArgTy = IsByRef ? Arg.getParamByRefType() : Arg.getType();
+ if (DL.getTypeAllocSize(ArgTy) == 0)
+ continue;
+
+ // Hardware preloads a contiguous prefix of the kernarg segment. Once a
+ // non-preloaded argument is reached, no later argument can be preloaded.
+ if (!InPreloadSequence || !Arg.hasInRegAttr())
+ break;
+
+ unsigned ArgIdx = Arg.getArgNo();
+ if (InIdx < SplitArgs.size() && SplitArgs[InIdx].OrigArgIndex != ArgIdx)
+ break;
+
+ for (; InIdx < SplitArgs.size() && SplitArgs[InIdx].OrigArgIndex == ArgIdx;
+ ++InIdx) {
+ const CCValAssign &ArgLoc = ArgLocs[InIdx];
+ assert(ArgLoc.isMemLoc());
+ const Align KernelArgBaseAlign = Align(16);
+ unsigned ArgOffset = ArgLoc.getLocMemOffset();
+ Align Alignment = commonAlignment(KernelArgBaseAlign, ArgOffset);
+ unsigned NumAllocSGPRs =
+ alignTo(ArgLoc.getLocVT().getFixedSizeInBits(), 32) / 32;
+
+ if (Arg.hasAttribute("amdgpu-hidden-argument")) {
+ // Hidden arguments are laid out after the explicit kernargs, aligned to
+ // the ABI-required implicit argument block boundary.
+ if (!AlignedForImplicitArgs) {
+ ImplicitArgOffset =
+ alignTo(LastExplicitArgOffset,
+ Subtarget.getAlignmentForImplicitArgPtr()) -
+ LastExplicitArgOffset;
+ AlignedForImplicitArgs = true;
+ }
+ ArgOffset += ImplicitArgOffset;
+ }
+
+ if (ArgLoc.getLocVT().getStoreSize() < 4 && Alignment < 4) {
+ // Packed sub-dword values share the previous 32-bit SGPR. The extract
+ // happens later when materializing the argument value.
+ assert(InIdx >= 1 && "No previous SGPR");
+ Info.getArgInfo().PreloadKernArgs[InIdx].Regs.push_back(
+ Info.getArgInfo().PreloadKernArgs[InIdx - 1].Regs[0]);
+ continue;
+ }
+
+ // Padding between the last preloaded byte and this argument still
+ // consumes preload SGPR slots because the hardware preload sequence is
+ // contiguous.
+ unsigned Padding = ArgOffset - LastExplicitArgOffset;
+ unsigned PaddingSGPRs = alignTo(Padding, 4) / 4;
+ if (PaddingSGPRs + NumAllocSGPRs > SGPRInfo.getNumFreeUserSGPRs()) {
+ InPreloadSequence = false;
+ break;
+ }
+
+ const TargetRegisterClass *RC =
+ TRI.getSGPRClassForBitWidth(NumAllocSGPRs * 32);
+ SmallVectorImpl<MCRegister> *PreloadRegs =
+ Info.addPreloadedKernArg(TRI, RC, NumAllocSGPRs, InIdx, PaddingSGPRs);
+
+ // Record each assigned physical SGPR as a live-in and reserve it from the
+ // calling convention allocator.
+ if (PreloadRegs->size() > 1)
+ RC = &AMDGPU::SGPR_32RegClass;
+ for (MCRegister Reg : *PreloadRegs) {
+ assert(Reg);
+ MF.addLiveIn(Reg, RC);
+ CCInfo.AllocateReg(Reg);
+ }
+
+ LastExplicitArgOffset = NumAllocSGPRs * 4 + ArgOffset;
+ }
+ }
+}
} // anonymous namespace
AMDGPUCallLowering::AMDGPUCallLowering(const TargetLowering &TLI)
@@ -451,46 +644,59 @@ void AMDGPUCallLowering::lowerParameterPtr(Register DstReg, MachineIRBuilder &B,
B.buildPtrAdd(DstReg, KernArgSegmentVReg, OffsetReg);
}
-void AMDGPUCallLowering::lowerParameter(MachineIRBuilder &B, ArgInfo &OrigArg,
- uint64_t Offset,
- Align Alignment) const {
+bool AMDGPUCallLowering::lowerParameter(MachineIRBuilder &B, ArgInfo &Arg,
+ const CCValAssign &ArgLoc,
+ Align Alignment,
+ unsigned InputArgIndex) const {
MachineFunction &MF = B.getMF();
const Function &F = MF.getFunction();
const DataLayout &DL = F.getDataLayout();
+ const Argument *IRArg = dyn_cast_if_present<Argument>(Arg.OrigValue);
+ const bool IsHiddenArg =
+ IRArg && IRArg->hasAttribute("amdgpu-hidden-argument");
+ const SIMachineFunctionInfo *Info = MF.getInfo<SIMachineFunctionInfo>();
+ const GCNSubtarget &Subtarget = MF.getSubtarget<GCNSubtarget>();
+ const SIRegisterInfo *TRI = Subtarget.getRegisterInfo();
const SITargetLowering &TLI = *getTLI<SITargetLowering>();
MachinePointerInfo PtrInfo = TLI.getKernargSegmentPtrInfo(MF);
LLT PtrTy = LLT::pointer(AMDGPUAS::CONSTANT_ADDRESS, 64);
+ LLT ArgTy = getLLTForType(*Arg.Ty, DL);
+ if (Arg.Flags[0].isPointer()) {
+ // Compensate for losing pointeriness in splitValueTypes.
+ LLT PtrTy = LLT::pointer(Arg.Flags[0].getPointerAddrSpace(),
+ ArgTy.getScalarSizeInBits());
+ ArgTy =
+ ArgTy.isVector() ? LLT::vector(ArgTy.getElementCount(), PtrTy) : PtrTy;
+ }
+
+ assert(Arg.Regs.size() == 1);
+
+ uint64_t Offset = ArgLoc.getLocMemOffset();
+ auto PreloadArg = Info->getArgInfo().PreloadKernArgs.find(InputArgIndex);
+ if (PreloadArg != Info->getArgInfo().PreloadKernArgs.end()) {
+ lowerPreloadedKernArg(B, Arg.Regs[0], ArgTy, Offset, Alignment,
+ PreloadArg->getSecond().Regs, *TRI);
+ return true;
+ }
- SmallVector<ArgInfo, 32> SplitArgs;
- SmallVector<TypeSize> FieldOffsets;
- splitToValueTypes(OrigArg, SplitArgs, DL, F.getCallingConv(), &FieldOffsets);
-
- unsigned Idx = 0;
- for (ArgInfo &SplitArg : SplitArgs) {
- Register PtrReg = B.getMRI()->createGenericVirtualRegister(PtrTy);
- lowerParameterPtr(PtrReg, B, Offset + FieldOffsets[Idx]);
-
- LLT ArgTy = getLLTForType(*SplitArg.Ty, DL);
- if (SplitArg.Flags[0].isPointer()) {
- // Compensate for losing pointeriness in splitValueTypes.
- LLT PtrTy = LLT::pointer(SplitArg.Flags[0].getPointerAddrSpace(),
- ArgTy.getScalarSizeInBits());
- ArgTy = ArgTy.isVector() ? LLT::vector(ArgTy.getElementCount(), PtrTy)
- : PtrTy;
- }
+ if (IsHiddenArg) {
+ F.getContext().diagnose(DiagnosticInfoUnsupported(
+ F, "hidden argument in kernel signature was not preloaded"));
+ return false;
+ }
- MachineMemOperand *MMO = MF.getMachineMemOperand(
- PtrInfo,
- MachineMemOperand::MOLoad | MachineMemOperand::MODereferenceable |
- MachineMemOperand::MOInvariant,
- ArgTy, commonAlignment(Alignment, FieldOffsets[Idx]));
+ Register PtrReg = B.getMRI()->createGenericVirtualRegister(PtrTy);
+ lowerParameterPtr(PtrReg, B, Offset);
- assert(SplitArg.Regs.size() == 1);
+ MachineMemOperand *MMO = MF.getMachineMemOperand(
+ PtrInfo,
+ MachineMemOperand::MOLoad | MachineMemOperand::MODereferenceable |
+ MachineMemOperand::MOInvariant,
+ ArgTy, Alignment);
- B.buildLoad(SplitArg.Regs[0], PtrReg, *MMO);
- ++Idx;
- }
+ B.buildLoad(Arg.Regs[0], PtrReg, *MMO);
+ return true;
}
// Allocate special inputs passed in user SGPRs.
@@ -568,19 +774,14 @@ bool AMDGPUCallLowering::lowerFormalArgumentsKernel(
allocateHSAUserSGPRs(CCInfo, B, MF, *TRI, *Info);
- unsigned i = 0;
- const Align KernArgBaseAlign(16);
+ SmallVector<ArgInfo, 16> SplitArgs;
+ SmallVector<Align, 16> SplitArgAlignments;
+
+ unsigned VRegIdx = 0;
const unsigned BaseOffset = Subtarget->getExplicitKernelArgOffset();
uint64_t ExplicitArgOffset = 0;
- // TODO: Align down to dword alignment and extract bits for extending loads.
- for (auto &Arg : F.args()) {
- // TODO: Add support for kernarg preload.
- if (Arg.hasAttribute("amdgpu-hidden-argument")) {
- LLVM_DEBUG(dbgs() << "Preloading hidden arguments is not supported\n");
- return false;
- }
-
+ for (const Argument &Arg : F.args()) {
const bool IsByRef = Arg.hasByRefAttr();
Type *ArgTy = IsByRef ? Arg.getParamByRefType() : Arg.getType();
unsigned AllocSize = DL.getTypeAllocSize(ArgTy);
@@ -592,33 +793,84 @@ bool AMDGPUCallLowering::lowerFormalArgumentsKernel(
uint64_t ArgOffset = alignTo(ExplicitArgOffset, ABIAlign) + BaseOffset;
ExplicitArgOffset = alignTo(ExplicitArgOffset, ABIAlign) + AllocSize;
+ Align ArgBaseAlign = commonAlignment(Align(16), ArgOffset);
+
+ ArgInfo OrigArg(VRegs[VRegIdx], Arg, Arg.getArgNo());
+ const unsigned OrigArgIdx = Arg.getArgNo() + AttributeList::FirstArgIndex;
+ setArgFlags(OrigArg, OrigArgIdx, DL, F);
+
+ SmallVector<ArgInfo, 32> ArgSplitArgs;
+ SmallVector<TypeSize> FieldOffsets;
+ splitToValueTypes(OrigArg, ArgSplitArgs, DL, F.getCallingConv(),
+ &FieldOffsets);
+
+ for (unsigned SplitIdx = 0, NumSplitArgs = ArgSplitArgs.size();
+ SplitIdx != NumSplitArgs; ++SplitIdx) {
+ unsigned InputArgIndex = SplitArgs.size();
+ ArgSplitArgs[SplitIdx].OrigValue = &Arg;
+ MVT LocVT = getLocVTForSplitArg(TLI, DL, ArgSplitArgs[SplitIdx].Ty,
+ F.getContext());
+ CCInfo.addLoc(CCValAssign::getCustomMem(
+ InputArgIndex, LocVT,
+ ArgOffset + FieldOffsets[SplitIdx].getFixedValue(), LocVT,
+ CCValAssign::Full));
+ SplitArgAlignments.push_back(
+ commonAlignment(ArgBaseAlign, FieldOffsets[SplitIdx]));
+ SplitArgs.push_back(ArgSplitArgs[SplitIdx]);
+ }
+
+ ++VRegIdx;
+ }
+
+ if (Subtarget->hasKernargPreload())
+ allocatePreloadKernArgSGPRs(CCInfo, SplitArgs, ArgLocs, MF, *Subtarget,
+ *TRI, *Info);
+
+ unsigned i = 0;
+ unsigned InputArgIndex = 0;
+
+ // TODO: Align down to dword alignment and extract bits for extending loads.
+ for (auto &Arg : F.args()) {
+ const bool IsByRef = Arg.hasByRefAttr();
+ Type *ArgTy = IsByRef ? Arg.getParamByRefType() : Arg.getType();
+ unsigned AllocSize = DL.getTypeAllocSize(ArgTy);
+ if (AllocSize == 0)
+ continue;
+
+ unsigned FirstInputArgIndex = InputArgIndex;
+ while (InputArgIndex < SplitArgs.size() &&
+ SplitArgs[InputArgIndex].OrigArgIndex == Arg.getArgNo())
+ ++InputArgIndex;
+ unsigned NumInputParts = InputArgIndex - FirstInputArgIndex;
if (Arg.use_empty()) {
++i;
continue;
}
- Align Alignment = commonAlignment(KernArgBaseAlign, ArgOffset);
-
if (IsByRef) {
unsigned ByRefAS = cast<PointerType>(Arg.getType())->getAddressSpace();
assert(VRegs[i].size() == 1 &&
"expected only one register for byval pointers");
if (ByRefAS == AMDGPUAS::CONSTANT_ADDRESS) {
- lowerParameterPtr(VRegs[i][0], B, ArgOffset);
+ lowerParameterPtr(VRegs[i][0], B,
+ ArgLocs[FirstInputArgIndex].getLocMemOffset());
} else {
const LLT ConstPtrTy = LLT::pointer(AMDGPUAS::CONSTANT_ADDRESS, 64);
Register PtrReg = MRI.createGenericVirtualRegister(ConstPtrTy);
- lowerParameterPtr(PtrReg, B, ArgOffset);
+ lowerParameterPtr(PtrReg, B,
+ ArgLocs[FirstInputArgIndex].getLocMemOffset());
B.buildAddrSpaceCast(VRegs[i][0], PtrReg);
}
} else {
- ArgInfo OrigArg(VRegs[i], Arg, i);
- const unsigned OrigArgIdx = i + AttributeList::FirstArgIndex;
- setArgFlags(OrigArg, OrigArgIdx, DL, F);
- lowerParameter(B, OrigArg, ArgOffset, Alignment);
+ for (unsigned Part = 0; Part != NumInputParts; ++Part) {
+ unsigned InputArgIndex = FirstInputArgIndex + Part;
+ if (!lowerParameter(B, SplitArgs[InputArgIndex], ArgLocs[InputArgIndex],
+ SplitArgAlignments[InputArgIndex], InputArgIndex))
+ return false;
+ }
}
++i;
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUCallLowering.h b/llvm/lib/Target/AMDGPU/AMDGPUCallLowering.h
index 239300c1469b9..b775db64554f7 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUCallLowering.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPUCallLowering.h
@@ -26,8 +26,9 @@ class AMDGPUCallLowering final : public CallLowering {
void lowerParameterPtr(Register DstReg, MachineIRBuilder &B,
uint64_t Offset) const;
- void lowerParameter(MachineIRBuilder &B, ArgInfo &AI, uint64_t Offset,
- Align Alignment) const;
+ bool lowerParameter(MachineIRBuilder &B, ArgInfo &AI,
+ const CCValAssign &ArgLoc, Align Alignment,
+ unsigned InputArgIndex) const;
bool canLowerReturn(MachineFunction &MF, CallingConv::ID CallConv,
SmallVectorImpl<BaseArgInfo> &Outs,
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/preload-kernargs.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/preload-kernargs.ll
new file mode 100644
index 0000000000000..18ba715537465
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/preload-kernargs.ll
@@ -0,0 +1,100 @@
+; RUN: llc -verify-machineinstrs -global-isel=1 -global-isel-abort=2 -mtriple=amdgcn--amdhsa -mcpu=gfx90a -stop-after=irtranslator -o - < %s | FileCheck -check-prefix=MIR %s
+; RUN: llc -verify-machineinstrs -global-isel=1 -global-isel-abort=2 -mtriple=amdgcn--amdhsa -mcpu=gfx90a -o - < %s | FileCheck -check-prefix=ASM %s
+; RUN: llc -verify-machineinstrs -O1 -global-isel=1 -global-isel-abort=2 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a -o - < %s | FileCheck -check-prefix=ISSUE %s
+
+define amdgpu_kernel void @explicit_i32_inreg(i32 inreg %x,
+ ptr addrspace(1) %out) #1 {
+; MIR-LABEL: name: explicit_i32_inreg
+; MIR: firstKernArgPreloadReg
+; MIR: numKernargPreloadSGPRs: 1
+; MIR: body:
+; MIR: [[X:%[0-9]+]]:_(s32) = COPY %{{[0-9]+}}
+; MIR-NOT: [[X]]:_(s32) = G_LOAD
+; ASM-LABEL: explicit_i32_inreg:
+; ASM: .amdhsa_user_sgpr_kernarg_preload_length 1
+ store i32 %x, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+define amdgpu_kernel void @explicit_i64_inreg(i64 inreg %x,
+ ptr addrspace(1) %out) #1 {
+; MIR-LABEL: name: explicit_i64_inreg
+; MIR: firstKernArgPreloadReg
+; MIR: numKernargPreloadSGPRs: 2
+; MIR: body:
+; MIR: [[X:%[0-9]+]]:_(s64) = COPY %{{[0-9]+}}
+; MIR-NOT: [[X]]:_(s64) = G_LOAD
+; ASM-LABEL: explicit_i64_inreg:
+; ASM: .amdhsa_user_sgpr_kernarg_preload_length 2
+ store i64 %x, ptr addrspace(1) %out, align 8
+ ret void
+}
+
+define amdgpu_kernel void @explicit_ptr_inreg(ptr addrspace(1) inreg %p,
+ ptr addrspace(1) %out) #1 {
+; MIR-LABEL: name: explicit_ptr_inreg
+; MIR: firstKernArgPreloadReg
+; MIR: numKernargPreloadSGPRs: 2
+; MIR: body:
+; MIR: [[PINT:%[0-9]+]]:_(s64) = COPY %{{[0-9]+}}
+; MIR: [[P:%[0-9]+]]:_(p1) = G_INTTOPTR [[PINT]](s64)
+; MIR-NOT: [[P]]:_(p1) = G_LOAD
+; ASM-LABEL: explicit_ptr_inreg:
+; ASM: .amdhsa_user_sgpr_kernarg_preload_length 2
+ store ptr addrspace(1) %p, ptr addrspace(1) %out, align 8
+ ret void
+}
+
+define amdgpu_kernel void @packed_i16_inreg(i16 inreg %a, i16 inreg %b,
+ ptr addrspace(1) %out) #1 {
+; MIR-LABEL: name: packed_i16_inreg
+; MIR: firstKernArgPreloadReg
+; MIR: numKernargPreloadSGPRs: 1
+; MIR: body:
+; MIR: [[PACKED:%[0-9]+]]:_(s32) = COPY %{{[0-9]+}}
+; MIR: [[SHIFT:%[0-9]+]]:_(s32) = G_LSHR [[PACKED]], %{{[0-9]+}}(s32)
+; MIR: [[B:%[0-9]+]]:_(s16) = G_TRUNC [[SHIFT]](s32)
+; MIR-NOT: [[B]]:_(s16) = G_LOAD
+; ASM-LABEL: packed_i16_inreg:
+; ASM: .amdhsa_user_sgpr_kernarg_preload_length 1
+ store i16 %b, ptr addrspace(1) %out, align 2
+ ret void
+}
+
+define amdgpu_kernel void @preload_block_count_x(
+ ptr addrspace(1) inreg %out,
+ i32 inreg "amdgpu-hidden-argument" %_hidden_block_count_x) #0 {
+; MIR-LABEL: name: preload_block_count_x
+; MIR: firstKernArgPreloadReg
+; MIR: numKernargPreloadSGPRs: 3
+; MIR: body:
+; MIR: {{%[0-9]+}}:_(p1) = G_INTTOPTR {{%[0-9]+}}(s64)
+; MIR: {{%[0-9]+}}:_(s32) = COPY %{{[0-9]+}}
+; ASM-LABEL: preload_block_count_x:
+; ASM: .amdhsa_user_sgpr_kernarg_preload_length 3
+ store i32 %_hidden_block_count_x, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+define amdgpu_kernel void @_Z15hipsmith_kernelv() local_unnamed_addr #0 {
+entry:
+; ISSUE-LABEL: _Z15hipsmith_kernelv:
+; ISSUE: .amdhsa_user_sgpr_kernarg_preload_length 1
+ %0 = tail call dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
+ %1 = load i32, ptr addrspace(4) %0, align 4
+ %.not.i.not = icmp eq i32 %1, 0
+ br i1 %.not.i.not, label %lor.end.i, label %common.ret1
+
+common.ret1:
+ ret void
+
+lor.end.i:
+ store ptr null, ptr inttoptr (i64 32 to ptr), align 32
+ br label %common.ret1
+}
+
+declare noundef align 4 ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr() #2
+
+attributes #0 = { "amdgpu-agpr-alloc"="0" "amdgpu-no-cluster-id-x" "amdgpu-no-cluster-id-y" "amdgpu-no-cluster-id-z" "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-lds-kernel-id" "amdgpu-no-multigrid-sync-arg" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "amdgpu-no-wwm" }
+attributes #1 = { "target-cpu"="gfx90a" }
+attributes #2 = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.sat.pk.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.sat.pk.ll
index 08dccdf5872d0..4db018812a8fe 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.sat.pk.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.sat.pk.ll
@@ -84,24 +84,20 @@ define amdgpu_kernel void @sat_pk4_i4_i8_f32_s(i32 inreg %src, ptr %out) #1 {
; GISEL-REAL16-LABEL: sat_pk4_i4_i8_f32_s:
; GISEL-REAL16: ; %bb.0:
; GISEL-REAL16-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
-; GISEL-REAL16-NEXT: s_clause 0x1
-; GISEL-REAL16-NEXT: s_load_b32 s2, s[4:5], 0x0 nv
; GISEL-REAL16-NEXT: s_load_b64 s[0:1], s[4:5], 0x8 nv
+; GISEL-REAL16-NEXT: v_sat_pk4_i4_i8_e32 v0.l, s8
; GISEL-REAL16-NEXT: v_mov_b32_e32 v1, 0
; GISEL-REAL16-NEXT: s_wait_kmcnt 0x0
-; GISEL-REAL16-NEXT: v_sat_pk4_i4_i8_e32 v0.l, s2
; GISEL-REAL16-NEXT: flat_store_b16 v1, v0, s[0:1]
; GISEL-REAL16-NEXT: s_endpgm
;
; GISEL-FAKE16-LABEL: sat_pk4_i4_i8_f32_s:
; GISEL-FAKE16: ; %bb.0:
; GISEL-FAKE16-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
-; GISEL-FAKE16-NEXT: s_clause 0x1
-; GISEL-FAKE16-NEXT: s_load_b32 s2, s[4:5], 0x0 nv
; GISEL-FAKE16-NEXT: s_load_b64 s[0:1], s[4:5], 0x8 nv
+; GISEL-FAKE16-NEXT: v_sat_pk4_i4_i8_e32 v0, s8
; GISEL-FAKE16-NEXT: v_mov_b32_e32 v1, 0
; GISEL-FAKE16-NEXT: s_wait_kmcnt 0x0
-; GISEL-FAKE16-NEXT: v_sat_pk4_i4_i8_e32 v0, s2
; GISEL-FAKE16-NEXT: flat_store_b16 v1, v0, s[0:1]
; GISEL-FAKE16-NEXT: s_endpgm
%cvt = call i16 @llvm.amdgcn.sat.pk4.i4.i8(i32 %src) #0
@@ -231,24 +227,20 @@ define amdgpu_kernel void @sat_pk4_u4_u8_f32_s(i32 inreg %src, ptr %out) #1 {
; GISEL-REAL16-LABEL: sat_pk4_u4_u8_f32_s:
; GISEL-REAL16: ; %bb.0:
; GISEL-REAL16-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
-; GISEL-REAL16-NEXT: s_clause 0x1
-; GISEL-REAL16-NEXT: s_load_b32 s2, s[4:5], 0x0 nv
; GISEL-REAL16-NEXT: s_load_b64 s[0:1], s[4:5], 0x8 nv
+; GISEL-REAL16-NEXT: v_sat_pk4_u4_u8_e32 v0.l, s8
; GISEL-REAL16-NEXT: v_mov_b32_e32 v1, 0
; GISEL-REAL16-NEXT: s_wait_kmcnt 0x0
-; GISEL-REAL16-NEXT: v_sat_pk4_u4_u8_e32 v0.l, s2
; GISEL-REAL16-NEXT: flat_store_b16 v1, v0, s[0:1]
; GISEL-REAL16-NEXT: s_endpgm
;
; GISEL-FAKE16-LABEL: sat_pk4_u4_u8_f32_s:
; GISEL-FAKE16: ; %bb.0:
; GISEL-FAKE16-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
-; GISEL-FAKE16-NEXT: s_clause 0x1
-; GISEL-FAKE16-NEXT: s_load_b32 s2, s[4:5], 0x0 nv
; GISEL-FAKE16-NEXT: s_load_b64 s[0:1], s[4:5], 0x8 nv
+; GISEL-FAKE16-NEXT: v_sat_pk4_u4_u8_e32 v0, s8
; GISEL-FAKE16-NEXT: v_mov_b32_e32 v1, 0
; GISEL-FAKE16-NEXT: s_wait_kmcnt 0x0
-; GISEL-FAKE16-NEXT: v_sat_pk4_u4_u8_e32 v0, s2
; GISEL-FAKE16-NEXT: flat_store_b16 v1, v0, s[0:1]
; GISEL-FAKE16-NEXT: s_endpgm
%cvt = call i16 @llvm.amdgcn.sat.pk4.u4.u8(i32 %src) #0
>From 066161fd71321532628ee61db364cf9b4a6b82d3 Mon Sep 17 00:00:00 2001
From: Keshav Vinayak Jha <keshavvinayakjha at gmail.com>
Date: Tue, 23 Jun 2026 22:33:05 +0530
Subject: [PATCH 2/2] AMDGPU: Add GlobalISel kernarg preload coverage
Add full-pipeline GlobalISel run lines to the existing kernarg preload tests and move MIR coverage into a dedicated IRTranslator test.
Also handle vector destinations when adjusting preloaded argument types.
Co-authored-by: GPT-5 Codex <noreply at openai.com>
Signed-off-by: Keshav Vinayak Jha <keshavvinayakjha at gmail.com>
---
llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp | 7 +-
.../irtranslator-preload-kernargs.ll | 104 +++
.../AMDGPU/GlobalISel/preload-kernargs.ll | 100 ---
.../AMDGPU/preload-implicit-kernargs.ll | 436 +++++++++++++
llvm/test/CodeGen/AMDGPU/preload-kernargs.ll | 593 ++++++++++++++++++
5 files changed, 1139 insertions(+), 101 deletions(-)
create mode 100644 llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-preload-kernargs.ll
delete mode 100644 llvm/test/CodeGen/AMDGPU/GlobalISel/preload-kernargs.ll
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp b/llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp
index b3331d76df7f0..92e90a516e02c 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp
@@ -317,7 +317,9 @@ static Register adjustPreloadedArgType(MachineIRBuilder &B, Register Src,
if (SrcTy == DstTy)
return Src;
- LLT IntDstTy = DstTy.isPointer() ? LLT::scalar(DstTy.getSizeInBits()) : DstTy;
+ LLT IntDstTy = (DstTy.isPointer() || DstTy.isVector())
+ ? LLT::scalar(DstTy.getSizeInBits())
+ : DstTy;
if (SrcTy.getSizeInBits() != IntDstTy.getSizeInBits())
Src = B.buildAnyExtOrTrunc(IntDstTy, Src).getReg(0);
else if (SrcTy != IntDstTy)
@@ -326,6 +328,9 @@ static Register adjustPreloadedArgType(MachineIRBuilder &B, Register Src,
if (DstTy.isPointer())
return B.buildIntToPtr(DstTy, Src).getReg(0);
+ if (DstTy.isVector())
+ return B.buildBitcast(DstTy, Src).getReg(0);
+
return Src;
}
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-preload-kernargs.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-preload-kernargs.ll
new file mode 100644
index 0000000000000..70541447e49e9
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-preload-kernargs.ll
@@ -0,0 +1,104 @@
+; NOTE: Assertions have been autogenerated by utils/update_mir_test_checks.py UTC_ARGS: --version 6
+; RUN: llc -global-isel -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a -stop-after=irtranslator -o - < %s | FileCheck -check-prefix=GFX90A %s
+
+define amdgpu_kernel void @preload_i32_inreg(i32 inreg %x, ptr addrspace(1) %out) #0 {
+ ; GFX90A-LABEL: name: preload_i32_inreg
+ ; GFX90A: bb.1 (%ir-block.0):
+ ; GFX90A-NEXT: liveins: $sgpr6, $sgpr4_sgpr5
+ ; GFX90A-NEXT: {{ $}}
+ ; GFX90A-NEXT: [[COPY:%[0-9]+]]:sreg_32 = COPY $sgpr6
+ ; GFX90A-NEXT: [[COPY1:%[0-9]+]]:_(p4) = COPY $sgpr4_sgpr5
+ ; GFX90A-NEXT: [[COPY2:%[0-9]+]]:_(s32) = COPY [[COPY]]
+ ; GFX90A-NEXT: [[COPY3:%[0-9]+]]:_(s32) = COPY [[COPY2]](s32)
+ ; GFX90A-NEXT: [[INT:%[0-9]+]]:_(p4) = G_INTRINSIC intrinsic(@llvm.amdgcn.kernarg.segment.ptr)
+ ; GFX90A-NEXT: [[C:%[0-9]+]]:_(s64) = G_CONSTANT i64 8
+ ; GFX90A-NEXT: [[PTR_ADD:%[0-9]+]]:_(p4) = nuw nusw inbounds G_PTR_ADD [[INT]], [[C]](s64)
+ ; GFX90A-NEXT: [[LOAD:%[0-9]+]]:_(p1) = G_LOAD [[PTR_ADD]](p4) :: (dereferenceable invariant load (p1) from %ir.out.kernarg.offset, addrspace 4)
+ ; GFX90A-NEXT: G_STORE [[COPY3]](s32), [[LOAD]](p1) :: (store (s32) into %ir.out.load, addrspace 1)
+ ; GFX90A-NEXT: S_ENDPGM 0
+ store i32 %x, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+define amdgpu_kernel void @preload_i64_inreg(i64 inreg %x, ptr addrspace(1) %out) #0 {
+ ; GFX90A-LABEL: name: preload_i64_inreg
+ ; GFX90A: bb.1 (%ir-block.0):
+ ; GFX90A-NEXT: liveins: $sgpr4_sgpr5, $sgpr6_sgpr7
+ ; GFX90A-NEXT: {{ $}}
+ ; GFX90A-NEXT: [[COPY:%[0-9]+]]:sreg_64 = COPY $sgpr6_sgpr7
+ ; GFX90A-NEXT: [[COPY1:%[0-9]+]]:_(p4) = COPY $sgpr4_sgpr5
+ ; GFX90A-NEXT: [[COPY2:%[0-9]+]]:_(s64) = COPY [[COPY]]
+ ; GFX90A-NEXT: [[COPY3:%[0-9]+]]:_(s64) = COPY [[COPY2]](s64)
+ ; GFX90A-NEXT: [[INT:%[0-9]+]]:_(p4) = G_INTRINSIC intrinsic(@llvm.amdgcn.kernarg.segment.ptr)
+ ; GFX90A-NEXT: [[C:%[0-9]+]]:_(s64) = G_CONSTANT i64 8
+ ; GFX90A-NEXT: [[PTR_ADD:%[0-9]+]]:_(p4) = nuw nusw inbounds G_PTR_ADD [[INT]], [[C]](s64)
+ ; GFX90A-NEXT: [[LOAD:%[0-9]+]]:_(p1) = G_LOAD [[PTR_ADD]](p4) :: (dereferenceable invariant load (p1) from %ir.out.kernarg.offset, addrspace 4)
+ ; GFX90A-NEXT: G_STORE [[COPY3]](s64), [[LOAD]](p1) :: (store (s64) into %ir.out.load, addrspace 1)
+ ; GFX90A-NEXT: S_ENDPGM 0
+ store i64 %x, ptr addrspace(1) %out, align 8
+ ret void
+}
+
+define amdgpu_kernel void @preload_ptr_inreg(ptr addrspace(1) inreg %p, ptr addrspace(1) %out) #0 {
+ ; GFX90A-LABEL: name: preload_ptr_inreg
+ ; GFX90A: bb.1 (%ir-block.0):
+ ; GFX90A-NEXT: liveins: $sgpr4_sgpr5, $sgpr6_sgpr7
+ ; GFX90A-NEXT: {{ $}}
+ ; GFX90A-NEXT: [[COPY:%[0-9]+]]:sreg_64 = COPY $sgpr6_sgpr7
+ ; GFX90A-NEXT: [[COPY1:%[0-9]+]]:_(p4) = COPY $sgpr4_sgpr5
+ ; GFX90A-NEXT: [[COPY2:%[0-9]+]]:_(s64) = COPY [[COPY]]
+ ; GFX90A-NEXT: [[INTTOPTR:%[0-9]+]]:_(p1) = G_INTTOPTR [[COPY2]](s64)
+ ; GFX90A-NEXT: [[COPY3:%[0-9]+]]:_(p1) = COPY [[INTTOPTR]](p1)
+ ; GFX90A-NEXT: [[INT:%[0-9]+]]:_(p4) = G_INTRINSIC intrinsic(@llvm.amdgcn.kernarg.segment.ptr)
+ ; GFX90A-NEXT: [[C:%[0-9]+]]:_(s64) = G_CONSTANT i64 8
+ ; GFX90A-NEXT: [[PTR_ADD:%[0-9]+]]:_(p4) = nuw nusw inbounds G_PTR_ADD [[INT]], [[C]](s64)
+ ; GFX90A-NEXT: [[LOAD:%[0-9]+]]:_(p1) = G_LOAD [[PTR_ADD]](p4) :: (dereferenceable invariant load (p1) from %ir.out.kernarg.offset, addrspace 4)
+ ; GFX90A-NEXT: G_STORE [[COPY3]](p1), [[LOAD]](p1) :: (store (p1) into %ir.out.load, addrspace 1)
+ ; GFX90A-NEXT: S_ENDPGM 0
+ store ptr addrspace(1) %p, ptr addrspace(1) %out, align 8
+ ret void
+}
+
+define amdgpu_kernel void @preload_packed_i16_inreg(i16 inreg %a, i16 inreg %b, ptr addrspace(1) %out) #0 {
+ ; GFX90A-LABEL: name: preload_packed_i16_inreg
+ ; GFX90A: bb.1 (%ir-block.0):
+ ; GFX90A-NEXT: liveins: $sgpr6, $sgpr4_sgpr5
+ ; GFX90A-NEXT: {{ $}}
+ ; GFX90A-NEXT: [[COPY:%[0-9]+]]:sreg_32 = COPY $sgpr6
+ ; GFX90A-NEXT: [[COPY1:%[0-9]+]]:_(p4) = COPY $sgpr4_sgpr5
+ ; GFX90A-NEXT: [[COPY2:%[0-9]+]]:_(s32) = COPY [[COPY]]
+ ; GFX90A-NEXT: [[C:%[0-9]+]]:_(s32) = G_CONSTANT i32 16
+ ; GFX90A-NEXT: [[LSHR:%[0-9]+]]:_(s32) = G_LSHR [[COPY2]], [[C]](s32)
+ ; GFX90A-NEXT: [[TRUNC:%[0-9]+]]:_(s16) = G_TRUNC [[LSHR]](s32)
+ ; GFX90A-NEXT: [[COPY3:%[0-9]+]]:_(s16) = COPY [[TRUNC]](s16)
+ ; GFX90A-NEXT: [[INT:%[0-9]+]]:_(p4) = G_INTRINSIC intrinsic(@llvm.amdgcn.kernarg.segment.ptr)
+ ; GFX90A-NEXT: [[C1:%[0-9]+]]:_(s64) = G_CONSTANT i64 8
+ ; GFX90A-NEXT: [[PTR_ADD:%[0-9]+]]:_(p4) = nuw nusw inbounds G_PTR_ADD [[INT]], [[C1]](s64)
+ ; GFX90A-NEXT: [[LOAD:%[0-9]+]]:_(p1) = G_LOAD [[PTR_ADD]](p4) :: (dereferenceable invariant load (p1) from %ir.out.kernarg.offset, addrspace 4)
+ ; GFX90A-NEXT: G_STORE [[COPY3]](s16), [[LOAD]](p1) :: (store (s16) into %ir.out.load, addrspace 1)
+ ; GFX90A-NEXT: S_ENDPGM 0
+ store i16 %b, ptr addrspace(1) %out, align 2
+ ret void
+}
+
+define amdgpu_kernel void @preload_hidden_block_count_x(ptr addrspace(1) inreg %out, i32 inreg "amdgpu-hidden-argument" %_hidden_block_count_x) #0 {
+ ; GFX90A-LABEL: name: preload_hidden_block_count_x
+ ; GFX90A: bb.1 (%ir-block.0):
+ ; GFX90A-NEXT: liveins: $sgpr8, $sgpr4_sgpr5, $sgpr6_sgpr7
+ ; GFX90A-NEXT: {{ $}}
+ ; GFX90A-NEXT: [[COPY:%[0-9]+]]:sreg_32 = COPY $sgpr8
+ ; GFX90A-NEXT: [[COPY1:%[0-9]+]]:sreg_64 = COPY $sgpr6_sgpr7
+ ; GFX90A-NEXT: [[COPY2:%[0-9]+]]:_(p4) = COPY $sgpr4_sgpr5
+ ; GFX90A-NEXT: [[COPY3:%[0-9]+]]:_(s64) = COPY [[COPY1]]
+ ; GFX90A-NEXT: [[INTTOPTR:%[0-9]+]]:_(p1) = G_INTTOPTR [[COPY3]](s64)
+ ; GFX90A-NEXT: [[COPY4:%[0-9]+]]:_(p1) = COPY [[INTTOPTR]](p1)
+ ; GFX90A-NEXT: [[COPY5:%[0-9]+]]:_(s32) = COPY [[COPY]]
+ ; GFX90A-NEXT: [[COPY6:%[0-9]+]]:_(s32) = COPY [[COPY5]](s32)
+ ; GFX90A-NEXT: [[INT:%[0-9]+]]:_(p4) = G_INTRINSIC intrinsic(@llvm.amdgcn.kernarg.segment.ptr)
+ ; GFX90A-NEXT: G_STORE [[COPY6]](s32), [[COPY4]](p1) :: (store (s32) into %ir.out, addrspace 1)
+ ; GFX90A-NEXT: S_ENDPGM 0
+ store i32 %_hidden_block_count_x, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+attributes #0 = { "amdgpu-agpr-alloc"="0" "amdgpu-no-cluster-id-x" "amdgpu-no-cluster-id-y" "amdgpu-no-cluster-id-z" "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-lds-kernel-id" "amdgpu-no-multigrid-sync-arg" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "amdgpu-no-wwm" "target-cpu"="gfx90a" }
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/preload-kernargs.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/preload-kernargs.ll
deleted file mode 100644
index 18ba715537465..0000000000000
--- a/llvm/test/CodeGen/AMDGPU/GlobalISel/preload-kernargs.ll
+++ /dev/null
@@ -1,100 +0,0 @@
-; RUN: llc -verify-machineinstrs -global-isel=1 -global-isel-abort=2 -mtriple=amdgcn--amdhsa -mcpu=gfx90a -stop-after=irtranslator -o - < %s | FileCheck -check-prefix=MIR %s
-; RUN: llc -verify-machineinstrs -global-isel=1 -global-isel-abort=2 -mtriple=amdgcn--amdhsa -mcpu=gfx90a -o - < %s | FileCheck -check-prefix=ASM %s
-; RUN: llc -verify-machineinstrs -O1 -global-isel=1 -global-isel-abort=2 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a -o - < %s | FileCheck -check-prefix=ISSUE %s
-
-define amdgpu_kernel void @explicit_i32_inreg(i32 inreg %x,
- ptr addrspace(1) %out) #1 {
-; MIR-LABEL: name: explicit_i32_inreg
-; MIR: firstKernArgPreloadReg
-; MIR: numKernargPreloadSGPRs: 1
-; MIR: body:
-; MIR: [[X:%[0-9]+]]:_(s32) = COPY %{{[0-9]+}}
-; MIR-NOT: [[X]]:_(s32) = G_LOAD
-; ASM-LABEL: explicit_i32_inreg:
-; ASM: .amdhsa_user_sgpr_kernarg_preload_length 1
- store i32 %x, ptr addrspace(1) %out, align 4
- ret void
-}
-
-define amdgpu_kernel void @explicit_i64_inreg(i64 inreg %x,
- ptr addrspace(1) %out) #1 {
-; MIR-LABEL: name: explicit_i64_inreg
-; MIR: firstKernArgPreloadReg
-; MIR: numKernargPreloadSGPRs: 2
-; MIR: body:
-; MIR: [[X:%[0-9]+]]:_(s64) = COPY %{{[0-9]+}}
-; MIR-NOT: [[X]]:_(s64) = G_LOAD
-; ASM-LABEL: explicit_i64_inreg:
-; ASM: .amdhsa_user_sgpr_kernarg_preload_length 2
- store i64 %x, ptr addrspace(1) %out, align 8
- ret void
-}
-
-define amdgpu_kernel void @explicit_ptr_inreg(ptr addrspace(1) inreg %p,
- ptr addrspace(1) %out) #1 {
-; MIR-LABEL: name: explicit_ptr_inreg
-; MIR: firstKernArgPreloadReg
-; MIR: numKernargPreloadSGPRs: 2
-; MIR: body:
-; MIR: [[PINT:%[0-9]+]]:_(s64) = COPY %{{[0-9]+}}
-; MIR: [[P:%[0-9]+]]:_(p1) = G_INTTOPTR [[PINT]](s64)
-; MIR-NOT: [[P]]:_(p1) = G_LOAD
-; ASM-LABEL: explicit_ptr_inreg:
-; ASM: .amdhsa_user_sgpr_kernarg_preload_length 2
- store ptr addrspace(1) %p, ptr addrspace(1) %out, align 8
- ret void
-}
-
-define amdgpu_kernel void @packed_i16_inreg(i16 inreg %a, i16 inreg %b,
- ptr addrspace(1) %out) #1 {
-; MIR-LABEL: name: packed_i16_inreg
-; MIR: firstKernArgPreloadReg
-; MIR: numKernargPreloadSGPRs: 1
-; MIR: body:
-; MIR: [[PACKED:%[0-9]+]]:_(s32) = COPY %{{[0-9]+}}
-; MIR: [[SHIFT:%[0-9]+]]:_(s32) = G_LSHR [[PACKED]], %{{[0-9]+}}(s32)
-; MIR: [[B:%[0-9]+]]:_(s16) = G_TRUNC [[SHIFT]](s32)
-; MIR-NOT: [[B]]:_(s16) = G_LOAD
-; ASM-LABEL: packed_i16_inreg:
-; ASM: .amdhsa_user_sgpr_kernarg_preload_length 1
- store i16 %b, ptr addrspace(1) %out, align 2
- ret void
-}
-
-define amdgpu_kernel void @preload_block_count_x(
- ptr addrspace(1) inreg %out,
- i32 inreg "amdgpu-hidden-argument" %_hidden_block_count_x) #0 {
-; MIR-LABEL: name: preload_block_count_x
-; MIR: firstKernArgPreloadReg
-; MIR: numKernargPreloadSGPRs: 3
-; MIR: body:
-; MIR: {{%[0-9]+}}:_(p1) = G_INTTOPTR {{%[0-9]+}}(s64)
-; MIR: {{%[0-9]+}}:_(s32) = COPY %{{[0-9]+}}
-; ASM-LABEL: preload_block_count_x:
-; ASM: .amdhsa_user_sgpr_kernarg_preload_length 3
- store i32 %_hidden_block_count_x, ptr addrspace(1) %out, align 4
- ret void
-}
-
-define amdgpu_kernel void @_Z15hipsmith_kernelv() local_unnamed_addr #0 {
-entry:
-; ISSUE-LABEL: _Z15hipsmith_kernelv:
-; ISSUE: .amdhsa_user_sgpr_kernarg_preload_length 1
- %0 = tail call dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
- %1 = load i32, ptr addrspace(4) %0, align 4
- %.not.i.not = icmp eq i32 %1, 0
- br i1 %.not.i.not, label %lor.end.i, label %common.ret1
-
-common.ret1:
- ret void
-
-lor.end.i:
- store ptr null, ptr inttoptr (i64 32 to ptr), align 32
- br label %common.ret1
-}
-
-declare noundef align 4 ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr() #2
-
-attributes #0 = { "amdgpu-agpr-alloc"="0" "amdgpu-no-cluster-id-x" "amdgpu-no-cluster-id-y" "amdgpu-no-cluster-id-z" "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-lds-kernel-id" "amdgpu-no-multigrid-sync-arg" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "amdgpu-no-wwm" }
-attributes #1 = { "target-cpu"="gfx90a" }
-attributes #2 = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
diff --git a/llvm/test/CodeGen/AMDGPU/preload-implicit-kernargs.ll b/llvm/test/CodeGen/AMDGPU/preload-implicit-kernargs.ll
index 344d1f5d0b854..bbadf150e1e14 100644
--- a/llvm/test/CodeGen/AMDGPU/preload-implicit-kernargs.ll
+++ b/llvm/test/CodeGen/AMDGPU/preload-implicit-kernargs.ll
@@ -1,6 +1,7 @@
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx942 < %s | FileCheck -check-prefixes=GFX942 %s
; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx90a < %s | FileCheck -check-prefixes=GFX90a %s
+; RUN: llc -global-isel -mtriple=amdgcn--amdhsa -mcpu=gfx90a < %s | FileCheck -check-prefixes=GISEL-GFX90A %s
; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx1250 -mattr=-real-true16 < %s | FileCheck -check-prefixes=GFX1250,GFX1250-FAKE16 %s
; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx1250 -mattr=+real-true16 < %s | FileCheck -check-prefixes=GFX1250,GFX1250-REAL16 %s
@@ -33,6 +34,20 @@ define amdgpu_kernel void @preload_block_count_x(ptr addrspace(1) inreg %out) #0
; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: preload_block_count_x:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[8:9], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_load_dword s10, s[4:5], 0x8
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB0_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB0_0:
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s10
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: global_store_dword v1, v0, s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: preload_block_count_x:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -75,6 +90,20 @@ define amdgpu_kernel void @preload_unused_arg_block_count_x(ptr addrspace(1) inr
; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: preload_unused_arg_block_count_x:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_load_dword s12, s[4:5], 0x10
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB1_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB1_0:
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s12
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: global_store_dword v1, v0, s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: preload_unused_arg_block_count_x:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -118,6 +147,21 @@ define amdgpu_kernel void @no_free_sgprs_block_count_x(ptr addrspace(1) inreg %o
; GFX90a-NEXT: global_store_dword v0, v1, s[14:15]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: no_free_sgprs_block_count_x:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[14:15], s[8:9], 0x0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB2_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB2_0:
+; GISEL-GFX90A-NEXT: s_load_dword s0, s[8:9], 0x28
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s0
+; GISEL-GFX90A-NEXT: global_store_dword v1, v0, s[14:15]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: no_free_sgprs_block_count_x:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -151,6 +195,16 @@ define amdgpu_kernel void @no_inreg_block_count_x(ptr addrspace(1) %out) #0 {
; GFX90a-NEXT: global_store_dword v0, v1, s[0:1]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: no_inreg_block_count_x:
+; GISEL-GFX90A: ; %bb.0:
+; GISEL-GFX90A-NEXT: s_load_dword s2, s[4:5], 0x8
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s2
+; GISEL-GFX90A-NEXT: global_store_dword v1, v0, s[0:1]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: no_inreg_block_count_x:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -189,6 +243,16 @@ define amdgpu_kernel void @mixed_inreg_block_count_x(ptr addrspace(1) %out, i32
; GFX90a-NEXT: global_store_dword v0, v1, s[0:1]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: mixed_inreg_block_count_x:
+; GISEL-GFX90A: ; %bb.0:
+; GISEL-GFX90A-NEXT: s_load_dword s2, s[4:5], 0x10
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s2
+; GISEL-GFX90A-NEXT: global_store_dword v1, v0, s[0:1]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: mixed_inreg_block_count_x:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -236,6 +300,21 @@ define amdgpu_kernel void @incorrect_type_i64_block_count_x(ptr addrspace(1) inr
; GFX90a-NEXT: global_store_dwordx2 v0, v[2:3], s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: incorrect_type_i64_block_count_x:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[8:9], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB5_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB5_0:
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x8
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v2, 0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: v_pk_mov_b32 v[0:1], s[0:1], s[0:1] op_sel:[0,1]
+; GISEL-GFX90A-NEXT: global_store_dwordx2 v2, v[0:1], s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: incorrect_type_i64_block_count_x:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -282,6 +361,21 @@ define amdgpu_kernel void @incorrect_type_i16_block_count_x(ptr addrspace(1) inr
; GFX90a-NEXT: global_store_short v0, v1, s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: incorrect_type_i16_block_count_x:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[8:9], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB6_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB6_0:
+; GISEL-GFX90A-NEXT: s_load_dword s0, s[4:5], 0x8
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s0
+; GISEL-GFX90A-NEXT: global_store_short v1, v0, s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-FAKE16-LABEL: incorrect_type_i16_block_count_x:
; GFX1250-FAKE16: ; %bb.0:
; GFX1250-FAKE16-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -333,6 +427,19 @@ define amdgpu_kernel void @preload_block_count_y(ptr addrspace(1) inreg %out) #0
; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: preload_block_count_y:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB7_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB7_0:
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s11
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: global_store_dword v1, v0, s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: preload_block_count_y:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -377,6 +484,21 @@ define amdgpu_kernel void @random_incorrect_offset(ptr addrspace(1) inreg %out)
; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: random_incorrect_offset:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[8:9], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB8_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB8_0:
+; GISEL-GFX90A-NEXT: s_load_dword s0, s[4:5], 0xa
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s0
+; GISEL-GFX90A-NEXT: global_store_dword v1, v0, s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: random_incorrect_offset:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -422,6 +544,20 @@ define amdgpu_kernel void @preload_block_count_z(ptr addrspace(1) inreg %out) #0
; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: preload_block_count_z:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_load_dword s12, s[4:5], 0x10
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB9_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB9_0:
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s12
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: global_store_dword v1, v0, s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: preload_block_count_z:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -469,6 +605,22 @@ define amdgpu_kernel void @preload_block_count_x_imparg_align_ptr_i8(ptr addrspa
; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: preload_block_count_x_imparg_align_ptr_i8:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_load_dword s12, s[4:5], 0x10
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB10_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB10_0:
+; GISEL-GFX90A-NEXT: s_and_b32 s0, s10, 0xff
+; GISEL-GFX90A-NEXT: s_add_i32 s0, s12, s0
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s0
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: global_store_dword v1, v0, s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: preload_block_count_x_imparg_align_ptr_i8:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -520,6 +672,25 @@ define amdgpu_kernel void @preload_block_count_xyz(ptr addrspace(1) inreg %out)
; GFX90a-NEXT: global_store_dwordx3 v3, v[0:2], s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: preload_block_count_xyz:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_load_dword s12, s[4:5], 0x10
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB11_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB11_0:
+; GISEL-GFX90A-NEXT: s_mov_b32 s0, s10
+; GISEL-GFX90A-NEXT: s_mov_b32 s2, s12
+; GISEL-GFX90A-NEXT: s_mov_b32 s1, s11
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s0
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, s1
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v2, s2
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v3, 0
+; GISEL-GFX90A-NEXT: global_store_dwordx3 v3, v[0:2], s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: preload_block_count_xyz:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -572,6 +743,21 @@ define amdgpu_kernel void @preload_workgroup_size_x(ptr addrspace(1) inreg %out)
; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: preload_workgroup_size_x:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[12:13], s[4:5], 0x10
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB12_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB12_0:
+; GISEL-GFX90A-NEXT: s_and_b32 s0, s13, 0xffff
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s0
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: global_store_dword v1, v0, s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: preload_workgroup_size_x:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -619,6 +805,21 @@ define amdgpu_kernel void @preload_workgroup_size_y(ptr addrspace(1) inreg %out)
; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: preload_workgroup_size_y:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[12:13], s[4:5], 0x10
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB13_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB13_0:
+; GISEL-GFX90A-NEXT: s_lshr_b32 s0, s13, 16
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s0
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: global_store_dword v1, v0, s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: preload_workgroup_size_y:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -668,6 +869,22 @@ define amdgpu_kernel void @preload_workgroup_size_z(ptr addrspace(1) inreg %out)
; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: preload_workgroup_size_z:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[12:13], s[4:5], 0x10
+; GISEL-GFX90A-NEXT: s_load_dword s14, s[4:5], 0x18
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB14_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB14_0:
+; GISEL-GFX90A-NEXT: s_and_b32 s0, s14, 0xffff
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s0
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: global_store_dword v1, v0, s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: preload_workgroup_size_z:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -725,6 +942,26 @@ define amdgpu_kernel void @preload_workgroup_size_xyz(ptr addrspace(1) inreg %ou
; GFX90a-NEXT: global_store_dwordx3 v3, v[0:2], s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: preload_workgroup_size_xyz:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[12:13], s[4:5], 0x10
+; GISEL-GFX90A-NEXT: s_load_dword s14, s[4:5], 0x18
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB15_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB15_0:
+; GISEL-GFX90A-NEXT: s_and_b32 s0, s13, 0xffff
+; GISEL-GFX90A-NEXT: s_lshr_b32 s1, s13, 16
+; GISEL-GFX90A-NEXT: s_and_b32 s2, s14, 0xffff
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s0
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, s1
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v2, s2
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v3, 0
+; GISEL-GFX90A-NEXT: global_store_dwordx3 v3, v[0:2], s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: preload_workgroup_size_xyz:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -785,6 +1022,22 @@ define amdgpu_kernel void @preload_remainder_x(ptr addrspace(1) inreg %out) #0 {
; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: preload_remainder_x:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[12:13], s[4:5], 0x10
+; GISEL-GFX90A-NEXT: s_load_dword s14, s[4:5], 0x18
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB16_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB16_0:
+; GISEL-GFX90A-NEXT: s_lshr_b32 s0, s14, 16
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s0
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: global_store_dword v1, v0, s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: preload_remainder_x:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -832,6 +1085,20 @@ define amdgpu_kernel void @preloadremainder_y(ptr addrspace(1) inreg %out) #0 {
; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: preloadremainder_y:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB17_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB17_0:
+; GISEL-GFX90A-NEXT: s_and_b32 s0, s15, 0xffff
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s0
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: global_store_dword v1, v0, s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: preloadremainder_y:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -879,6 +1146,20 @@ define amdgpu_kernel void @preloadremainder_z(ptr addrspace(1) inreg %out) #0 {
; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: preloadremainder_z:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB18_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB18_0:
+; GISEL-GFX90A-NEXT: s_lshr_b32 s0, s15, 16
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s0
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: global_store_dword v1, v0, s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: preloadremainder_z:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -934,6 +1215,24 @@ define amdgpu_kernel void @preloadremainder_xyz(ptr addrspace(1) inreg %out) #0
; GFX90a-NEXT: global_store_dwordx3 v3, v[0:2], s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: preloadremainder_xyz:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB19_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB19_0:
+; GISEL-GFX90A-NEXT: s_lshr_b32 s0, s14, 16
+; GISEL-GFX90A-NEXT: s_lshr_b32 s2, s15, 16
+; GISEL-GFX90A-NEXT: s_and_b32 s1, s15, 0xffff
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s0
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, s1
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v2, s2
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v3, 0
+; GISEL-GFX90A-NEXT: global_store_dwordx3 v3, v[0:2], s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: preloadremainder_xyz:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -992,6 +1291,22 @@ define amdgpu_kernel void @no_free_sgprs_preloadremainder_z(ptr addrspace(1) inr
; GFX90a-NEXT: global_store_dword v0, v1, s[14:15]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: no_free_sgprs_preloadremainder_z:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[14:15], s[8:9], 0x0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB20_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB20_0:
+; GISEL-GFX90A-NEXT: s_load_dword s0, s[8:9], 0x1c
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_lshr_b32 s0, s0, 16
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s0
+; GISEL-GFX90A-NEXT: global_store_dword v1, v0, s[14:15]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: no_free_sgprs_preloadremainder_z:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -1041,6 +1356,21 @@ define amdgpu_kernel void @preload_block_max_user_sgprs(ptr addrspace(1) inreg %
; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: preload_block_max_user_sgprs:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[8:9], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB21_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB21_0:
+; GISEL-GFX90A-NEXT: s_load_dword s0, s[4:5], 0x28
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s0
+; GISEL-GFX90A-NEXT: global_store_dword v1, v0, s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: preload_block_max_user_sgprs:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -1090,6 +1420,24 @@ define amdgpu_kernel void @preload_block_count_z_workgroup_size_z_remainder_z(pt
; GFX90a-NEXT: global_store_dwordx3 v3, v[0:2], s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: preload_block_count_z_workgroup_size_z_remainder_z:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB22_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB22_0:
+; GISEL-GFX90A-NEXT: s_mov_b32 s0, s14
+; GISEL-GFX90A-NEXT: s_lshr_b32 s14, s15, 16
+; GISEL-GFX90A-NEXT: s_and_b32 s13, s0, 0xffff
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s12
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, s13
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v2, s14
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v3, 0
+; GISEL-GFX90A-NEXT: global_store_dwordx3 v3, v[0:2], s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: preload_block_count_z_workgroup_size_z_remainder_z:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -1115,4 +1463,92 @@ define amdgpu_kernel void @preload_block_count_z_workgroup_size_z_remainder_z(pt
ret void
}
+define amdgpu_kernel void @preload_block_count_x_no_explicit_args() local_unnamed_addr #0 {
+; GFX942-LABEL: preload_block_count_x_no_explicit_args:
+; GFX942: ; %bb.3:
+; GFX942-NEXT: s_load_dword s2, s[0:1], 0x0
+; GFX942-NEXT: s_waitcnt lgkmcnt(0)
+; GFX942-NEXT: s_branch .LBB23_0
+; GFX942-NEXT: .p2align 8
+; GFX942-NEXT: ; %bb.4:
+; GFX942-NEXT: .LBB23_0: ; %entry
+; GFX942-NEXT: s_cmp_lg_u32 s2, 0
+; GFX942-NEXT: s_cbranch_scc0 .LBB23_2
+; GFX942-NEXT: ; %bb.1: ; %ret
+; GFX942-NEXT: s_endpgm
+; GFX942-NEXT: .LBB23_2: ; %store
+; GFX942-NEXT: v_mov_b64_e32 v[0:1], 0
+; GFX942-NEXT: v_mov_b64_e32 v[2:3], 32
+; GFX942-NEXT: flat_store_dwordx2 v[2:3], v[0:1]
+; GFX942-NEXT: s_endpgm
+;
+; GFX90a-LABEL: preload_block_count_x_no_explicit_args:
+; GFX90a: ; %bb.3:
+; GFX90a-NEXT: s_load_dword s8, s[4:5], 0x0
+; GFX90a-NEXT: s_waitcnt lgkmcnt(0)
+; GFX90a-NEXT: s_branch .LBB23_0
+; GFX90a-NEXT: .p2align 8
+; GFX90a-NEXT: ; %bb.4:
+; GFX90a-NEXT: .LBB23_0: ; %entry
+; GFX90a-NEXT: s_add_u32 flat_scratch_lo, s6, s12
+; GFX90a-NEXT: s_addc_u32 flat_scratch_hi, s7, 0
+; GFX90a-NEXT: s_cmp_lg_u32 s8, 0
+; GFX90a-NEXT: s_cbranch_scc0 .LBB23_2
+; GFX90a-NEXT: ; %bb.1: ; %ret
+; GFX90a-NEXT: s_endpgm
+; GFX90a-NEXT: .LBB23_2: ; %store
+; GFX90a-NEXT: v_mov_b32_e32 v0, 0
+; GFX90a-NEXT: v_mov_b32_e32 v1, v0
+; GFX90a-NEXT: v_mov_b32_e32 v2, 32
+; GFX90a-NEXT: v_mov_b32_e32 v3, 0
+; GFX90a-NEXT: flat_store_dwordx2 v[2:3], v[0:1]
+; GFX90a-NEXT: s_endpgm
+;
+; GISEL-GFX90A-LABEL: preload_block_count_x_no_explicit_args:
+; GISEL-GFX90A: ; %bb.3:
+; GISEL-GFX90A-NEXT: s_load_dword s8, s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB23_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.4:
+; GISEL-GFX90A-NEXT: .LBB23_0: ; %entry
+; GISEL-GFX90A-NEXT: s_add_u32 flat_scratch_lo, s6, s12
+; GISEL-GFX90A-NEXT: s_addc_u32 flat_scratch_hi, s7, 0
+; GISEL-GFX90A-NEXT: s_cmp_lg_u32 s8, 0
+; GISEL-GFX90A-NEXT: s_cbranch_scc0 .LBB23_2
+; GISEL-GFX90A-NEXT: ; %bb.1: ; %ret
+; GISEL-GFX90A-NEXT: s_endpgm
+; GISEL-GFX90A-NEXT: .LBB23_2: ; %store
+; GISEL-GFX90A-NEXT: v_pk_mov_b32 v[0:1], 0, 0
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v2, 32
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v3, 0
+; GISEL-GFX90A-NEXT: flat_store_dwordx2 v[2:3], v[0:1]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
+; GFX1250-LABEL: preload_block_count_x_no_explicit_args:
+; GFX1250: ; %bb.0: ; %entry
+; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX1250-NEXT: s_cmp_lg_u32 s2, 0
+; GFX1250-NEXT: s_cbranch_scc0 .LBB23_2
+; GFX1250-NEXT: ; %bb.1: ; %ret
+; GFX1250-NEXT: s_endpgm
+; GFX1250-NEXT: .LBB23_2: ; %store
+; GFX1250-NEXT: v_mov_b64_e32 v[0:1], 0
+; GFX1250-NEXT: v_mov_b64_e32 v[2:3], 32
+; GFX1250-NEXT: flat_store_b64 v[2:3], v[0:1]
+; GFX1250-NEXT: s_endpgm
+entry:
+ %imp_arg_ptr = tail call dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
+ %load = load i32, ptr addrspace(4) %imp_arg_ptr, align 4
+ %is_zero = icmp eq i32 %load, 0
+ br i1 %is_zero, label %store, label %ret
+
+store:
+ store ptr null, ptr inttoptr (i64 32 to ptr), align 32
+ br label %ret
+
+ret:
+ 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" }
diff --git a/llvm/test/CodeGen/AMDGPU/preload-kernargs.ll b/llvm/test/CodeGen/AMDGPU/preload-kernargs.ll
index 70d08930e3b9c..7833f163ead4d 100644
--- a/llvm/test/CodeGen/AMDGPU/preload-kernargs.ll
+++ b/llvm/test/CodeGen/AMDGPU/preload-kernargs.ll
@@ -1,6 +1,7 @@
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 4
; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx942 < %s | FileCheck -check-prefixes=GFX942 %s
; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx90a < %s | FileCheck -check-prefixes=GFX90a %s
+; RUN: llc -global-isel -mtriple=amdgcn--amdhsa -mcpu=gfx90a < %s | FileCheck -check-prefixes=GISEL-GFX90A %s
; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx1250 < %s | FileCheck -check-prefixes=GFX1250 %s
define amdgpu_kernel void @ptr1_i8(ptr addrspace(1) inreg %out, i8 inreg %arg0) #0 {
@@ -34,6 +35,21 @@ define amdgpu_kernel void @ptr1_i8(ptr addrspace(1) inreg %out, i8 inreg %arg0)
; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: ptr1_i8:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[8:9], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_load_dword s10, s[4:5], 0x8
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB0_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB0_0:
+; GISEL-GFX90A-NEXT: s_and_b32 s0, s10, 0xff
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s0
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: global_store_dword v1, v0, s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: ptr1_i8:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -78,6 +94,21 @@ define amdgpu_kernel void @ptr1_i8_zext_arg(ptr addrspace(1) inreg %out, i8 zero
; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: ptr1_i8_zext_arg:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[8:9], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_load_dword s10, s[4:5], 0x8
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB1_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB1_0:
+; GISEL-GFX90A-NEXT: s_and_b32 s0, s10, 0xff
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s0
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: global_store_dword v1, v0, s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: ptr1_i8_zext_arg:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -122,6 +153,21 @@ define amdgpu_kernel void @ptr1_i16_preload_arg(ptr addrspace(1) inreg %out, i16
; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: ptr1_i16_preload_arg:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[8:9], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_load_dword s10, s[4:5], 0x8
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB2_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB2_0:
+; GISEL-GFX90A-NEXT: s_and_b32 s0, s10, 0xffff
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s0
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: global_store_dword v1, v0, s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: ptr1_i16_preload_arg:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -164,6 +210,20 @@ define amdgpu_kernel void @ptr1_i32_preload_arg(ptr addrspace(1) inreg %out, i32
; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: ptr1_i32_preload_arg:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[8:9], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_load_dword s10, s[4:5], 0x8
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB3_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB3_0:
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s10
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: global_store_dword v1, v0, s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: ptr1_i32_preload_arg:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -207,6 +267,21 @@ define amdgpu_kernel void @i32_ptr1_i32_preload_arg(i32 inreg %arg0, ptr addrspa
; GFX90a-NEXT: global_store_dword v0, v1, s[10:11]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: i32_ptr1_i32_preload_arg:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_load_dword s12, s[4:5], 0x10
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB4_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB4_0:
+; GISEL-GFX90A-NEXT: s_add_i32 s0, s8, s12
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s0
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: global_store_dword v1, v0, s[10:11]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: i32_ptr1_i32_preload_arg:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -255,6 +330,23 @@ define amdgpu_kernel void @ptr1_i16_i16_preload_arg(ptr addrspace(1) inreg %out,
; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: ptr1_i16_i16_preload_arg:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[8:9], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_load_dword s10, s[4:5], 0x8
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB5_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB5_0:
+; GISEL-GFX90A-NEXT: s_lshr_b32 s0, s10, 16
+; GISEL-GFX90A-NEXT: s_and_b32 s1, s10, 0xffff
+; GISEL-GFX90A-NEXT: s_add_i32 s0, s1, s0
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s0
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: global_store_dword v1, v0, s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: ptr1_i16_i16_preload_arg:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -301,6 +393,20 @@ define amdgpu_kernel void @ptr1_v2i8_preload_arg(ptr addrspace(1) inreg %out, <2
; GFX90a-NEXT: global_store_short v0, v1, s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: ptr1_v2i8_preload_arg:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[8:9], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_load_dword s10, s[4:5], 0x8
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB6_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB6_0:
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s10
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: global_store_short v1, v0, s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: ptr1_v2i8_preload_arg:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -351,6 +457,25 @@ define amdgpu_kernel void @byref_preload_arg(ptr addrspace(1) inreg %out, ptr ad
; GFX90a-NEXT: s_waitcnt vmcnt(0)
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: byref_preload_arg:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[8:9], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB7_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB7_0:
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x100
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, 0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, s0
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v2, s1
+; GISEL-GFX90A-NEXT: global_store_dword v0, v1, s[8:9]
+; GISEL-GFX90A-NEXT: s_waitcnt vmcnt(0)
+; GISEL-GFX90A-NEXT: global_store_dword v0, v2, s[8:9]
+; GISEL-GFX90A-NEXT: s_waitcnt vmcnt(0)
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: byref_preload_arg:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -411,6 +536,25 @@ define amdgpu_kernel void @byref_staggered_preload_arg(ptr addrspace(1) inreg %o
; GFX90a-NEXT: s_waitcnt vmcnt(0)
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: byref_staggered_preload_arg:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[8:9], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB8_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB8_0:
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x100
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, 0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, s0
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v2, s1
+; GISEL-GFX90A-NEXT: global_store_dword v0, v1, s[8:9]
+; GISEL-GFX90A-NEXT: s_waitcnt vmcnt(0)
+; GISEL-GFX90A-NEXT: global_store_dword v0, v2, s[8:9]
+; GISEL-GFX90A-NEXT: s_waitcnt vmcnt(0)
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: byref_staggered_preload_arg:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -480,6 +624,25 @@ define amdgpu_kernel void @v8i32_arg(ptr addrspace(1) nocapture inreg %out, <8 x
; GFX90a-NEXT: global_store_dwordx4 v4, v[0:3], s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: v8i32_arg:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[8:9], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB9_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB9_0:
+; GISEL-GFX90A-NEXT: s_load_dwordx8 s[12:19], s[4:5], 0x20
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v8, 0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: v_pk_mov_b32 v[0:1], s[12:13], s[12:13] op_sel:[0,1]
+; GISEL-GFX90A-NEXT: v_pk_mov_b32 v[2:3], s[14:15], s[14:15] op_sel:[0,1]
+; GISEL-GFX90A-NEXT: v_pk_mov_b32 v[4:5], s[16:17], s[16:17] op_sel:[0,1]
+; GISEL-GFX90A-NEXT: v_pk_mov_b32 v[6:7], s[18:19], s[18:19] op_sel:[0,1]
+; GISEL-GFX90A-NEXT: global_store_dwordx4 v8, v[0:3], s[8:9]
+; GISEL-GFX90A-NEXT: global_store_dwordx4 v8, v[4:7], s[8:9] offset:16
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: v8i32_arg:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -528,6 +691,24 @@ define amdgpu_kernel void @v3i16_preload_arg(ptr addrspace(1) nocapture inreg %o
; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: v3i16_preload_arg:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB10_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB10_0:
+; GISEL-GFX90A-NEXT: s_lshr_b32 s0, s10, 16
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s10
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: global_store_short v1, v0, s[8:9]
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s0
+; GISEL-GFX90A-NEXT: global_store_short v1, v0, s[8:9] offset:2
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s11
+; GISEL-GFX90A-NEXT: global_store_short v1, v0, s[8:9] offset:4
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: v3i16_preload_arg:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -574,6 +755,21 @@ define amdgpu_kernel void @v3i32_preload_arg(ptr addrspace(1) nocapture inreg %o
; GFX90a-NEXT: global_store_dwordx3 v3, v[0:2], s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: v3i32_preload_arg:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB11_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB11_0:
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s12
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, s13
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v2, s14
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v3, 0
+; GISEL-GFX90A-NEXT: global_store_dwordx3 v3, v[0:2], s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: v3i32_preload_arg:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -618,6 +814,21 @@ define amdgpu_kernel void @v3f32_preload_arg(ptr addrspace(1) nocapture inreg %o
; GFX90a-NEXT: global_store_dwordx3 v3, v[0:2], s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: v3f32_preload_arg:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB12_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB12_0:
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s12
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, s13
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v2, s14
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v3, 0
+; GISEL-GFX90A-NEXT: global_store_dwordx3 v3, v[0:2], s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: v3f32_preload_arg:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -675,6 +886,31 @@ define amdgpu_kernel void @v5i8_preload_arg(ptr addrspace(1) nocapture inreg %ou
; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: v5i8_preload_arg:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB13_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB13_0:
+; GISEL-GFX90A-NEXT: s_and_b32 s1, 0xffff, s10
+; GISEL-GFX90A-NEXT: s_lshr_b32 s1, s1, 8
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s10
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: s_lshr_b32 s0, s10, 16
+; GISEL-GFX90A-NEXT: global_store_byte v1, v0, s[8:9]
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s1
+; GISEL-GFX90A-NEXT: s_lshr_b32 s2, s10, 24
+; GISEL-GFX90A-NEXT: global_store_byte v1, v0, s[8:9] offset:1
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s0
+; GISEL-GFX90A-NEXT: global_store_byte v1, v0, s[8:9] offset:2
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s2
+; GISEL-GFX90A-NEXT: global_store_byte v1, v0, s[8:9] offset:3
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s11
+; GISEL-GFX90A-NEXT: global_store_byte v1, v0, s[8:9] offset:4
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: v5i8_preload_arg:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -747,6 +983,28 @@ define amdgpu_kernel void @v5f64_arg(ptr addrspace(1) nocapture inreg %out, <5 x
; GFX90a-NEXT: global_store_dwordx4 v4, v[0:3], s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: v5f64_arg:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[8:9], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB14_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB14_0:
+; GISEL-GFX90A-NEXT: s_load_dwordx8 s[12:19], s[4:5], 0x40
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x60
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v8, 0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: v_pk_mov_b32 v[0:1], s[12:13], s[12:13] op_sel:[0,1]
+; GISEL-GFX90A-NEXT: v_pk_mov_b32 v[2:3], s[14:15], s[14:15] op_sel:[0,1]
+; GISEL-GFX90A-NEXT: v_pk_mov_b32 v[4:5], s[16:17], s[16:17] op_sel:[0,1]
+; GISEL-GFX90A-NEXT: v_pk_mov_b32 v[6:7], s[18:19], s[18:19] op_sel:[0,1]
+; GISEL-GFX90A-NEXT: global_store_dwordx4 v8, v[0:3], s[8:9]
+; GISEL-GFX90A-NEXT: global_store_dwordx4 v8, v[4:7], s[8:9] offset:16
+; GISEL-GFX90A-NEXT: v_pk_mov_b32 v[0:1], s[0:1], s[0:1] op_sel:[0,1]
+; GISEL-GFX90A-NEXT: global_store_dwordx2 v8, v[0:1], s[8:9] offset:32
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: v5f64_arg:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -827,6 +1085,19 @@ define amdgpu_kernel void @v8i8_preload_arg(ptr addrspace(1) inreg %out, <8 x i8
; GFX90a-NEXT: global_store_dwordx2 v2, v[0:1], s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: v8i8_preload_arg:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB15_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB15_0:
+; GISEL-GFX90A-NEXT: v_pk_mov_b32 v[0:1], s[10:11], s[10:11] op_sel:[0,1]
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v2, 0
+; GISEL-GFX90A-NEXT: global_store_dwordx2 v2, v[0:1], s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: v8i8_preload_arg:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -873,6 +1144,19 @@ define amdgpu_kernel void @i64_kernel_preload_arg(ptr addrspace(1) inreg %out, i
; GFX90a-NEXT: global_store_dwordx2 v0, v[2:3], s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: i64_kernel_preload_arg:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB16_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB16_0:
+; GISEL-GFX90A-NEXT: v_pk_mov_b32 v[0:1], s[10:11], s[10:11] op_sel:[0,1]
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v2, 0
+; GISEL-GFX90A-NEXT: global_store_dwordx2 v2, v[0:1], s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: i64_kernel_preload_arg:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -912,6 +1196,19 @@ define amdgpu_kernel void @f64_kernel_preload_arg(ptr addrspace(1) inreg %out, d
; GFX90a-NEXT: global_store_dwordx2 v0, v[2:3], s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: f64_kernel_preload_arg:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB17_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB17_0:
+; GISEL-GFX90A-NEXT: v_pk_mov_b32 v[0:1], s[10:11], s[10:11] op_sel:[0,1]
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v2, 0
+; GISEL-GFX90A-NEXT: global_store_dwordx2 v2, v[0:1], s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: f64_kernel_preload_arg:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -952,6 +1249,20 @@ define amdgpu_kernel void @half_kernel_preload_arg(ptr addrspace(1) inreg %out,
; GFX90a-NEXT: global_store_short v0, v1, s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: half_kernel_preload_arg:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[8:9], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_load_dword s10, s[4:5], 0x8
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB18_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB18_0:
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s10
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: global_store_short v1, v0, s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: half_kernel_preload_arg:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -991,6 +1302,20 @@ define amdgpu_kernel void @bfloat_kernel_preload_arg(ptr addrspace(1) inreg %out
; GFX90a-NEXT: global_store_short v0, v1, s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: bfloat_kernel_preload_arg:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[8:9], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_load_dword s10, s[4:5], 0x8
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB19_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB19_0:
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s10
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: global_store_short v1, v0, s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: bfloat_kernel_preload_arg:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -1030,6 +1355,20 @@ define amdgpu_kernel void @v2bfloat_kernel_preload_arg(ptr addrspace(1) inreg %o
; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: v2bfloat_kernel_preload_arg:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[8:9], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_load_dword s10, s[4:5], 0x8
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB20_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB20_0:
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s10
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: global_store_dword v1, v0, s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: v2bfloat_kernel_preload_arg:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -1072,6 +1411,24 @@ define amdgpu_kernel void @v3bfloat_kernel_preload_arg(ptr addrspace(1) inreg %o
; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: v3bfloat_kernel_preload_arg:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB21_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB21_0:
+; GISEL-GFX90A-NEXT: s_lshr_b32 s0, s10, 16
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s10
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: global_store_short v1, v0, s[8:9]
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s0
+; GISEL-GFX90A-NEXT: global_store_short v1, v0, s[8:9] offset:2
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s11
+; GISEL-GFX90A-NEXT: global_store_short v1, v0, s[8:9] offset:4
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: v3bfloat_kernel_preload_arg:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -1118,6 +1475,21 @@ define amdgpu_kernel void @v6bfloat_kernel_preload_arg(ptr addrspace(1) inreg %o
; GFX90a-NEXT: global_store_dwordx3 v3, v[0:2], s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: v6bfloat_kernel_preload_arg:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB22_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB22_0:
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s12
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, s13
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v2, s14
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v3, 0
+; GISEL-GFX90A-NEXT: global_store_dwordx3 v3, v[0:2], s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: v6bfloat_kernel_preload_arg:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -1171,6 +1543,38 @@ define amdgpu_kernel void @half_v7bfloat_kernel_preload_arg(ptr addrspace(1) inr
; GFX90a-NEXT: global_store_dwordx3 v3, v[0:2], s[0:1]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: half_v7bfloat_kernel_preload_arg:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB23_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB23_0:
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x20
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s10
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: s_lshr_b32 s2, s12, 16
+; GISEL-GFX90A-NEXT: global_store_short v1, v0, s[8:9]
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s12
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: global_store_short v1, v0, s[0:1]
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s2
+; GISEL-GFX90A-NEXT: s_lshr_b32 s3, s13, 16
+; GISEL-GFX90A-NEXT: global_store_short v1, v0, s[0:1] offset:2
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s13
+; GISEL-GFX90A-NEXT: global_store_short v1, v0, s[0:1] offset:4
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s3
+; GISEL-GFX90A-NEXT: s_lshr_b32 s6, s14, 16
+; GISEL-GFX90A-NEXT: global_store_short v1, v0, s[0:1] offset:6
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s14
+; GISEL-GFX90A-NEXT: global_store_short v1, v0, s[0:1] offset:8
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s6
+; GISEL-GFX90A-NEXT: global_store_short v1, v0, s[0:1] offset:10
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s15
+; GISEL-GFX90A-NEXT: global_store_short v1, v0, s[0:1] offset:12
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: half_v7bfloat_kernel_preload_arg:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -1218,6 +1622,21 @@ define amdgpu_kernel void @i1_kernel_preload_arg(ptr addrspace(1) inreg %out, i1
; GFX90a-NEXT: global_store_byte v0, v1, s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: i1_kernel_preload_arg:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[8:9], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_load_dword s10, s[4:5], 0x8
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB24_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB24_0:
+; GISEL-GFX90A-NEXT: s_and_b32 s0, s10, 1
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s0
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: global_store_byte v1, v0, s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: i1_kernel_preload_arg:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -1265,6 +1684,20 @@ define amdgpu_kernel void @fp128_kernel_preload_arg(ptr addrspace(1) inreg %out,
; GFX90a-NEXT: global_store_dwordx4 v4, v[0:3], s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: fp128_kernel_preload_arg:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB25_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB25_0:
+; GISEL-GFX90A-NEXT: v_pk_mov_b32 v[0:1], s[12:13], s[12:13] op_sel:[0,1]
+; GISEL-GFX90A-NEXT: v_pk_mov_b32 v[2:3], s[14:15], s[14:15] op_sel:[0,1]
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v4, 0
+; GISEL-GFX90A-NEXT: global_store_dwordx4 v4, v[0:3], s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: fp128_kernel_preload_arg:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -1325,6 +1758,38 @@ define amdgpu_kernel void @v7i8_kernel_preload_arg(ptr addrspace(1) inreg %out,
; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: v7i8_kernel_preload_arg:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB26_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB26_0:
+; GISEL-GFX90A-NEXT: s_and_b32 s2, 0xffff, s10
+; GISEL-GFX90A-NEXT: s_lshr_b32 s2, s2, 8
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s10
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: s_lshr_b32 s0, s10, 16
+; GISEL-GFX90A-NEXT: global_store_byte v1, v0, s[8:9]
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s2
+; GISEL-GFX90A-NEXT: s_lshr_b32 s3, s10, 24
+; GISEL-GFX90A-NEXT: global_store_byte v1, v0, s[8:9] offset:1
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s0
+; GISEL-GFX90A-NEXT: s_and_b32 s4, 0xffff, s11
+; GISEL-GFX90A-NEXT: global_store_byte v1, v0, s[8:9] offset:2
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s3
+; GISEL-GFX90A-NEXT: s_lshr_b32 s4, s4, 8
+; GISEL-GFX90A-NEXT: global_store_byte v1, v0, s[8:9] offset:3
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s11
+; GISEL-GFX90A-NEXT: s_lshr_b32 s1, s11, 16
+; GISEL-GFX90A-NEXT: global_store_byte v1, v0, s[8:9] offset:4
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s4
+; GISEL-GFX90A-NEXT: global_store_byte v1, v0, s[8:9] offset:5
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s1
+; GISEL-GFX90A-NEXT: global_store_byte v1, v0, s[8:9] offset:6
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: v7i8_kernel_preload_arg:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -1380,6 +1845,34 @@ define amdgpu_kernel void @v7half_kernel_preload_arg(ptr addrspace(1) inreg %out
; GFX90a-NEXT: global_store_dwordx3 v3, v[0:2], s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: v7half_kernel_preload_arg:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB27_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB27_0:
+; GISEL-GFX90A-NEXT: s_lshr_b32 s0, s12, 16
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s12
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: global_store_short v1, v0, s[8:9]
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s0
+; GISEL-GFX90A-NEXT: s_lshr_b32 s1, s13, 16
+; GISEL-GFX90A-NEXT: global_store_short v1, v0, s[8:9] offset:2
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s13
+; GISEL-GFX90A-NEXT: global_store_short v1, v0, s[8:9] offset:4
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s1
+; GISEL-GFX90A-NEXT: s_lshr_b32 s2, s14, 16
+; GISEL-GFX90A-NEXT: global_store_short v1, v0, s[8:9] offset:6
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s14
+; GISEL-GFX90A-NEXT: global_store_short v1, v0, s[8:9] offset:8
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s2
+; GISEL-GFX90A-NEXT: global_store_short v1, v0, s[8:9] offset:10
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s15
+; GISEL-GFX90A-NEXT: global_store_short v1, v0, s[8:9] offset:12
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: v7half_kernel_preload_arg:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -1427,6 +1920,22 @@ define amdgpu_kernel void @i16_i32_kernel_preload_arg(ptr addrspace(1) inreg %ou
; GFX90a-NEXT: global_store_dword v0, v1, s[12:13]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: i16_i32_kernel_preload_arg:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[12:13], s[4:5], 0x10
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB28_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB28_0:
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s10
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: global_store_short v1, v0, s[8:9]
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s11
+; GISEL-GFX90A-NEXT: global_store_dword v1, v0, s[12:13]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: i16_i32_kernel_preload_arg:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -1479,6 +1988,25 @@ define amdgpu_kernel void @i16_v3i32_kernel_preload_arg(ptr addrspace(1) inreg %
; GFX90a-NEXT: global_store_dwordx3 v3, v[0:2], s[0:1]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: i16_v3i32_kernel_preload_arg:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB29_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB29_0:
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x20
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s10
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v3, 0
+; GISEL-GFX90A-NEXT: global_store_short v3, v0, s[8:9]
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s12
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, s13
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v2, s14
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: global_store_dwordx3 v3, v[0:2], s[0:1]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: i16_v3i32_kernel_preload_arg:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -1525,6 +2053,23 @@ define amdgpu_kernel void @i16_i16_kernel_preload_arg(ptr addrspace(1) inreg %ou
; GFX90a-NEXT: global_store_short_d16_hi v0, v1, s[12:13]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: i16_i16_kernel_preload_arg:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[12:13], s[4:5], 0x10
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB30_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB30_0:
+; GISEL-GFX90A-NEXT: s_lshr_b32 s0, s10, 16
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s10
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: global_store_short v1, v0, s[8:9]
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s0
+; GISEL-GFX90A-NEXT: global_store_short v1, v0, s[12:13]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: i16_i16_kernel_preload_arg:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -1579,6 +2124,23 @@ define amdgpu_kernel void @i16_v2i8_kernel_preload_arg(ptr addrspace(1) inreg %o
; GFX90a-NEXT: global_store_short v0, v1, s[12:13]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: i16_v2i8_kernel_preload_arg:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[12:13], s[4:5], 0x10
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB31_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB31_0:
+; GISEL-GFX90A-NEXT: s_lshr_b32 s0, s10, 16
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s10
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: global_store_short v1, v0, s[8:9]
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s0
+; GISEL-GFX90A-NEXT: global_store_short v1, v0, s[12:13]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: i16_v2i8_kernel_preload_arg:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -1629,6 +2191,23 @@ define amdgpu_kernel void @i32_ptr1_i32_staggered_preload_arg(i32 inreg %arg0, p
; GFX90a-NEXT: global_store_dword v0, v1, s[0:1]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: i32_ptr1_i32_staggered_preload_arg:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dword s8, s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB32_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB32_0:
+; GISEL-GFX90A-NEXT: s_load_dword s2, s[4:5], 0x10
+; GISEL-GFX90A-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x8
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_add_i32 s2, s8, s2
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s2
+; GISEL-GFX90A-NEXT: global_store_dword v1, v0, s[0:1]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: i32_ptr1_i32_staggered_preload_arg:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
@@ -1674,6 +2253,20 @@ define amdgpu_kernel void @ptr1_i8_trailing_unused(ptr addrspace(1) inreg %out,
; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]
; GFX90a-NEXT: s_endpgm
;
+; GISEL-GFX90A-LABEL: ptr1_i8_trailing_unused:
+; GISEL-GFX90A: ; %bb.1:
+; GISEL-GFX90A-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x0
+; GISEL-GFX90A-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-GFX90A-NEXT: s_branch .LBB33_0
+; GISEL-GFX90A-NEXT: .p2align 8
+; GISEL-GFX90A-NEXT: ; %bb.2:
+; GISEL-GFX90A-NEXT: .LBB33_0:
+; GISEL-GFX90A-NEXT: s_and_b32 s0, s10, 0xff
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v0, s0
+; GISEL-GFX90A-NEXT: v_mov_b32_e32 v1, 0
+; GISEL-GFX90A-NEXT: global_store_dword v1, v0, s[8:9]
+; GISEL-GFX90A-NEXT: s_endpgm
+;
; GFX1250-LABEL: ptr1_i8_trailing_unused:
; GFX1250: ; %bb.0:
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
More information about the llvm-commits
mailing list