[AMDGPU] Fix hidden kernarg preload count inconsistency #116759

kerbowa · 2024-11-19T07:40:49Z

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.

llvmbot · 2024-11-19T07:41:23Z

@llvm/pr-subscribers-backend-amdgpu

Author: Austin Kerbow (kerbowa)

Changes

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.

Full diff: https://github.com/llvm/llvm-project/pull/116759.diff

4 Files Affected:

(modified) llvm/lib/Target/AMDGPU/AMDGPULowerKernelArguments.cpp (+6-6)
(modified) llvm/lib/Target/AMDGPU/SIISelLowering.cpp (+8)
(modified) llvm/test/CodeGen/AMDGPU/preload-implicit-kernargs.ll (+44-15)
(modified) llvm/test/CodeGen/AMDGPU/preload-kernargs-IR-lowering.ll (+4-2)

diff --git a/llvm/lib/Target/AMDGPU/AMDGPULowerKernelArguments.cpp b/llvm/lib/Target/AMDGPU/AMDGPULowerKernelArguments.cpp
index 380633f22a1781..5ef4d3555a4eb4 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPULowerKernelArguments.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPULowerKernelArguments.cpp
@@ -143,17 +143,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.
@@ -169,6 +169,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);
@@ -214,7 +215,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.
@@ -474,7 +474,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 1e261f4256c93b..379a5b9bbf84fc 100644
--- a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
+++ b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
@@ -2999,6 +2999,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
 ;

github-actions · 2024-11-19T07:44:14Z

✅ With the latest revision this PR passed the C/C++ code formatter.

arsenm

It is possible that the number of hidden arguments that are selected to be preloaded in AMDGPULowerKernel arguments and isel can differ.

This should not happen. The introduction of arguments to preload should be precise

arsenm · 2024-11-21T04:27:25Z

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

+          assert(!OrigArg->hasAttribute("amdgpu-hidden-argument") &&
+                 "Hidden arguments should be preloaded");


No error on valid IR can be an assert. This should report a proper error (via DiagnosticInfo)

arsenm · 2024-11-21T04:27:56Z

llvm/lib/Target/AMDGPU/AMDGPULowerKernelArguments.cpp

-
-    NumFreeUserSGPRs = MaxUserSGPRs - UserSGPRInfo.getNumUsedUserSGPRs();
+    NumFreeUserSGPRs =
+        UserSGPRInfo.getNumFreeUserSGPRs() - 1 /* Synthetic SGPRs*/;


Not sure what synthetic means here

We call LDSKernelId a synthetic SGPR which is allocated after normal user SGPRs but before system SGPRs. I just realized this doesn't actually matter for kernels since it should never be used there.

arsenm · 2024-12-03T21:48:12Z

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


Typo 'preloded'

arsenm · 2024-12-03T21:48:23Z

llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp

+    if (Arg.hasAttribute("amdgpu-hidden-argument")) {
+      DiagnosticInfoUnsupported NonPreloadHiddenArg(
+          *Arg.getParent(),
+          "Hidden argument in kernel signature was not preloaded");


Start error messages with lowercase

arsenm · 2024-12-03T21:50:08Z

llvm/lib/Target/AMDGPU/GCNSubtarget.cpp

+  if (!AMDGPU::isGraphics(CC) && !IsKernel &&
+      !F.hasFnAttribute("amdgpu-no-lds-kernel-id"))
+    LDSKernelId = true;
+


The kernel ID has to exist for all callable functions, !graphics will be wrong for amdgpu_gfx. Should be not entry function cc?

I've just removed the refactoring for LDSKernelID since it's not relevant for kernels, but FYI this just copied the existing logic from SIMachineFunction.cpp, so if it's incorrect we should update it there.

arsenm · 2024-12-03T21:51:10Z

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

+          if (OrigArg->hasAttribute("amdgpu-hidden-argument")) {
+            DiagnosticInfoUnsupported NonPreloadHiddenArg(
+                *OrigArg->getParent(),
+                "Hidden argument in kernel signature was not preloaded",


Start messages with lowercase

arsenm · 2024-12-03T21:51:48Z

llvm/lib/Target/AMDGPU/GCNSubtarget.h

@@ -1692,11 +1694,12 @@ class GCNUserSGPRUsageInfo {
    KernargSegmentPtrID = 4,
    DispatchIdID = 5,
    FlatScratchInitID = 6,
-    PrivateSegmentSizeID = 7
+    PrivateSegmentSizeID = 7,
+    LDSKernelIdID = 8


Comment this isn't really a user SGPR? Maybe should rename the enum

arsenm · 2024-12-03T21:52:56Z

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

+; 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


Shouldn't need the -start-after

arsenm · 2024-12-03T21:54:53Z

llvm/lib/Target/AMDGPU/GCNSubtarget.cpp

+  if (hasLDSKernelId())
+    NumSyntheticSGPRs += getNumUserSGPRForField(LDSKernelIdID);


I'm not sure I understand why anything is changing with the kernel ID. Is this just changing bookkeeping to split out the kernel's user SGPRs. vs. the implicit SGPRs in a function?

arsenm · 2024-12-03T21:55:04Z

llvm/lib/Target/AMDGPU/GCNSubtarget.h

  unsigned getNumKernargPreloadSGPRs() const { return NumKernargPreloadSGPRs; }

  unsigned getNumUsedUserSGPRs() const { return NumUsedUserSGPRs; }

-  unsigned getNumFreeUserSGPRs();
+  unsigned getNumFreeKernargPreloadSGPRs();


arsenm · 2024-12-03T21:55:42Z

llvm/lib/Target/AMDGPU/GCNSubtarget.h

@@ -1676,11 +1676,13 @@ class GCNUserSGPRUsageInfo {

  bool hasPrivateSegmentSize() const { return PrivateSegmentSize; }

+  bool hasLDSKernelId() const { return LDSKernelId; }


I don't know why this is defined in GCNSubtarget.h, it's not related to the subtarget

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.

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.

arsenm · 2024-12-06T16:24:40Z

llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp

+    // 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");
+      F.getContext().diagnose(NonPreloadHiddenArg);
+    }


This is an unconditonal error, and untested. Currently the IR kernel argument lowering is treated as optional, both selectors are supposed to fully handle kernel lowering independently with -amdgpu-ir-lower-kernel-arguments=0.

I also don't understand the second part of the comment. If they're loaded through the implicit pointer, they weren't preloaded?

We had decided a while back that hidden arguments should not always be in the kernel signature. They will only be present if they are preloaded into user SGPRs. This diagnostic is saying we should never see the attribute in gisel since preloading kernargs isn't implemented there yet, but I'm going to add that support next.

If a hidden argument in the signature and has the attribute amdgpu-hidden-argument uses of that argument must use the preloaded SGPR. There is nothing stopping other instances of the same argument from being accessed via the implicit_arg ptr in a callee for example.

We could easily make it so that isel can handle these even if they are not preloaded, but since it is not intended and it indicates a mismatch between isel and AMDGPULowerKernelArguments we error out instead.

I would expect globalisel to be at feature parity, and hit the fallback for unsupported features, not hard error

arsenm · 2024-12-06T16:25:35Z

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 < %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
+


Also test with the ir lowering disabled

arsenm · 2024-12-06T16:25:44Z

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

+; 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


arsenm · 2024-12-06T21:31:50Z

llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp

+      DiagnosticInfoUnsupported NonPreloadHiddenArg(
+          *Arg.getParent(),
+          "hidden argument in kernel signature was not preloaded");
+      F.getContext().diagnose(NonPreloadHiddenArg);


return false and comment (with debug print) this is just because it's not implemented in globalisel

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. Change-Id: Ib474daf86ee84913ed1f864f52c399f0a1480710

kerbowa requested a review from arsenm November 19, 2024 07:40

llvmbot added the backend:AMDGPU label Nov 19, 2024

arsenm reviewed Nov 21, 2024

View reviewed changes

kerbowa force-pushed the preload-hidden-arg-mismatch branch from b1a57c3 to f501081 Compare December 2, 2024 04:41

arsenm reviewed Dec 3, 2024

View reviewed changes

kerbowa added 3 commits December 5, 2024 15:00

Diagnostic error and move LDSKernelId.

85311bd

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.

temp

2b5ad12

kerbowa force-pushed the preload-hidden-arg-mismatch branch from f501081 to 2b5ad12 Compare December 6, 2024 16:15

arsenm reviewed Dec 6, 2024

View reviewed changes

Address comments.

ec97a8b

arsenm reviewed Dec 6, 2024

View reviewed changes

Abort in gisel.

4a3ed82

arsenm approved these changes Dec 8, 2024

View reviewed changes

kerbowa merged commit b1d4246 into llvm:main Dec 8, 2024
8 checks passed

		assert(!OrigArg->hasAttribute("amdgpu-hidden-argument") &&
		"Hidden arguments should be preloaded");

		; 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

		if (hasLDSKernelId())
		NumSyntheticSGPRs += getNumUserSGPRForField(LDSKernelIdID);

		@@ -1676,11 +1676,13 @@ class GCNUserSGPRUsageInfo {

		bool hasPrivateSegmentSize() const { return PrivateSegmentSize; }

		bool hasLDSKernelId() const { return LDSKernelId; }

		@@ -0,0 +1,17 @@
		; 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

[AMDGPU] Fix hidden kernarg preload count inconsistency #116759

[AMDGPU] Fix hidden kernarg preload count inconsistency #116759

Uh oh!

Conversation

kerbowa commented Nov 19, 2024

Uh oh!

llvmbot commented Nov 19, 2024

Uh oh!

github-actions bot commented Nov 19, 2024 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

arsenm left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

github-actions bot commented Nov 19, 2024 •

edited

Loading