[llvm] [AMDGPU] Fix hidden kernarg preload count inconsistency (PR #116759)
Austin Kerbow via llvm-commits
llvm-commits at lists.llvm.org
Fri Dec 6 13:29:07 PST 2024
https://github.com/kerbowa updated https://github.com/llvm/llvm-project/pull/116759
>From 69313b64cde16f7e8aee2113f38a341a25783fcb Mon Sep 17 00:00:00 2001
From: Austin Kerbow <Austin.Kerbow at amd.com>
Date: Sun, 10 Nov 2024 21:40:50 -0800
Subject: [PATCH 1/4] [AMDGPU] Fix hidden kernarg preload count inconsistency
It is possible that the number of hidden arguments that are selected to
be preloaded in AMDGPULowerKernel arguments and isel can differ. This
isn't an issue with explicit arguments since isel can lower the argument
correctly either way, but with hidden arguments we may have alignment
issues if we try to load these hidden arguments that were added to the
kernel signature.
The reason for the mismatch is that isel reserves an extra synthetic
user SGPR for module LDS.
Instead of teaching lowerFormalArguments how to handle these properly it
makes more sense and is less expensive to fix the mismatch and assert if
we ever run into this issue again. We should never be trying to lower
these in the normal way.
In a future change we probably want to revise how we track "synthetic"
user SGPRs and unify the handling in GCNUserSGPRUsageInfo. Sometimes
synthetic SGPRSs are considered user SGPRs and sometimes they are not.
Until then this patch resolves the inconsistency, fixes the bug, and is
otherwise a NFC.
---
.../AMDGPU/AMDGPULowerKernelArguments.cpp | 12 ++--
llvm/lib/Target/AMDGPU/SIISelLowering.cpp | 8 +++
.../AMDGPU/preload-implicit-kernargs.ll | 59 ++++++++++++++-----
.../AMDGPU/preload-kernargs-IR-lowering.ll | 6 +-
4 files changed, 62 insertions(+), 23 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPULowerKernelArguments.cpp b/llvm/lib/Target/AMDGPU/AMDGPULowerKernelArguments.cpp
index 9de4cf82d0faca..507f1f008b9831 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPULowerKernelArguments.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPULowerKernelArguments.cpp
@@ -144,17 +144,17 @@ class PreloadKernelArgInfo {
// Returns the maximum number of user SGPRs that we have available to preload
// arguments.
void setInitialFreeUserSGPRsCount() {
- const unsigned MaxUserSGPRs = ST.getMaxNumUserSGPRs();
GCNUserSGPRUsageInfo UserSGPRInfo(F, ST);
-
- NumFreeUserSGPRs = MaxUserSGPRs - UserSGPRInfo.getNumUsedUserSGPRs();
+ NumFreeUserSGPRs =
+ UserSGPRInfo.getNumFreeUserSGPRs() - 1 /* Synthetic SGPRs*/;
}
bool tryAllocPreloadSGPRs(unsigned AllocSize, uint64_t ArgOffset,
uint64_t LastExplicitArgOffset) {
// Check if this argument may be loaded into the same register as the
// previous argument.
- if (!isAligned(Align(4), ArgOffset) && AllocSize < 4)
+ if (ArgOffset == LastExplicitArgOffset && !isAligned(Align(4), ArgOffset) &&
+ AllocSize < 4)
return true;
// Pad SGPRs for kernarg alignment.
@@ -170,6 +170,7 @@ class PreloadKernelArgInfo {
// Try to allocate SGPRs to preload implicit kernel arguments.
void tryAllocImplicitArgPreloadSGPRs(uint64_t ImplicitArgsBaseOffset,
+ uint64_t LastExplicitArgOffset,
IRBuilder<> &Builder) {
Function *ImplicitArgPtr = Intrinsic::getDeclarationIfExists(
F.getParent(), Intrinsic::amdgcn_implicitarg_ptr);
@@ -215,7 +216,6 @@ class PreloadKernelArgInfo {
// argument can actually be preloaded.
std::sort(ImplicitArgLoads.begin(), ImplicitArgLoads.end(), less_second());
- uint64_t LastExplicitArgOffset = ImplicitArgsBaseOffset;
// If we fail to preload any implicit argument we know we don't have SGPRs
// to preload any subsequent ones with larger offsets. Find the first
// argument that we cannot preload.
@@ -485,7 +485,7 @@ static bool lowerKernelArguments(Function &F, const TargetMachine &TM) {
uint64_t ImplicitArgsBaseOffset =
alignTo(ExplicitArgOffset, ST.getAlignmentForImplicitArgPtr()) +
BaseOffset;
- PreloadInfo.tryAllocImplicitArgPreloadSGPRs(ImplicitArgsBaseOffset,
+ PreloadInfo.tryAllocImplicitArgPreloadSGPRs(ImplicitArgsBaseOffset, ExplicitArgOffset,
Builder);
}
diff --git a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
index 70230b5abc5171..058e596c514de2 100644
--- a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
+++ b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
@@ -3024,6 +3024,14 @@ SDValue SITargetLowering::LowerFormalArguments(
NewArg = DAG.getMergeValues({NewArg, Chain}, DL);
}
} else {
+#ifndef NDEBUG
+ if (Arg.isOrigArg()) {
+ Argument *OrigArg = Fn.getArg(Arg.getOrigArgIndex());
+ assert(!OrigArg->hasAttribute("amdgpu-hidden-argument") &&
+ "Hidden arguments should be preloaded");
+ }
+#endif // NDEBUG
+
NewArg =
lowerKernargMemParameter(DAG, VT, MemVT, DL, Chain, Offset,
Alignment, Ins[i].Flags.isSExt(), &Ins[i]);
diff --git a/llvm/test/CodeGen/AMDGPU/preload-implicit-kernargs.ll b/llvm/test/CodeGen/AMDGPU/preload-implicit-kernargs.ll
index 5b8acc31b22cfd..0eb1f1df02a9ae 100644
--- a/llvm/test/CodeGen/AMDGPU/preload-implicit-kernargs.ll
+++ b/llvm/test/CodeGen/AMDGPU/preload-implicit-kernargs.ll
@@ -439,13 +439,13 @@ define amdgpu_kernel void @preload_workgroup_size_xyz(ptr addrspace(1) inreg %ou
; GFX90a: s_trap 2 ; Kernarg preload header. Trap with incompatible firmware that doesn't support preloading kernel arguments.
; GFX90a-NEXT: .fill 63, 4, 0xbf800000 ; s_nop 0
; GFX90a-NEXT: ; %bb.0:
+; GFX90a-NEXT: v_mov_b32_e32 v3, 0
+; GFX90a-NEXT: global_load_ushort v2, v3, s[4:5] offset:24
; GFX90a-NEXT: s_lshr_b32 s0, s11, 16
; GFX90a-NEXT: s_and_b32 s1, s11, 0xffff
-; GFX90a-NEXT: s_and_b32 s2, s12, 0xffff
-; GFX90a-NEXT: v_mov_b32_e32 v3, 0
; GFX90a-NEXT: v_mov_b32_e32 v0, s1
; GFX90a-NEXT: v_mov_b32_e32 v1, s0
-; GFX90a-NEXT: v_mov_b32_e32 v2, s2
+; GFX90a-NEXT: s_waitcnt vmcnt(0)
; GFX90a-NEXT: global_store_dwordx3 v3, v[0:2], s[6:7]
; GFX90a-NEXT: s_endpgm
%imp_arg_ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
@@ -554,13 +554,13 @@ define amdgpu_kernel void @preloadremainder_xyz(ptr addrspace(1) inreg %out) #0
; GFX940: s_trap 2 ; Kernarg preload header. Trap with incompatible firmware that doesn't support preloading kernel arguments.
; GFX940-NEXT: .fill 63, 4, 0xbf800000 ; s_nop 0
; GFX940-NEXT: ; %bb.0:
-; GFX940-NEXT: s_lshr_b32 s0, s9, 16
-; GFX940-NEXT: s_lshr_b32 s1, s8, 16
-; GFX940-NEXT: s_and_b32 s4, s9, 0xffff
; GFX940-NEXT: v_mov_b32_e32 v3, 0
-; GFX940-NEXT: v_mov_b32_e32 v0, s1
-; GFX940-NEXT: v_mov_b32_e32 v1, s4
-; GFX940-NEXT: v_mov_b32_e32 v2, s0
+; GFX940-NEXT: global_load_ushort v2, v3, s[0:1] offset:30
+; GFX940-NEXT: s_lshr_b32 s0, s8, 16
+; GFX940-NEXT: s_and_b32 s1, s9, 0xffff
+; GFX940-NEXT: v_mov_b32_e32 v0, s0
+; GFX940-NEXT: v_mov_b32_e32 v1, s1
+; GFX940-NEXT: s_waitcnt vmcnt(0)
; GFX940-NEXT: global_store_dwordx3 v3, v[0:2], s[2:3] sc0 sc1
; GFX940-NEXT: s_endpgm
;
@@ -568,13 +568,14 @@ define amdgpu_kernel void @preloadremainder_xyz(ptr addrspace(1) inreg %out) #0
; GFX90a: s_trap 2 ; Kernarg preload header. Trap with incompatible firmware that doesn't support preloading kernel arguments.
; GFX90a-NEXT: .fill 63, 4, 0xbf800000 ; s_nop 0
; GFX90a-NEXT: ; %bb.0:
-; GFX90a-NEXT: s_lshr_b32 s0, s13, 16
-; GFX90a-NEXT: s_lshr_b32 s1, s12, 16
-; GFX90a-NEXT: s_and_b32 s2, s13, 0xffff
; GFX90a-NEXT: v_mov_b32_e32 v3, 0
-; GFX90a-NEXT: v_mov_b32_e32 v0, s1
-; GFX90a-NEXT: v_mov_b32_e32 v1, s2
-; GFX90a-NEXT: v_mov_b32_e32 v2, s0
+; GFX90a-NEXT: global_load_dword v0, v3, s[4:5] offset:26
+; GFX90a-NEXT: global_load_ushort v2, v3, s[4:5] offset:30
+; GFX90a-NEXT: s_lshr_b32 s0, s12, 16
+; GFX90a-NEXT: s_waitcnt vmcnt(1)
+; GFX90a-NEXT: v_lshrrev_b32_e32 v1, 16, v0
+; GFX90a-NEXT: v_mov_b32_e32 v0, s0
+; GFX90a-NEXT: s_waitcnt vmcnt(0)
; GFX90a-NEXT: global_store_dwordx3 v3, v[0:2], s[6:7]
; GFX90a-NEXT: s_endpgm
%imp_arg_ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
@@ -626,4 +627,32 @@ define amdgpu_kernel void @no_free_sgprs_preloadremainder_z(ptr addrspace(1) inr
ret void
}
+; Check for consistency between isel and earlier passes preload SGPR accounting.
+
+define amdgpu_kernel void @preload_block_max_user_sgprs(ptr addrspace(1) inreg %out, i192 inreg %t0, i32 inreg %t1) #0 {
+; GFX940-LABEL: preload_block_max_user_sgprs:
+; GFX940: s_trap 2 ; Kernarg preload header. Trap with incompatible firmware that doesn't support preloading kernel arguments.
+; GFX940-NEXT: .fill 63, 4, 0xbf800000 ; s_nop 0
+; GFX940-NEXT: ; %bb.0:
+; GFX940-NEXT: v_mov_b32_e32 v0, 0
+; GFX940-NEXT: v_mov_b32_e32 v1, s12
+; GFX940-NEXT: global_store_dword v0, v1, s[2:3] sc0 sc1
+; GFX940-NEXT: s_endpgm
+;
+; GFX90a-LABEL: preload_block_max_user_sgprs:
+; GFX90a: s_trap 2 ; Kernarg preload header. Trap with incompatible firmware that doesn't support preloading kernel arguments.
+; GFX90a-NEXT: .fill 63, 4, 0xbf800000 ; s_nop 0
+; GFX90a-NEXT: ; %bb.0:
+; GFX90a-NEXT: s_load_dword s0, s[4:5], 0x28
+; GFX90a-NEXT: v_mov_b32_e32 v0, 0
+; GFX90a-NEXT: s_waitcnt lgkmcnt(0)
+; GFX90a-NEXT: v_mov_b32_e32 v1, s0
+; GFX90a-NEXT: global_store_dword v0, v1, s[6:7]
+; GFX90a-NEXT: s_endpgm
+ %imp_arg_ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
+ %load = load i32, ptr addrspace(4) %imp_arg_ptr
+ store i32 %load, ptr addrspace(1) %out
+ ret void
+}
+
attributes #0 = { "amdgpu-no-agpr" "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-heap-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-multigrid-sync-arg" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
diff --git a/llvm/test/CodeGen/AMDGPU/preload-kernargs-IR-lowering.ll b/llvm/test/CodeGen/AMDGPU/preload-kernargs-IR-lowering.ll
index ab0fb7584d50ce..a1dd8060720832 100644
--- a/llvm/test/CodeGen/AMDGPU/preload-kernargs-IR-lowering.ll
+++ b/llvm/test/CodeGen/AMDGPU/preload-kernargs-IR-lowering.ll
@@ -187,15 +187,17 @@ define amdgpu_kernel void @test_preload_IR_lowering_kernel_8(ptr addrspace(1) %i
; PRELOAD-8-LABEL: define {{[^@]+}}@test_preload_IR_lowering_kernel_8
; PRELOAD-8-SAME: (ptr addrspace(1) inreg [[IN:%.*]], ptr addrspace(1) inreg [[IN1:%.*]], ptr addrspace(1) inreg [[IN2:%.*]], ptr addrspace(1) inreg [[IN3:%.*]], ptr addrspace(1) inreg [[OUT:%.*]], ptr addrspace(1) inreg [[OUT1:%.*]], ptr addrspace(1) inreg [[OUT2:%.*]], ptr addrspace(1) inreg [[OUT3:%.*]]) #[[ATTR0]] {
; PRELOAD-8-NEXT: [[TEST_PRELOAD_IR_LOWERING_KERNEL_8_KERNARG_SEGMENT:%.*]] = call nonnull align 16 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.kernarg.segment.ptr()
+; PRELOAD-8-NEXT: [[OUT2_KERNARG_OFFSET:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[TEST_PRELOAD_IR_LOWERING_KERNEL_8_KERNARG_SEGMENT]], i64 48
+; PRELOAD-8-NEXT: [[OUT2_LOAD:%.*]] = load ptr addrspace(1), ptr addrspace(4) [[OUT2_KERNARG_OFFSET]], align 16, !invariant.load [[META0:![0-9]+]]
; PRELOAD-8-NEXT: [[OUT3_KERNARG_OFFSET:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[TEST_PRELOAD_IR_LOWERING_KERNEL_8_KERNARG_SEGMENT]], i64 56
-; PRELOAD-8-NEXT: [[OUT3_LOAD:%.*]] = load ptr addrspace(1), ptr addrspace(4) [[OUT3_KERNARG_OFFSET]], align 8, !invariant.load [[META0:![0-9]+]]
+; PRELOAD-8-NEXT: [[OUT3_LOAD:%.*]] = load ptr addrspace(1), ptr addrspace(4) [[OUT3_KERNARG_OFFSET]], align 8, !invariant.load [[META0]]
; PRELOAD-8-NEXT: [[LOAD:%.*]] = load i32, ptr addrspace(1) [[IN]], align 4
; PRELOAD-8-NEXT: [[LOAD1:%.*]] = load i32, ptr addrspace(1) [[IN1]], align 4
; PRELOAD-8-NEXT: [[LOAD2:%.*]] = load i32, ptr addrspace(1) [[IN2]], align 4
; PRELOAD-8-NEXT: [[LOAD3:%.*]] = load i32, ptr addrspace(1) [[IN3]], align 4
; PRELOAD-8-NEXT: store i32 [[LOAD]], ptr addrspace(1) [[OUT]], align 4
; PRELOAD-8-NEXT: store i32 [[LOAD1]], ptr addrspace(1) [[OUT1]], align 4
-; PRELOAD-8-NEXT: store i32 [[LOAD2]], ptr addrspace(1) [[OUT2]], align 4
+; PRELOAD-8-NEXT: store i32 [[LOAD2]], ptr addrspace(1) [[OUT2_LOAD]], align 4
; PRELOAD-8-NEXT: store i32 [[LOAD3]], ptr addrspace(1) [[OUT3_LOAD]], align 4
; PRELOAD-8-NEXT: ret void
;
>From 85311bd3a47e01a5006402583573da226e8b6fc7 Mon Sep 17 00:00:00 2001
From: Austin Kerbow <Austin.Kerbow at amd.com>
Date: Sun, 1 Dec 2024 14:02:25 -0800
Subject: [PATCH 2/4] Diagnostic error and move LDSKernelId.
Use diagnostic error. Unify preload kernarg SGPR accounting and move
LDSKernelId to UserSGPRInfo so that synthetic SGPRs can be properly
tracked when calculating the number of available registers for preloads.
---
llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp | 11 +++
.../AMDGPU/AMDGPULowerKernelArguments.cpp | 11 +--
llvm/lib/Target/AMDGPU/GCNSubtarget.cpp | 11 ++-
llvm/lib/Target/AMDGPU/GCNSubtarget.h | 15 +++-
llvm/lib/Target/AMDGPU/SIISelLowering.cpp | 26 +++---
.../Target/AMDGPU/SIMachineFunctionInfo.cpp | 5 +-
.../lib/Target/AMDGPU/SIMachineFunctionInfo.h | 3 -
...alid-hidden-kernarg-in-kernel-signature.ll | 17 ++++
.../AMDGPU/preload-implicit-kernargs.ll | 79 ++++++++++++++-----
.../AMDGPU/preload-kernargs-IR-lowering.ll | 6 +-
10 files changed, 134 insertions(+), 50 deletions(-)
create mode 100644 llvm/test/CodeGen/AMDGPU/invalid-hidden-kernarg-in-kernel-signature.ll
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp b/llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp
index 3e24f7875ac898..64f0e773a1f27b 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp
@@ -520,6 +520,17 @@ bool AMDGPUCallLowering::lowerFormalArgumentsKernel(
// TODO: Align down to dword alignment and extract bits for extending loads.
for (auto &Arg : F.args()) {
+ // Hidden arguments that are in the kernel signature must be preloded to
+ // user SGPRs, or loaded via the implicit_arg ptr. Print a diagnostic
+ // error if a hidden argument is in the argument list and is not
+ // preloaded.
+ if (Arg.hasAttribute("amdgpu-hidden-argument")) {
+ DiagnosticInfoUnsupported NonPreloadHiddenArg(
+ *Arg.getParent(),
+ "Hidden argument in kernel signature was not preloaded");
+ F.getContext().diagnose(NonPreloadHiddenArg);
+ }
+
const bool IsByRef = Arg.hasByRefAttr();
Type *ArgTy = IsByRef ? Arg.getParamByRefType() : Arg.getType();
unsigned AllocSize = DL.getTypeAllocSize(ArgTy);
diff --git a/llvm/lib/Target/AMDGPU/AMDGPULowerKernelArguments.cpp b/llvm/lib/Target/AMDGPU/AMDGPULowerKernelArguments.cpp
index 507f1f008b9831..3b3f8a72cc4a3c 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPULowerKernelArguments.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPULowerKernelArguments.cpp
@@ -145,19 +145,19 @@ class PreloadKernelArgInfo {
// arguments.
void setInitialFreeUserSGPRsCount() {
GCNUserSGPRUsageInfo UserSGPRInfo(F, ST);
- NumFreeUserSGPRs =
- UserSGPRInfo.getNumFreeUserSGPRs() - 1 /* Synthetic SGPRs*/;
+ NumFreeUserSGPRs = UserSGPRInfo.getNumFreeKernargPreloadSGPRs();
}
bool tryAllocPreloadSGPRs(unsigned AllocSize, uint64_t ArgOffset,
uint64_t LastExplicitArgOffset) {
// Check if this argument may be loaded into the same register as the
// previous argument.
- if (ArgOffset == LastExplicitArgOffset && !isAligned(Align(4), ArgOffset) &&
- AllocSize < 4)
+ if (ArgOffset - LastExplicitArgOffset < 4 &&
+ !isAligned(Align(4), ArgOffset))
return true;
// Pad SGPRs for kernarg alignment.
+ ArgOffset = alignDown(ArgOffset, 4);
unsigned Padding = ArgOffset - LastExplicitArgOffset;
unsigned PaddingSGPRs = alignTo(Padding, 4) / 4;
unsigned NumPreloadSGPRs = alignTo(AllocSize, 4) / 4;
@@ -229,7 +229,8 @@ class PreloadKernelArgInfo {
LastExplicitArgOffset))
return true;
- LastExplicitArgOffset = LoadOffset + LoadSize;
+ LastExplicitArgOffset =
+ ImplicitArgsBaseOffset + LoadOffset + LoadSize;
return false;
});
diff --git a/llvm/lib/Target/AMDGPU/GCNSubtarget.cpp b/llvm/lib/Target/AMDGPU/GCNSubtarget.cpp
index 51361b75940560..4e65e6aa67d9a8 100644
--- a/llvm/lib/Target/AMDGPU/GCNSubtarget.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNSubtarget.cpp
@@ -748,6 +748,10 @@ GCNUserSGPRUsageInfo::GCNUserSGPRUsageInfo(const Function &F,
FlatScratchInit = true;
}
+ if (!AMDGPU::isGraphics(CC) && !IsKernel &&
+ !F.hasFnAttribute("amdgpu-no-lds-kernel-id"))
+ LDSKernelId = true;
+
if (hasImplicitBufferPtr())
NumUsedUserSGPRs += getNumUserSGPRForField(ImplicitBufferPtrID);
@@ -771,6 +775,9 @@ GCNUserSGPRUsageInfo::GCNUserSGPRUsageInfo(const Function &F,
if (hasPrivateSegmentSize())
NumUsedUserSGPRs += getNumUserSGPRForField(PrivateSegmentSizeID);
+
+ if (hasLDSKernelId())
+ NumSyntheticSGPRs += getNumUserSGPRForField(LDSKernelIdID);
}
void GCNUserSGPRUsageInfo::allocKernargPreloadSGPRs(unsigned NumSGPRs) {
@@ -779,6 +786,6 @@ void GCNUserSGPRUsageInfo::allocKernargPreloadSGPRs(unsigned NumSGPRs) {
NumUsedUserSGPRs += NumSGPRs;
}
-unsigned GCNUserSGPRUsageInfo::getNumFreeUserSGPRs() {
- return AMDGPU::getMaxNumUserSGPRs(ST) - NumUsedUserSGPRs;
+unsigned GCNUserSGPRUsageInfo::getNumFreeKernargPreloadSGPRs() {
+ return AMDGPU::getMaxNumUserSGPRs(ST) - (NumUsedUserSGPRs + NumSyntheticSGPRs);
}
diff --git a/llvm/lib/Target/AMDGPU/GCNSubtarget.h b/llvm/lib/Target/AMDGPU/GCNSubtarget.h
index 5cecaf6349c883..a179880dfae8c8 100644
--- a/llvm/lib/Target/AMDGPU/GCNSubtarget.h
+++ b/llvm/lib/Target/AMDGPU/GCNSubtarget.h
@@ -1678,11 +1678,13 @@ class GCNUserSGPRUsageInfo {
bool hasPrivateSegmentSize() const { return PrivateSegmentSize; }
+ bool hasLDSKernelId() const { return LDSKernelId; }
+
unsigned getNumKernargPreloadSGPRs() const { return NumKernargPreloadSGPRs; }
unsigned getNumUsedUserSGPRs() const { return NumUsedUserSGPRs; }
- unsigned getNumFreeUserSGPRs();
+ unsigned getNumFreeKernargPreloadSGPRs();
void allocKernargPreloadSGPRs(unsigned NumSGPRs);
@@ -1694,11 +1696,12 @@ class GCNUserSGPRUsageInfo {
KernargSegmentPtrID = 4,
DispatchIdID = 5,
FlatScratchInitID = 6,
- PrivateSegmentSizeID = 7
+ PrivateSegmentSizeID = 7,
+ LDSKernelIdID = 8
};
// Returns the size in number of SGPRs for preload user SGPR field.
- static unsigned getNumUserSGPRForField(UserSGPRID ID) {
+ static constexpr unsigned getNumUserSGPRForField(UserSGPRID ID) {
switch (ID) {
case ImplicitBufferPtrID:
return 2;
@@ -1716,6 +1719,8 @@ class GCNUserSGPRUsageInfo {
return 2;
case PrivateSegmentSizeID:
return 1;
+ case LDSKernelIdID:
+ return 1;
}
llvm_unreachable("Unknown UserSGPRID.");
}
@@ -1744,9 +1749,13 @@ class GCNUserSGPRUsageInfo {
bool PrivateSegmentSize = false;
+ bool LDSKernelId = false;
+
unsigned NumKernargPreloadSGPRs = 0;
unsigned NumUsedUserSGPRs = 0;
+
+ unsigned NumSyntheticSGPRs = 0;
};
} // end namespace llvm
diff --git a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
index 058e596c514de2..e6d1e96bb2a1df 100644
--- a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
+++ b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
@@ -2422,7 +2422,7 @@ void SITargetLowering::allocateSpecialInputSGPRs(
if (Info.hasWorkGroupIDZ())
allocateSGPR32Input(CCInfo, ArgInfo.WorkGroupIDZ);
- if (Info.hasLDSKernelId())
+ if (UserSGPRInfo.hasLDSKernelId())
allocateSGPR32Input(CCInfo, ArgInfo.LDSKernelId);
}
@@ -2545,8 +2545,8 @@ void SITargetLowering::allocatePreloadKernArgSGPRs(
unsigned Padding = ArgOffset - LastExplicitArgOffset;
unsigned PaddingSGPRs = alignTo(Padding, 4) / 4;
// Check for free user SGPRs for preloading.
- if (PaddingSGPRs + NumAllocSGPRs + 1 /*Synthetic SGPRs*/ >
- SGPRInfo.getNumFreeUserSGPRs()) {
+ if (PaddingSGPRs + NumAllocSGPRs >
+ SGPRInfo.getNumFreeKernargPreloadSGPRs()) {
InPreloadSequence = false;
break;
}
@@ -2574,7 +2574,8 @@ void SITargetLowering::allocateLDSKernelId(CCState &CCInfo, MachineFunction &MF,
const SIRegisterInfo &TRI,
SIMachineFunctionInfo &Info) const {
// Always allocate this last since it is a synthetic preload.
- if (Info.hasLDSKernelId()) {
+ const GCNUserSGPRUsageInfo &UserSGPRInfo = Info.getUserSGPRInfo();
+ if (UserSGPRInfo.hasLDSKernelId()) {
Register Reg = Info.addLDSKernelId();
MF.addLiveIn(Reg, &AMDGPU::SGPR_32RegClass);
CCInfo.AllocateReg(Reg);
@@ -2824,7 +2825,7 @@ SDValue SITargetLowering::LowerFormalArguments(
const GCNUserSGPRUsageInfo &UserSGPRInfo = Info->getUserSGPRInfo();
assert(!UserSGPRInfo.hasDispatchPtr() &&
!UserSGPRInfo.hasKernargSegmentPtr() && !Info->hasWorkGroupInfo() &&
- !Info->hasLDSKernelId() && !Info->hasWorkItemIDX() &&
+ !UserSGPRInfo.hasLDSKernelId() && !Info->hasWorkItemIDX() &&
!Info->hasWorkItemIDY() && !Info->hasWorkItemIDZ());
(void)UserSGPRInfo;
if (!Subtarget->enableFlatScratch())
@@ -3024,13 +3025,20 @@ SDValue SITargetLowering::LowerFormalArguments(
NewArg = DAG.getMergeValues({NewArg, Chain}, DL);
}
} else {
-#ifndef NDEBUG
+ // Hidden arguments that are in the kernel signature must be preloded to
+ // user SGPRs, or loaded via the implicit_arg ptr. Print a diagnostic
+ // error if a hidden argument is in the argument list and is not
+ // preloaded.
if (Arg.isOrigArg()) {
Argument *OrigArg = Fn.getArg(Arg.getOrigArgIndex());
- assert(!OrigArg->hasAttribute("amdgpu-hidden-argument") &&
- "Hidden arguments should be preloaded");
+ if (OrigArg->hasAttribute("amdgpu-hidden-argument")) {
+ DiagnosticInfoUnsupported NonPreloadHiddenArg(
+ *OrigArg->getParent(),
+ "Hidden argument in kernel signature was not preloaded",
+ DL.getDebugLoc());
+ DAG.getContext()->diagnose(NonPreloadHiddenArg);
+ }
}
-#endif // NDEBUG
NewArg =
lowerKernargMemParameter(DAG, VT, MemVT, DL, Chain, Offset,
diff --git a/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp b/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp
index 1e43d2727a00da..3013230e9de78c 100644
--- a/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp
@@ -38,7 +38,7 @@ SIMachineFunctionInfo::SIMachineFunctionInfo(const Function &F,
const GCNSubtarget *STI)
: AMDGPUMachineFunction(F, *STI), Mode(F, *STI), GWSResourcePSV(getTM(STI)),
UserSGPRInfo(F, *STI), WorkGroupIDX(false), WorkGroupIDY(false),
- WorkGroupIDZ(false), WorkGroupInfo(false), LDSKernelId(false),
+ WorkGroupIDZ(false), WorkGroupInfo(false),
PrivateSegmentWaveByteOffset(false), WorkItemIDX(false),
WorkItemIDY(false), WorkItemIDZ(false), ImplicitArgPtr(false),
GITPtrHigh(0xffffffff), HighBitsOf32BitAddress(0) {
@@ -131,9 +131,6 @@ SIMachineFunctionInfo::SIMachineFunctionInfo(const Function &F,
if (!F.hasFnAttribute("amdgpu-no-workitem-id-z") &&
ST.getMaxWorkitemID(F, 2) != 0)
WorkItemIDZ = true;
-
- if (!IsKernel && !F.hasFnAttribute("amdgpu-no-lds-kernel-id"))
- LDSKernelId = true;
}
if (isEntryFunction()) {
diff --git a/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h b/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h
index 2a754680fdc8ca..2877a3c4fe38ca 100644
--- a/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h
+++ b/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h
@@ -461,7 +461,6 @@ class SIMachineFunctionInfo final : public AMDGPUMachineFunction,
bool WorkGroupIDY : 1;
bool WorkGroupIDZ : 1;
bool WorkGroupInfo : 1;
- bool LDSKernelId : 1;
bool PrivateSegmentWaveByteOffset : 1;
bool WorkItemIDX : 1; // Always initialized.
@@ -822,8 +821,6 @@ class SIMachineFunctionInfo final : public AMDGPUMachineFunction,
return ArgInfo.WorkGroupInfo.getRegister();
}
- bool hasLDSKernelId() const { return LDSKernelId; }
-
// Add special VGPR inputs
void setWorkItemIDX(ArgDescriptor Arg) {
ArgInfo.WorkItemIDX = Arg;
diff --git a/llvm/test/CodeGen/AMDGPU/invalid-hidden-kernarg-in-kernel-signature.ll b/llvm/test/CodeGen/AMDGPU/invalid-hidden-kernarg-in-kernel-signature.ll
new file mode 100644
index 00000000000000..b90da58de95818
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/invalid-hidden-kernarg-in-kernel-signature.ll
@@ -0,0 +1,17 @@
+; RUN: not llc -global-isel=1 -mtriple=amdgcn--amdhsa -mcpu=gfx940 -start-after=amdgpu-lower-kernel-arguments < %s 2>&1 | FileCheck -check-prefix=ERROR %s
+; RUN: not llc -global-isel=0 -mtriple=amdgcn--amdhsa -mcpu=gfx940 -start-after=amdgpu-lower-kernel-arguments < %s 2>&1 | FileCheck -check-prefix=ERROR %s
+
+define amdgpu_kernel void @no_free_sgprs_block_count_x_no_preload_diag(ptr addrspace(1) inreg %out, i512 inreg, i32 inreg "amdgpu-hidden-argument" %_hidden_block_count_x) #0 {
+; ERROR: error: <unknown>:0:0: in function no_free_sgprs_block_count_x_no_preload_diag void (ptr addrspace(1), i512, i32): Hidden argument in kernel signature was not preloaded
+ store i32 %_hidden_block_count_x, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_kernel void @preloadremainder_z_no_preload_diag(ptr addrspace(1) inreg %out, i256 inreg, i32 inreg "amdgpu-hidden-argument" %_hidden_block_count_x, i32 inreg "amdgpu-hidden-argument" %_hidden_block_count_y, i32 inreg "amdgpu-hidden-argument" %_hidden_block_count_z, i16 inreg "amdgpu-hidden-argument" %_hidden_group_size_x, i16 inreg "amdgpu-hidden-argument" %_hidden_group_size_y, i16 inreg "amdgpu-hidden-argument" %_hidden_group_size_z, i16 inreg "amdgpu-hidden-argument" %_hidden_remainder_x, i16 inreg "amdgpu-hidden-argument" %_hidden_remainder_y, i16 inreg "amdgpu-hidden-argument" %_hidden_remainder_z) #0 {
+; ERROR: error: <unknown>:0:0: in function preloadremainder_z_no_preload_diag void (ptr addrspace(1), i256, i32, i32, i32, i16, i16, i16, i16, i16, i16): Hidden argument in kernel signature was not preloaded
+ %conv = zext i16 %_hidden_remainder_z to i32
+ store i32 %conv, ptr addrspace(1) %out
+ ret void
+}
+
+attributes #0 = { "amdgpu-no-agpr" "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-heap-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-multigrid-sync-arg" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
diff --git a/llvm/test/CodeGen/AMDGPU/preload-implicit-kernargs.ll b/llvm/test/CodeGen/AMDGPU/preload-implicit-kernargs.ll
index 0eb1f1df02a9ae..0c6d8dce193da6 100644
--- a/llvm/test/CodeGen/AMDGPU/preload-implicit-kernargs.ll
+++ b/llvm/test/CodeGen/AMDGPU/preload-implicit-kernargs.ll
@@ -439,13 +439,13 @@ define amdgpu_kernel void @preload_workgroup_size_xyz(ptr addrspace(1) inreg %ou
; GFX90a: s_trap 2 ; Kernarg preload header. Trap with incompatible firmware that doesn't support preloading kernel arguments.
; GFX90a-NEXT: .fill 63, 4, 0xbf800000 ; s_nop 0
; GFX90a-NEXT: ; %bb.0:
-; GFX90a-NEXT: v_mov_b32_e32 v3, 0
-; GFX90a-NEXT: global_load_ushort v2, v3, s[4:5] offset:24
; GFX90a-NEXT: s_lshr_b32 s0, s11, 16
; GFX90a-NEXT: s_and_b32 s1, s11, 0xffff
+; GFX90a-NEXT: s_and_b32 s2, s12, 0xffff
+; GFX90a-NEXT: v_mov_b32_e32 v3, 0
; GFX90a-NEXT: v_mov_b32_e32 v0, s1
; GFX90a-NEXT: v_mov_b32_e32 v1, s0
-; GFX90a-NEXT: s_waitcnt vmcnt(0)
+; GFX90a-NEXT: v_mov_b32_e32 v2, s2
; GFX90a-NEXT: global_store_dwordx3 v3, v[0:2], s[6:7]
; GFX90a-NEXT: s_endpgm
%imp_arg_ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
@@ -554,13 +554,13 @@ define amdgpu_kernel void @preloadremainder_xyz(ptr addrspace(1) inreg %out) #0
; GFX940: s_trap 2 ; Kernarg preload header. Trap with incompatible firmware that doesn't support preloading kernel arguments.
; GFX940-NEXT: .fill 63, 4, 0xbf800000 ; s_nop 0
; GFX940-NEXT: ; %bb.0:
+; GFX940-NEXT: s_lshr_b32 s0, s9, 16
+; GFX940-NEXT: s_lshr_b32 s1, s8, 16
+; GFX940-NEXT: s_and_b32 s4, s9, 0xffff
; GFX940-NEXT: v_mov_b32_e32 v3, 0
-; GFX940-NEXT: global_load_ushort v2, v3, s[0:1] offset:30
-; GFX940-NEXT: s_lshr_b32 s0, s8, 16
-; GFX940-NEXT: s_and_b32 s1, s9, 0xffff
-; GFX940-NEXT: v_mov_b32_e32 v0, s0
-; GFX940-NEXT: v_mov_b32_e32 v1, s1
-; GFX940-NEXT: s_waitcnt vmcnt(0)
+; GFX940-NEXT: v_mov_b32_e32 v0, s1
+; GFX940-NEXT: v_mov_b32_e32 v1, s4
+; GFX940-NEXT: v_mov_b32_e32 v2, s0
; GFX940-NEXT: global_store_dwordx3 v3, v[0:2], s[2:3] sc0 sc1
; GFX940-NEXT: s_endpgm
;
@@ -568,14 +568,13 @@ define amdgpu_kernel void @preloadremainder_xyz(ptr addrspace(1) inreg %out) #0
; GFX90a: s_trap 2 ; Kernarg preload header. Trap with incompatible firmware that doesn't support preloading kernel arguments.
; GFX90a-NEXT: .fill 63, 4, 0xbf800000 ; s_nop 0
; GFX90a-NEXT: ; %bb.0:
+; GFX90a-NEXT: s_lshr_b32 s0, s13, 16
+; GFX90a-NEXT: s_lshr_b32 s1, s12, 16
+; GFX90a-NEXT: s_and_b32 s2, s13, 0xffff
; GFX90a-NEXT: v_mov_b32_e32 v3, 0
-; GFX90a-NEXT: global_load_dword v0, v3, s[4:5] offset:26
-; GFX90a-NEXT: global_load_ushort v2, v3, s[4:5] offset:30
-; GFX90a-NEXT: s_lshr_b32 s0, s12, 16
-; GFX90a-NEXT: s_waitcnt vmcnt(1)
-; GFX90a-NEXT: v_lshrrev_b32_e32 v1, 16, v0
-; GFX90a-NEXT: v_mov_b32_e32 v0, s0
-; GFX90a-NEXT: s_waitcnt vmcnt(0)
+; GFX90a-NEXT: v_mov_b32_e32 v0, s1
+; GFX90a-NEXT: v_mov_b32_e32 v1, s2
+; GFX90a-NEXT: v_mov_b32_e32 v2, s0
; GFX90a-NEXT: global_store_dwordx3 v3, v[0:2], s[6:7]
; GFX90a-NEXT: s_endpgm
%imp_arg_ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
@@ -600,10 +599,8 @@ define amdgpu_kernel void @no_free_sgprs_preloadremainder_z(ptr addrspace(1) inr
; GFX940: s_trap 2 ; Kernarg preload header. Trap with incompatible firmware that doesn't support preloading kernel arguments.
; GFX940-NEXT: .fill 63, 4, 0xbf800000 ; s_nop 0
; GFX940-NEXT: ; %bb.0:
-; GFX940-NEXT: s_load_dword s0, s[4:5], 0x1c
+; GFX940-NEXT: s_lshr_b32 s0, s15, 16
; GFX940-NEXT: v_mov_b32_e32 v0, 0
-; GFX940-NEXT: s_waitcnt lgkmcnt(0)
-; GFX940-NEXT: s_lshr_b32 s0, s0, 16
; GFX940-NEXT: v_mov_b32_e32 v1, s0
; GFX940-NEXT: global_store_dword v0, v1, s[8:9] sc0 sc1
; GFX940-NEXT: s_endpgm
@@ -627,7 +624,7 @@ define amdgpu_kernel void @no_free_sgprs_preloadremainder_z(ptr addrspace(1) inr
ret void
}
-; Check for consistency between isel and earlier passes preload SGPR accounting.
+; Check for consistency between isel and earlier passes preload SGPR accounting with max preload SGPRs.
define amdgpu_kernel void @preload_block_max_user_sgprs(ptr addrspace(1) inreg %out, i192 inreg %t0, i32 inreg %t1) #0 {
; GFX940-LABEL: preload_block_max_user_sgprs:
@@ -655,4 +652,46 @@ define amdgpu_kernel void @preload_block_max_user_sgprs(ptr addrspace(1) inreg %
ret void
}
+define amdgpu_kernel void @preload_block_count_z_workgroup_size_z_remainder_z(ptr addrspace(1) inreg %out) #0 {
+; GFX940-LABEL: preload_block_count_z_workgroup_size_z_remainder_z:
+; GFX940: s_trap 2 ; Kernarg preload header. Trap with incompatible firmware that doesn't support preloading kernel arguments.
+; GFX940-NEXT: .fill 63, 4, 0xbf800000 ; s_nop 0
+; GFX940-NEXT: ; %bb.0:
+; GFX940-NEXT: s_lshr_b32 s0, s9, 16
+; GFX940-NEXT: s_and_b32 s1, s8, 0xffff
+; GFX940-NEXT: v_mov_b32_e32 v3, 0
+; GFX940-NEXT: v_mov_b32_e32 v0, s6
+; GFX940-NEXT: v_mov_b32_e32 v1, s1
+; GFX940-NEXT: v_mov_b32_e32 v2, s0
+; GFX940-NEXT: global_store_dwordx3 v3, v[0:2], s[2:3] sc0 sc1
+; GFX940-NEXT: s_endpgm
+;
+; GFX90a-LABEL: preload_block_count_z_workgroup_size_z_remainder_z:
+; GFX90a: s_trap 2 ; Kernarg preload header. Trap with incompatible firmware that doesn't support preloading kernel arguments.
+; GFX90a-NEXT: .fill 63, 4, 0xbf800000 ; s_nop 0
+; GFX90a-NEXT: ; %bb.0:
+; GFX90a-NEXT: s_lshr_b32 s0, s13, 16
+; GFX90a-NEXT: s_and_b32 s1, s12, 0xffff
+; GFX90a-NEXT: v_mov_b32_e32 v3, 0
+; GFX90a-NEXT: v_mov_b32_e32 v0, s10
+; GFX90a-NEXT: v_mov_b32_e32 v1, s1
+; GFX90a-NEXT: v_mov_b32_e32 v2, s0
+; GFX90a-NEXT: global_store_dwordx3 v3, v[0:2], s[6:7]
+; GFX90a-NEXT: s_endpgm
+ %imp_arg_ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
+ %gep0 = getelementptr i8, ptr addrspace(4) %imp_arg_ptr, i32 8
+ %gep1 = getelementptr i8, ptr addrspace(4) %imp_arg_ptr, i32 16
+ %gep2 = getelementptr i8, ptr addrspace(4) %imp_arg_ptr, i32 22
+ %load0 = load i32, ptr addrspace(4) %gep0
+ %load1 = load i16, ptr addrspace(4) %gep1
+ %load2 = load i16, ptr addrspace(4) %gep2
+ %conv1 = zext i16 %load1 to i32
+ %conv2 = zext i16 %load2 to i32
+ %ins.0 = insertelement <3 x i32> poison, i32 %load0, i32 0
+ %ins.1 = insertelement <3 x i32> %ins.0, i32 %conv1, i32 1
+ %ins.2 = insertelement <3 x i32> %ins.1, i32 %conv2, i32 2
+ store <3 x i32> %ins.2, ptr addrspace(1) %out
+ ret void
+}
+
attributes #0 = { "amdgpu-no-agpr" "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-heap-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-multigrid-sync-arg" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
diff --git a/llvm/test/CodeGen/AMDGPU/preload-kernargs-IR-lowering.ll b/llvm/test/CodeGen/AMDGPU/preload-kernargs-IR-lowering.ll
index a1dd8060720832..ab0fb7584d50ce 100644
--- a/llvm/test/CodeGen/AMDGPU/preload-kernargs-IR-lowering.ll
+++ b/llvm/test/CodeGen/AMDGPU/preload-kernargs-IR-lowering.ll
@@ -187,17 +187,15 @@ define amdgpu_kernel void @test_preload_IR_lowering_kernel_8(ptr addrspace(1) %i
; PRELOAD-8-LABEL: define {{[^@]+}}@test_preload_IR_lowering_kernel_8
; PRELOAD-8-SAME: (ptr addrspace(1) inreg [[IN:%.*]], ptr addrspace(1) inreg [[IN1:%.*]], ptr addrspace(1) inreg [[IN2:%.*]], ptr addrspace(1) inreg [[IN3:%.*]], ptr addrspace(1) inreg [[OUT:%.*]], ptr addrspace(1) inreg [[OUT1:%.*]], ptr addrspace(1) inreg [[OUT2:%.*]], ptr addrspace(1) inreg [[OUT3:%.*]]) #[[ATTR0]] {
; PRELOAD-8-NEXT: [[TEST_PRELOAD_IR_LOWERING_KERNEL_8_KERNARG_SEGMENT:%.*]] = call nonnull align 16 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.kernarg.segment.ptr()
-; PRELOAD-8-NEXT: [[OUT2_KERNARG_OFFSET:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[TEST_PRELOAD_IR_LOWERING_KERNEL_8_KERNARG_SEGMENT]], i64 48
-; PRELOAD-8-NEXT: [[OUT2_LOAD:%.*]] = load ptr addrspace(1), ptr addrspace(4) [[OUT2_KERNARG_OFFSET]], align 16, !invariant.load [[META0:![0-9]+]]
; PRELOAD-8-NEXT: [[OUT3_KERNARG_OFFSET:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[TEST_PRELOAD_IR_LOWERING_KERNEL_8_KERNARG_SEGMENT]], i64 56
-; PRELOAD-8-NEXT: [[OUT3_LOAD:%.*]] = load ptr addrspace(1), ptr addrspace(4) [[OUT3_KERNARG_OFFSET]], align 8, !invariant.load [[META0]]
+; PRELOAD-8-NEXT: [[OUT3_LOAD:%.*]] = load ptr addrspace(1), ptr addrspace(4) [[OUT3_KERNARG_OFFSET]], align 8, !invariant.load [[META0:![0-9]+]]
; PRELOAD-8-NEXT: [[LOAD:%.*]] = load i32, ptr addrspace(1) [[IN]], align 4
; PRELOAD-8-NEXT: [[LOAD1:%.*]] = load i32, ptr addrspace(1) [[IN1]], align 4
; PRELOAD-8-NEXT: [[LOAD2:%.*]] = load i32, ptr addrspace(1) [[IN2]], align 4
; PRELOAD-8-NEXT: [[LOAD3:%.*]] = load i32, ptr addrspace(1) [[IN3]], align 4
; PRELOAD-8-NEXT: store i32 [[LOAD]], ptr addrspace(1) [[OUT]], align 4
; PRELOAD-8-NEXT: store i32 [[LOAD1]], ptr addrspace(1) [[OUT1]], align 4
-; PRELOAD-8-NEXT: store i32 [[LOAD2]], ptr addrspace(1) [[OUT2_LOAD]], align 4
+; PRELOAD-8-NEXT: store i32 [[LOAD2]], ptr addrspace(1) [[OUT2]], align 4
; PRELOAD-8-NEXT: store i32 [[LOAD3]], ptr addrspace(1) [[OUT3_LOAD]], align 4
; PRELOAD-8-NEXT: ret void
;
>From 2b5ad125a8f100afea9a4b37a6e1c01c8c766f47 Mon Sep 17 00:00:00 2001
From: Austin Kerbow <Austin.Kerbow at amd.com>
Date: Thu, 5 Dec 2024 15:11:29 -0800
Subject: [PATCH 3/4] temp
---
llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp | 9 ++++-----
.../Target/AMDGPU/AMDGPULowerKernelArguments.cpp | 6 +++---
llvm/lib/Target/AMDGPU/GCNSubtarget.cpp | 11 ++---------
llvm/lib/Target/AMDGPU/GCNSubtarget.h | 15 +++------------
llvm/lib/Target/AMDGPU/SIISelLowering.cpp | 16 +++++++---------
llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp | 5 ++++-
llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h | 3 +++
...invalid-hidden-kernarg-in-kernel-signature.ll | 8 ++++----
8 files changed, 30 insertions(+), 43 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp b/llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp
index 64f0e773a1f27b..9bc02ba01880bc 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp
@@ -520,14 +520,13 @@ bool AMDGPUCallLowering::lowerFormalArgumentsKernel(
// TODO: Align down to dword alignment and extract bits for extending loads.
for (auto &Arg : F.args()) {
- // Hidden arguments that are in the kernel signature must be preloded to
- // user SGPRs, or loaded via the implicit_arg ptr. Print a diagnostic
- // error if a hidden argument is in the argument list and is not
- // preloaded.
+ // Hidden arguments that are in the kernel signature must be preloaded to
+ // user SGPRs, or loaded via the implicit_arg ptr. Print a diagnostic error
+ // if a hidden argument is in the argument list and is not preloaded.
if (Arg.hasAttribute("amdgpu-hidden-argument")) {
DiagnosticInfoUnsupported NonPreloadHiddenArg(
*Arg.getParent(),
- "Hidden argument in kernel signature was not preloaded");
+ "hidden argument in kernel signature was not preloaded");
F.getContext().diagnose(NonPreloadHiddenArg);
}
diff --git a/llvm/lib/Target/AMDGPU/AMDGPULowerKernelArguments.cpp b/llvm/lib/Target/AMDGPU/AMDGPULowerKernelArguments.cpp
index 3b3f8a72cc4a3c..e9d009baa20af2 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPULowerKernelArguments.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPULowerKernelArguments.cpp
@@ -145,7 +145,7 @@ class PreloadKernelArgInfo {
// arguments.
void setInitialFreeUserSGPRsCount() {
GCNUserSGPRUsageInfo UserSGPRInfo(F, ST);
- NumFreeUserSGPRs = UserSGPRInfo.getNumFreeKernargPreloadSGPRs();
+ NumFreeUserSGPRs = UserSGPRInfo.getNumFreeUserSGPRs();
}
bool tryAllocPreloadSGPRs(unsigned AllocSize, uint64_t ArgOffset,
@@ -486,8 +486,8 @@ static bool lowerKernelArguments(Function &F, const TargetMachine &TM) {
uint64_t ImplicitArgsBaseOffset =
alignTo(ExplicitArgOffset, ST.getAlignmentForImplicitArgPtr()) +
BaseOffset;
- PreloadInfo.tryAllocImplicitArgPreloadSGPRs(ImplicitArgsBaseOffset, ExplicitArgOffset,
- Builder);
+ PreloadInfo.tryAllocImplicitArgPreloadSGPRs(ImplicitArgsBaseOffset,
+ ExplicitArgOffset, Builder);
}
return true;
diff --git a/llvm/lib/Target/AMDGPU/GCNSubtarget.cpp b/llvm/lib/Target/AMDGPU/GCNSubtarget.cpp
index 4e65e6aa67d9a8..51361b75940560 100644
--- a/llvm/lib/Target/AMDGPU/GCNSubtarget.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNSubtarget.cpp
@@ -748,10 +748,6 @@ GCNUserSGPRUsageInfo::GCNUserSGPRUsageInfo(const Function &F,
FlatScratchInit = true;
}
- if (!AMDGPU::isGraphics(CC) && !IsKernel &&
- !F.hasFnAttribute("amdgpu-no-lds-kernel-id"))
- LDSKernelId = true;
-
if (hasImplicitBufferPtr())
NumUsedUserSGPRs += getNumUserSGPRForField(ImplicitBufferPtrID);
@@ -775,9 +771,6 @@ GCNUserSGPRUsageInfo::GCNUserSGPRUsageInfo(const Function &F,
if (hasPrivateSegmentSize())
NumUsedUserSGPRs += getNumUserSGPRForField(PrivateSegmentSizeID);
-
- if (hasLDSKernelId())
- NumSyntheticSGPRs += getNumUserSGPRForField(LDSKernelIdID);
}
void GCNUserSGPRUsageInfo::allocKernargPreloadSGPRs(unsigned NumSGPRs) {
@@ -786,6 +779,6 @@ void GCNUserSGPRUsageInfo::allocKernargPreloadSGPRs(unsigned NumSGPRs) {
NumUsedUserSGPRs += NumSGPRs;
}
-unsigned GCNUserSGPRUsageInfo::getNumFreeKernargPreloadSGPRs() {
- return AMDGPU::getMaxNumUserSGPRs(ST) - (NumUsedUserSGPRs + NumSyntheticSGPRs);
+unsigned GCNUserSGPRUsageInfo::getNumFreeUserSGPRs() {
+ return AMDGPU::getMaxNumUserSGPRs(ST) - NumUsedUserSGPRs;
}
diff --git a/llvm/lib/Target/AMDGPU/GCNSubtarget.h b/llvm/lib/Target/AMDGPU/GCNSubtarget.h
index a179880dfae8c8..5cecaf6349c883 100644
--- a/llvm/lib/Target/AMDGPU/GCNSubtarget.h
+++ b/llvm/lib/Target/AMDGPU/GCNSubtarget.h
@@ -1678,13 +1678,11 @@ class GCNUserSGPRUsageInfo {
bool hasPrivateSegmentSize() const { return PrivateSegmentSize; }
- bool hasLDSKernelId() const { return LDSKernelId; }
-
unsigned getNumKernargPreloadSGPRs() const { return NumKernargPreloadSGPRs; }
unsigned getNumUsedUserSGPRs() const { return NumUsedUserSGPRs; }
- unsigned getNumFreeKernargPreloadSGPRs();
+ unsigned getNumFreeUserSGPRs();
void allocKernargPreloadSGPRs(unsigned NumSGPRs);
@@ -1696,12 +1694,11 @@ class GCNUserSGPRUsageInfo {
KernargSegmentPtrID = 4,
DispatchIdID = 5,
FlatScratchInitID = 6,
- PrivateSegmentSizeID = 7,
- LDSKernelIdID = 8
+ PrivateSegmentSizeID = 7
};
// Returns the size in number of SGPRs for preload user SGPR field.
- static constexpr unsigned getNumUserSGPRForField(UserSGPRID ID) {
+ static unsigned getNumUserSGPRForField(UserSGPRID ID) {
switch (ID) {
case ImplicitBufferPtrID:
return 2;
@@ -1719,8 +1716,6 @@ class GCNUserSGPRUsageInfo {
return 2;
case PrivateSegmentSizeID:
return 1;
- case LDSKernelIdID:
- return 1;
}
llvm_unreachable("Unknown UserSGPRID.");
}
@@ -1749,13 +1744,9 @@ class GCNUserSGPRUsageInfo {
bool PrivateSegmentSize = false;
- bool LDSKernelId = false;
-
unsigned NumKernargPreloadSGPRs = 0;
unsigned NumUsedUserSGPRs = 0;
-
- unsigned NumSyntheticSGPRs = 0;
};
} // end namespace llvm
diff --git a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
index e6d1e96bb2a1df..d0ccca6e1ea230 100644
--- a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
+++ b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
@@ -2422,7 +2422,7 @@ void SITargetLowering::allocateSpecialInputSGPRs(
if (Info.hasWorkGroupIDZ())
allocateSGPR32Input(CCInfo, ArgInfo.WorkGroupIDZ);
- if (UserSGPRInfo.hasLDSKernelId())
+ if (Info.hasLDSKernelId())
allocateSGPR32Input(CCInfo, ArgInfo.LDSKernelId);
}
@@ -2545,8 +2545,7 @@ void SITargetLowering::allocatePreloadKernArgSGPRs(
unsigned Padding = ArgOffset - LastExplicitArgOffset;
unsigned PaddingSGPRs = alignTo(Padding, 4) / 4;
// Check for free user SGPRs for preloading.
- if (PaddingSGPRs + NumAllocSGPRs >
- SGPRInfo.getNumFreeKernargPreloadSGPRs()) {
+ if (PaddingSGPRs + NumAllocSGPRs > SGPRInfo.getNumFreeUserSGPRs()) {
InPreloadSequence = false;
break;
}
@@ -2574,8 +2573,7 @@ void SITargetLowering::allocateLDSKernelId(CCState &CCInfo, MachineFunction &MF,
const SIRegisterInfo &TRI,
SIMachineFunctionInfo &Info) const {
// Always allocate this last since it is a synthetic preload.
- const GCNUserSGPRUsageInfo &UserSGPRInfo = Info.getUserSGPRInfo();
- if (UserSGPRInfo.hasLDSKernelId()) {
+ if (Info.hasLDSKernelId()) {
Register Reg = Info.addLDSKernelId();
MF.addLiveIn(Reg, &AMDGPU::SGPR_32RegClass);
CCInfo.AllocateReg(Reg);
@@ -2825,7 +2823,7 @@ SDValue SITargetLowering::LowerFormalArguments(
const GCNUserSGPRUsageInfo &UserSGPRInfo = Info->getUserSGPRInfo();
assert(!UserSGPRInfo.hasDispatchPtr() &&
!UserSGPRInfo.hasKernargSegmentPtr() && !Info->hasWorkGroupInfo() &&
- !UserSGPRInfo.hasLDSKernelId() && !Info->hasWorkItemIDX() &&
+ !Info->hasLDSKernelId() && !Info->hasWorkItemIDX() &&
!Info->hasWorkItemIDY() && !Info->hasWorkItemIDZ());
(void)UserSGPRInfo;
if (!Subtarget->enableFlatScratch())
@@ -3025,8 +3023,8 @@ SDValue SITargetLowering::LowerFormalArguments(
NewArg = DAG.getMergeValues({NewArg, Chain}, DL);
}
} else {
- // Hidden arguments that are in the kernel signature must be preloded to
- // user SGPRs, or loaded via the implicit_arg ptr. Print a diagnostic
+ // Hidden arguments that are in the kernel signature must be preloaded
+ // to user SGPRs, or loaded via the implicit_arg ptr. Print a diagnostic
// error if a hidden argument is in the argument list and is not
// preloaded.
if (Arg.isOrigArg()) {
@@ -3034,7 +3032,7 @@ SDValue SITargetLowering::LowerFormalArguments(
if (OrigArg->hasAttribute("amdgpu-hidden-argument")) {
DiagnosticInfoUnsupported NonPreloadHiddenArg(
*OrigArg->getParent(),
- "Hidden argument in kernel signature was not preloaded",
+ "hidden argument in kernel signature was not preloaded",
DL.getDebugLoc());
DAG.getContext()->diagnose(NonPreloadHiddenArg);
}
diff --git a/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp b/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp
index 3013230e9de78c..1e43d2727a00da 100644
--- a/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp
@@ -38,7 +38,7 @@ SIMachineFunctionInfo::SIMachineFunctionInfo(const Function &F,
const GCNSubtarget *STI)
: AMDGPUMachineFunction(F, *STI), Mode(F, *STI), GWSResourcePSV(getTM(STI)),
UserSGPRInfo(F, *STI), WorkGroupIDX(false), WorkGroupIDY(false),
- WorkGroupIDZ(false), WorkGroupInfo(false),
+ WorkGroupIDZ(false), WorkGroupInfo(false), LDSKernelId(false),
PrivateSegmentWaveByteOffset(false), WorkItemIDX(false),
WorkItemIDY(false), WorkItemIDZ(false), ImplicitArgPtr(false),
GITPtrHigh(0xffffffff), HighBitsOf32BitAddress(0) {
@@ -131,6 +131,9 @@ SIMachineFunctionInfo::SIMachineFunctionInfo(const Function &F,
if (!F.hasFnAttribute("amdgpu-no-workitem-id-z") &&
ST.getMaxWorkitemID(F, 2) != 0)
WorkItemIDZ = true;
+
+ if (!IsKernel && !F.hasFnAttribute("amdgpu-no-lds-kernel-id"))
+ LDSKernelId = true;
}
if (isEntryFunction()) {
diff --git a/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h b/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h
index 2877a3c4fe38ca..2a754680fdc8ca 100644
--- a/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h
+++ b/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h
@@ -461,6 +461,7 @@ class SIMachineFunctionInfo final : public AMDGPUMachineFunction,
bool WorkGroupIDY : 1;
bool WorkGroupIDZ : 1;
bool WorkGroupInfo : 1;
+ bool LDSKernelId : 1;
bool PrivateSegmentWaveByteOffset : 1;
bool WorkItemIDX : 1; // Always initialized.
@@ -821,6 +822,8 @@ class SIMachineFunctionInfo final : public AMDGPUMachineFunction,
return ArgInfo.WorkGroupInfo.getRegister();
}
+ bool hasLDSKernelId() const { return LDSKernelId; }
+
// Add special VGPR inputs
void setWorkItemIDX(ArgDescriptor Arg) {
ArgInfo.WorkItemIDX = Arg;
diff --git a/llvm/test/CodeGen/AMDGPU/invalid-hidden-kernarg-in-kernel-signature.ll b/llvm/test/CodeGen/AMDGPU/invalid-hidden-kernarg-in-kernel-signature.ll
index b90da58de95818..67a72ad9558636 100644
--- a/llvm/test/CodeGen/AMDGPU/invalid-hidden-kernarg-in-kernel-signature.ll
+++ b/llvm/test/CodeGen/AMDGPU/invalid-hidden-kernarg-in-kernel-signature.ll
@@ -1,14 +1,14 @@
-; RUN: not llc -global-isel=1 -mtriple=amdgcn--amdhsa -mcpu=gfx940 -start-after=amdgpu-lower-kernel-arguments < %s 2>&1 | FileCheck -check-prefix=ERROR %s
-; RUN: not llc -global-isel=0 -mtriple=amdgcn--amdhsa -mcpu=gfx940 -start-after=amdgpu-lower-kernel-arguments < %s 2>&1 | FileCheck -check-prefix=ERROR %s
+; RUN: not llc -global-isel=1 -mtriple=amdgcn--amdhsa -mcpu=gfx940 < %s 2>&1 | FileCheck -check-prefix=ERROR %s
+; RUN: not llc -global-isel=0 -mtriple=amdgcn--amdhsa -mcpu=gfx940 < %s 2>&1 | FileCheck -check-prefix=ERROR %s
define amdgpu_kernel void @no_free_sgprs_block_count_x_no_preload_diag(ptr addrspace(1) inreg %out, i512 inreg, i32 inreg "amdgpu-hidden-argument" %_hidden_block_count_x) #0 {
-; ERROR: error: <unknown>:0:0: in function no_free_sgprs_block_count_x_no_preload_diag void (ptr addrspace(1), i512, i32): Hidden argument in kernel signature was not preloaded
+; ERROR: error: <unknown>:0:0: in function no_free_sgprs_block_count_x_no_preload_diag void (ptr addrspace(1), i512, i32): hidden argument in kernel signature was not preloaded
store i32 %_hidden_block_count_x, ptr addrspace(1) %out
ret void
}
define amdgpu_kernel void @preloadremainder_z_no_preload_diag(ptr addrspace(1) inreg %out, i256 inreg, i32 inreg "amdgpu-hidden-argument" %_hidden_block_count_x, i32 inreg "amdgpu-hidden-argument" %_hidden_block_count_y, i32 inreg "amdgpu-hidden-argument" %_hidden_block_count_z, i16 inreg "amdgpu-hidden-argument" %_hidden_group_size_x, i16 inreg "amdgpu-hidden-argument" %_hidden_group_size_y, i16 inreg "amdgpu-hidden-argument" %_hidden_group_size_z, i16 inreg "amdgpu-hidden-argument" %_hidden_remainder_x, i16 inreg "amdgpu-hidden-argument" %_hidden_remainder_y, i16 inreg "amdgpu-hidden-argument" %_hidden_remainder_z) #0 {
-; ERROR: error: <unknown>:0:0: in function preloadremainder_z_no_preload_diag void (ptr addrspace(1), i256, i32, i32, i32, i16, i16, i16, i16, i16, i16): Hidden argument in kernel signature was not preloaded
+; ERROR: error: <unknown>:0:0: in function preloadremainder_z_no_preload_diag void (ptr addrspace(1), i256, i32, i32, i32, i16, i16, i16, i16, i16, i16): hidden argument in kernel signature was not preloaded
%conv = zext i16 %_hidden_remainder_z to i32
store i32 %conv, ptr addrspace(1) %out
ret void
>From ec97a8b9b703a5454376decfee8e3a2ded38c30a Mon Sep 17 00:00:00 2001
From: Austin Kerbow <Austin.Kerbow at amd.com>
Date: Fri, 6 Dec 2024 13:28:26 -0800
Subject: [PATCH 4/4] Address comments.
---
llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp | 4 ++--
llvm/lib/Target/AMDGPU/SIISelLowering.cpp | 5 ++---
.../AMDGPU/invalid-hidden-kernarg-in-kernel-signature.ll | 6 ++++--
3 files changed, 8 insertions(+), 7 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp b/llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp
index 9bc02ba01880bc..06ac2f41f4496c 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp
@@ -521,8 +521,8 @@ bool AMDGPUCallLowering::lowerFormalArgumentsKernel(
// TODO: Align down to dword alignment and extract bits for extending loads.
for (auto &Arg : F.args()) {
// Hidden arguments that are in the kernel signature must be preloaded to
- // user SGPRs, or loaded via the implicit_arg ptr. Print a diagnostic error
- // if a hidden argument is in the argument list and is not preloaded.
+ // user SGPRs. Print a diagnostic error if a hidden argument is in the
+ // argument list and is not preloaded.
if (Arg.hasAttribute("amdgpu-hidden-argument")) {
DiagnosticInfoUnsupported NonPreloadHiddenArg(
*Arg.getParent(),
diff --git a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
index d0ccca6e1ea230..f3592add5c5a37 100644
--- a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
+++ b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
@@ -3024,9 +3024,8 @@ SDValue SITargetLowering::LowerFormalArguments(
}
} else {
// Hidden arguments that are in the kernel signature must be preloaded
- // to user SGPRs, or loaded via the implicit_arg ptr. Print a diagnostic
- // error if a hidden argument is in the argument list and is not
- // preloaded.
+ // to user SGPRs. Print a diagnostic error if a hidden argument is in
+ // the argument list and is not preloaded.
if (Arg.isOrigArg()) {
Argument *OrigArg = Fn.getArg(Arg.getOrigArgIndex());
if (OrigArg->hasAttribute("amdgpu-hidden-argument")) {
diff --git a/llvm/test/CodeGen/AMDGPU/invalid-hidden-kernarg-in-kernel-signature.ll b/llvm/test/CodeGen/AMDGPU/invalid-hidden-kernarg-in-kernel-signature.ll
index 67a72ad9558636..b51072e28a2b9d 100644
--- a/llvm/test/CodeGen/AMDGPU/invalid-hidden-kernarg-in-kernel-signature.ll
+++ b/llvm/test/CodeGen/AMDGPU/invalid-hidden-kernarg-in-kernel-signature.ll
@@ -1,5 +1,7 @@
-; RUN: not llc -global-isel=1 -mtriple=amdgcn--amdhsa -mcpu=gfx940 < %s 2>&1 | FileCheck -check-prefix=ERROR %s
-; RUN: not llc -global-isel=0 -mtriple=amdgcn--amdhsa -mcpu=gfx940 < %s 2>&1 | FileCheck -check-prefix=ERROR %s
+; RUN: not llc -global-isel=1 -mtriple=amdgcn--amdhsa -mcpu=gfx942 < %s 2>&1 | FileCheck -check-prefix=ERROR %s
+; RUN: not llc -global-isel=0 -mtriple=amdgcn--amdhsa -mcpu=gfx942 < %s 2>&1 | FileCheck -check-prefix=ERROR %s
+; RUN: not llc -global-isel=1 -amdgpu-ir-lower-kernel-arguments=0 -mtriple=amdgcn--amdhsa -mcpu=gfx942 < %s 2>&1 | FileCheck -check-prefix=ERROR %s
+; RUN: not llc -global-isel=0 -amdgpu-ir-lower-kernel-arguments=0 -mtriple=amdgcn--amdhsa -mcpu=gfx942 < %s 2>&1 | FileCheck -check-prefix=ERROR %s
define amdgpu_kernel void @no_free_sgprs_block_count_x_no_preload_diag(ptr addrspace(1) inreg %out, i512 inreg, i32 inreg "amdgpu-hidden-argument" %_hidden_block_count_x) #0 {
; ERROR: error: <unknown>:0:0: in function no_free_sgprs_block_count_x_no_preload_diag void (ptr addrspace(1), i512, i32): hidden argument in kernel signature was not preloaded
More information about the llvm-commits
mailing list