[llvm] [AMDGPU][MC] update USER_SGPR_COUNT bits for GFX1250 (PR #192579)
Zeng Wu via llvm-commits
llvm-commits at lists.llvm.org
Wed Apr 22 22:49:17 PDT 2026
https://github.com/zwu-2025 updated https://github.com/llvm/llvm-project/pull/192579
>From e80756baf0365f1a63ff8f5fc516fbffd3548fb6 Mon Sep 17 00:00:00 2001
From: zwu-2025 <Zeng.Wu2 at amd.com>
Date: Fri, 17 Apr 2026 01:40:56 +0000
Subject: [PATCH 1/7] [AMDGPU-MC] update USER_SGPR_COUNT bits for GFX1250
---
llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp | 11 +++---
.../AMDGPU/AsmParser/AMDGPUAsmParser.cpp | 2 +-
llvm/lib/Target/AMDGPU/SIProgramInfo.cpp | 34 +++++++++++--------
llvm/lib/Target/AMDGPU/SIProgramInfo.h | 6 ++--
.../llvm.amdgcn.gfx125.user.sgpr.encoding.ll | 15 ++++++++
5 files changed, 46 insertions(+), 22 deletions(-)
create mode 100644 llvm/test/CodeGen/AMDGPU/llvm.amdgcn.gfx125.user.sgpr.encoding.ll
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp b/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
index 19ad73f4fb421..77457acfe1462 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
@@ -649,7 +649,7 @@ AMDGPUAsmPrinter::getAmdhsaKernelDescriptor(const MachineFunction &MF,
STM.getKernArgSegmentSize(F, MaxKernArgAlign), Ctx);
KernelDescriptor.compute_pgm_rsrc1 = PI.getComputePGMRSrc1(STM, Ctx);
- KernelDescriptor.compute_pgm_rsrc2 = PI.getComputePGMRSrc2(Ctx);
+ KernelDescriptor.compute_pgm_rsrc2 = PI.getComputePGMRSrc2(STM, Ctx);
KernelDescriptor.kernel_code_properties = getAmdhsaKernelCodeProperties(MF);
int64_t PGM_Rsrc3 = 1;
@@ -1371,7 +1371,8 @@ void AMDGPUAsmPrinter::EmitProgramInfoSI(const MachineFunction &MF,
/*Size=*/4);
OutStreamer->emitInt32(R_00B84C_COMPUTE_PGM_RSRC2);
- EmitResolvedOrExpr(CurrentProgramInfo.getComputePGMRSrc2(Ctx), /*Size=*/4);
+ EmitResolvedOrExpr(CurrentProgramInfo.getComputePGMRSrc2(STM, Ctx),
+ /*Size=*/4);
OutStreamer->emitInt32(R_00B860_COMPUTE_TMPRING_SIZE);
@@ -1497,7 +1498,7 @@ void AMDGPUAsmPrinter::EmitPALMetadata(const MachineFunction &MF,
if (MD->getPALMajorVersion() < 3) {
MD->setRsrc1(CC, CurrentProgramInfo.getPGMRSrc1(CC, STM, Ctx), Ctx);
if (AMDGPU::isCompute(CC)) {
- MD->setRsrc2(CC, CurrentProgramInfo.getComputePGMRSrc2(Ctx), Ctx);
+ MD->setRsrc2(CC, CurrentProgramInfo.getComputePGMRSrc2(STM, Ctx), Ctx);
} else {
const MCExpr *HasScratchBlocks =
MCBinaryExpr::createGT(CurrentProgramInfo.ScratchBlocks,
@@ -1579,7 +1580,7 @@ void AMDGPUAsmPrinter::emitPALFunctionMetadata(const MachineFunction &MF) {
CallingConv::AMDGPU_CS,
CurrentProgramInfo.getPGMRSrc1(CallingConv::AMDGPU_CS, ST, Ctx), Ctx);
MD->setRsrc2(CallingConv::AMDGPU_CS,
- CurrentProgramInfo.getComputePGMRSrc2(Ctx), Ctx);
+ CurrentProgramInfo.getComputePGMRSrc2(ST, Ctx), Ctx);
} else {
EmitPALMetadataCommon(
MD, CurrentProgramInfo, CallingConv::AMDGPU_CS, ST,
@@ -1622,7 +1623,7 @@ void AMDGPUAsmPrinter::getAmdKernelCode(AMDGPUMCKernelCodeT &Out,
Out.compute_pgm_resource1_registers =
CurrentProgramInfo.getComputePGMRSrc1(STM, Ctx);
Out.compute_pgm_resource2_registers =
- CurrentProgramInfo.getComputePGMRSrc2(Ctx);
+ CurrentProgramInfo.getComputePGMRSrc2(STM, Ctx);
Out.code_properties |= AMD_CODE_PROPERTY_IS_PTR64;
Out.is_dynamic_callstack = CurrentProgramInfo.DynamicCallStack;
diff --git a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
index 3777adc9790e8..ab6612e422541 100644
--- a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
+++ b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
@@ -6400,7 +6400,7 @@ bool AMDGPUAsmParser::ParseDirectiveAMDHSAKernel() {
COMPUTE_PGM_RSRC1_GRANULATED_WAVEFRONT_SGPR_COUNT, getContext());
if (ExplicitUserSGPRCount && ImpliedUserSGPRCount > *ExplicitUserSGPRCount)
- return TokError("amdgpu_user_sgpr_count smaller than than implied by "
+ return TokError("amdgpu_user_sgpr_count smaller than implied by "
"enabled user SGPRs");
if (isGFX1250Plus()) {
diff --git a/llvm/lib/Target/AMDGPU/SIProgramInfo.cpp b/llvm/lib/Target/AMDGPU/SIProgramInfo.cpp
index a3f261b87e80b..7a1120de95955 100644
--- a/llvm/lib/Target/AMDGPU/SIProgramInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/SIProgramInfo.cpp
@@ -136,17 +136,21 @@ static uint64_t getPGMRSrc1Reg(const SIProgramInfo &ProgInfo,
return Reg;
}
-static uint64_t getComputePGMRSrc2Reg(const SIProgramInfo &ProgInfo) {
- uint64_t Reg = S_00B84C_USER_SGPR(ProgInfo.UserSGPR) |
- S_00B84C_TRAP_HANDLER(ProgInfo.TrapHandlerEnable) |
- S_00B84C_TGID_X_EN(ProgInfo.TGIdXEnable) |
- S_00B84C_TGID_Y_EN(ProgInfo.TGIdYEnable) |
- S_00B84C_TGID_Z_EN(ProgInfo.TGIdZEnable) |
- S_00B84C_TG_SIZE_EN(ProgInfo.TGSizeEnable) |
- S_00B84C_TIDIG_COMP_CNT(ProgInfo.TIdIGCompCount) |
- S_00B84C_EXCP_EN_MSB(ProgInfo.EXCPEnMSB) |
- S_00B84C_LDS_SIZE(ProgInfo.LdsSize) |
- S_00B84C_EXCP_EN(ProgInfo.EXCPEnable);
+static uint64_t getComputePGMRSrc2Reg(const GCNSubtarget &ST,
+ const SIProgramInfo &ProgInfo) {
+ uint64_t Reg = ST.hasGFX1250Insts()
+ ? ((ProgInfo.UserSGPR & 0x3F) << 0x1)
+ : (S_00B84C_USER_SGPR(ProgInfo.UserSGPR) |
+ S_00B84C_TRAP_HANDLER(ProgInfo.TrapHandlerEnable));
+
+ Reg |= S_00B84C_TGID_X_EN(ProgInfo.TGIdXEnable) |
+ S_00B84C_TGID_Y_EN(ProgInfo.TGIdYEnable) |
+ S_00B84C_TGID_Z_EN(ProgInfo.TGIdZEnable) |
+ S_00B84C_TG_SIZE_EN(ProgInfo.TGSizeEnable) |
+ S_00B84C_TIDIG_COMP_CNT(ProgInfo.TIdIGCompCount) |
+ S_00B84C_EXCP_EN_MSB(ProgInfo.EXCPEnMSB) |
+ S_00B84C_LDS_SIZE(ProgInfo.LdsSize) |
+ S_00B84C_EXCP_EN(ProgInfo.EXCPEnable);
return Reg;
}
@@ -189,16 +193,18 @@ const MCExpr *SIProgramInfo::getPGMRSrc1(CallingConv::ID CC,
return MCBinaryExpr::createOr(RegExpr, Res, Ctx);
}
-const MCExpr *SIProgramInfo::getComputePGMRSrc2(MCContext &Ctx) const {
- uint64_t Reg = getComputePGMRSrc2Reg(*this);
+const MCExpr *SIProgramInfo::getComputePGMRSrc2(const GCNSubtarget &ST,
+ MCContext &Ctx) const {
+ uint64_t Reg = getComputePGMRSrc2Reg(ST, *this);
const MCExpr *RegExpr = MCConstantExpr::create(Reg, Ctx);
return MCBinaryExpr::createOr(ScratchEnable, RegExpr, Ctx);
}
const MCExpr *SIProgramInfo::getPGMRSrc2(CallingConv::ID CC,
+ const GCNSubtarget &ST,
MCContext &Ctx) const {
if (AMDGPU::isCompute(CC))
- return getComputePGMRSrc2(Ctx);
+ return getComputePGMRSrc2(ST, Ctx);
return MCConstantExpr::create(0, Ctx);
}
diff --git a/llvm/lib/Target/AMDGPU/SIProgramInfo.h b/llvm/lib/Target/AMDGPU/SIProgramInfo.h
index 171c4a313a53b..947b473142a1f 100644
--- a/llvm/lib/Target/AMDGPU/SIProgramInfo.h
+++ b/llvm/lib/Target/AMDGPU/SIProgramInfo.h
@@ -117,8 +117,10 @@ struct LLVM_EXTERNAL_VISIBILITY SIProgramInfo {
MCContext &Ctx) const;
/// Compute the value of the ComputePGMRsrc2 register.
- const MCExpr *getComputePGMRSrc2(MCContext &Ctx) const;
- const MCExpr *getPGMRSrc2(CallingConv::ID CC, MCContext &Ctx) const;
+ const MCExpr *getComputePGMRSrc2(const GCNSubtarget &ST,
+ MCContext &Ctx) const;
+ const MCExpr *getPGMRSrc2(CallingConv::ID CC, const GCNSubtarget &ST,
+ MCContext &Ctx) const;
};
} // namespace llvm
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.gfx125.user.sgpr.encoding.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.gfx125.user.sgpr.encoding.ll
new file mode 100644
index 0000000000000..ac0578b883b49
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.gfx125.user.sgpr.encoding.ll
@@ -0,0 +1,15 @@
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 < %s -o - | llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx1250 -filetype=obj -o - | llvm-readobj --notes -
+
+; CHECK: .sgpr_count: 32
+define amdgpu_kernel void @many_inreg_i32(
+ i32 inreg %a0, i32 inreg %a1, i32 inreg %a2, i32 inreg %a3,
+ i32 inreg %a4, i32 inreg %a5, i32 inreg %a6, i32 inreg %a7,
+ i32 inreg %a8, i32 inreg %a9, i32 inreg %a10, i32 inreg %a11,
+ i32 inreg %a12, i32 inreg %a13, i32 inreg %a14, i32 inreg %a15,
+ i32 inreg %a16, i32 inreg %a17, i32 inreg %a18, i32 inreg %a19,
+ i32 inreg %a20, i32 inreg %a21, i32 inreg %a22, i32 inreg %a23,
+ i32 inreg %a24, i32 inreg %a25, i32 inreg %a26, i32 inreg %a27,
+ i32 inreg %a28, i32 inreg %a29, i32 inreg %a30, i32 inreg %a31,
+ i32 inreg %a32, i32 inreg %a33, i32 inreg %a34, i32 inreg %a35) {
+ ret void
+}
>From 27eae83a105e17eed781a0bacc92b5a687004b99 Mon Sep 17 00:00:00 2001
From: zwu-2025 <Zeng.Wu2 at amd.com>
Date: Fri, 17 Apr 2026 16:52:15 +0000
Subject: [PATCH 2/7] comments
---
llvm/lib/Target/AMDGPU/SIDefines.h | 4 +
llvm/lib/Target/AMDGPU/SIProgramInfo.cpp | 2 +-
.../llvm.amdgcn.gfx125.user.sgpr.encoding.s | 326 ++++++++++++++++++
3 files changed, 331 insertions(+), 1 deletion(-)
create mode 100644 llvm/test/CodeGen/AMDGPU/llvm.amdgcn.gfx125.user.sgpr.encoding.s
diff --git a/llvm/lib/Target/AMDGPU/SIDefines.h b/llvm/lib/Target/AMDGPU/SIDefines.h
index 5612b0df7aa9f..9867f5b02d7a7 100644
--- a/llvm/lib/Target/AMDGPU/SIDefines.h
+++ b/llvm/lib/Target/AMDGPU/SIDefines.h
@@ -1211,6 +1211,10 @@ enum {
#define G_00B84C_EXCP_EN(x) (((x) >> 24) & 0x7F)
#define C_00B84C_EXCP_EN 0x80FFFFFF
+#define S_00B84C_USER_SGPR_GFX1250(x) (((x) & 0x3F) << 1)
+#define G_00B84C_USER_SGPR_GFX1250(x) (((x) >> 1) & 0x3F)
+#define C_00B84C_USER_SGPR_GFX1250 0xFFFFFF81
+
#define R_0286CC_SPI_PS_INPUT_ENA 0x0286CC
#define R_0286D0_SPI_PS_INPUT_ADDR 0x0286D0
diff --git a/llvm/lib/Target/AMDGPU/SIProgramInfo.cpp b/llvm/lib/Target/AMDGPU/SIProgramInfo.cpp
index 7a1120de95955..644e080db2ab5 100644
--- a/llvm/lib/Target/AMDGPU/SIProgramInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/SIProgramInfo.cpp
@@ -139,7 +139,7 @@ static uint64_t getPGMRSrc1Reg(const SIProgramInfo &ProgInfo,
static uint64_t getComputePGMRSrc2Reg(const GCNSubtarget &ST,
const SIProgramInfo &ProgInfo) {
uint64_t Reg = ST.hasGFX1250Insts()
- ? ((ProgInfo.UserSGPR & 0x3F) << 0x1)
+ ? S_00B84C_USER_SGPR_GFX1250(ProgInfo.UserSGPR)
: (S_00B84C_USER_SGPR(ProgInfo.UserSGPR) |
S_00B84C_TRAP_HANDLER(ProgInfo.TrapHandlerEnable));
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.gfx125.user.sgpr.encoding.s b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.gfx125.user.sgpr.encoding.s
new file mode 100644
index 0000000000000..81b4584ae785a
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.gfx125.user.sgpr.encoding.s
@@ -0,0 +1,326 @@
+ .amdgcn_target "amdgcn-amd-amdhsa--gfx1250"
+ .amdhsa_code_object_version 6
+ .text
+ .globl many_inreg_i32 ; -- Begin function many_inreg_i32
+ .p2align 8
+ .type many_inreg_i32, at function
+many_inreg_i32: ; @many_inreg_i32
+; %bb.0:
+ s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+ s_endpgm
+ .section .rodata,"a", at progbits
+ .p2align 6, 0x0
+ .amdhsa_kernel many_inreg_i32
+ .amdhsa_group_segment_fixed_size 0
+ .amdhsa_private_segment_fixed_size 0
+ .amdhsa_kernarg_size 400
+ .amdhsa_user_sgpr_count 32
+ .amdhsa_user_sgpr_dispatch_ptr 1
+ .amdhsa_user_sgpr_queue_ptr 1
+ .amdhsa_user_sgpr_kernarg_segment_ptr 1
+ .amdhsa_user_sgpr_dispatch_id 1
+ .amdhsa_user_sgpr_kernarg_preload_length 24
+ .amdhsa_user_sgpr_kernarg_preload_offset 0
+ .amdhsa_user_sgpr_private_segment_size 0
+ .amdhsa_wavefront_size32 1
+ .amdhsa_uses_dynamic_stack 0
+ .amdhsa_enable_private_segment 0
+ .amdhsa_system_sgpr_workgroup_id_x 1
+ .amdhsa_system_sgpr_workgroup_id_y 1
+ .amdhsa_system_sgpr_workgroup_id_z 1
+ .amdhsa_system_sgpr_workgroup_info 0
+ .amdhsa_system_vgpr_workitem_id 2
+ .amdhsa_next_free_vgpr 1
+ .amdhsa_next_free_sgpr 32
+ .amdhsa_named_barrier_count 0
+ .amdhsa_reserve_vcc 0
+ .amdhsa_float_round_mode_32 0
+ .amdhsa_float_round_mode_16_64 0
+ .amdhsa_float_denorm_mode_32 3
+ .amdhsa_float_denorm_mode_16_64 3
+ .amdhsa_fp16_overflow 0
+ .amdhsa_memory_ordered 1
+ .amdhsa_forward_progress 1
+ .amdhsa_inst_pref_size 1
+ .amdhsa_round_robin_scheduling 0
+ .amdhsa_exception_fp_ieee_invalid_op 0
+ .amdhsa_exception_fp_denorm_src 0
+ .amdhsa_exception_fp_ieee_div_zero 0
+ .amdhsa_exception_fp_ieee_overflow 0
+ .amdhsa_exception_fp_ieee_underflow 0
+ .amdhsa_exception_fp_ieee_inexact 0
+ .amdhsa_exception_int_div_zero 0
+ .end_amdhsa_kernel
+ .text
+.Lfunc_end0:
+ .size many_inreg_i32, .Lfunc_end0-many_inreg_i32
+ ; -- End function
+ .set many_inreg_i32.num_vgpr, 0
+ .set many_inreg_i32.num_agpr, 0
+ .set many_inreg_i32.numbered_sgpr, 0
+ .set many_inreg_i32.num_named_barrier, 0
+ .set many_inreg_i32.private_seg_size, 0
+ .set many_inreg_i32.uses_vcc, 0
+ .set many_inreg_i32.uses_flat_scratch, 0
+ .set many_inreg_i32.has_dyn_sized_stack, 0
+ .set many_inreg_i32.has_recursion, 0
+ .set many_inreg_i32.has_indirect_call, 0
+ .section .AMDGPU.csdata,"", at progbits
+; Kernel info:
+; codeLenInByte = 12
+; TotalNumSgprs: 32
+; NumVgprs: 0
+; ScratchSize: 0
+; MemoryBound: 0
+; FloatMode: 240
+; IeeeMode: 1
+; LDSByteSize: 0 bytes/workgroup (compile time only)
+; SGPRBlocks: 0
+; VGPRBlocks: 0
+; NumSGPRsForWavesPerEU: 32
+; NumVGPRsForWavesPerEU: 1
+; NamedBarCnt: 0
+; Occupancy: 16
+; WaveLimiterHint : 0
+; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
+; COMPUTE_PGM_RSRC2:USER_SGPR: 32
+; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
+; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
+; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
+; COMPUTE_PGM_RSRC2:TGID_Z_EN: 1
+; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 2
+ .text
+ .p2alignl 7, 3214868480
+ .fill 96, 4, 3214868480
+ .section .AMDGPU.gpr_maximums,"", at progbits
+ .set amdgpu.max_num_vgpr, 0
+ .set amdgpu.max_num_agpr, 0
+ .set amdgpu.max_num_sgpr, 0
+ .set amdgpu.max_num_named_barrier, 0
+ .text
+ .section ".note.GNU-stack","", at progbits
+ .amdgpu_metadata
+---
+amdhsa.kernels:
+ - .args:
+ - .name: a0
+ .offset: 0
+ .size: 4
+ .value_kind: by_value
+ - .name: a1
+ .offset: 4
+ .size: 4
+ .value_kind: by_value
+ - .name: a2
+ .offset: 8
+ .size: 4
+ .value_kind: by_value
+ - .name: a3
+ .offset: 12
+ .size: 4
+ .value_kind: by_value
+ - .name: a4
+ .offset: 16
+ .size: 4
+ .value_kind: by_value
+ - .name: a5
+ .offset: 20
+ .size: 4
+ .value_kind: by_value
+ - .name: a6
+ .offset: 24
+ .size: 4
+ .value_kind: by_value
+ - .name: a7
+ .offset: 28
+ .size: 4
+ .value_kind: by_value
+ - .name: a8
+ .offset: 32
+ .size: 4
+ .value_kind: by_value
+ - .name: a9
+ .offset: 36
+ .size: 4
+ .value_kind: by_value
+ - .name: a10
+ .offset: 40
+ .size: 4
+ .value_kind: by_value
+ - .name: a11
+ .offset: 44
+ .size: 4
+ .value_kind: by_value
+ - .name: a12
+ .offset: 48
+ .size: 4
+ .value_kind: by_value
+ - .name: a13
+ .offset: 52
+ .size: 4
+ .value_kind: by_value
+ - .name: a14
+ .offset: 56
+ .size: 4
+ .value_kind: by_value
+ - .name: a15
+ .offset: 60
+ .size: 4
+ .value_kind: by_value
+ - .name: a16
+ .offset: 64
+ .size: 4
+ .value_kind: by_value
+ - .name: a17
+ .offset: 68
+ .size: 4
+ .value_kind: by_value
+ - .name: a18
+ .offset: 72
+ .size: 4
+ .value_kind: by_value
+ - .name: a19
+ .offset: 76
+ .size: 4
+ .value_kind: by_value
+ - .name: a20
+ .offset: 80
+ .size: 4
+ .value_kind: by_value
+ - .name: a21
+ .offset: 84
+ .size: 4
+ .value_kind: by_value
+ - .name: a22
+ .offset: 88
+ .size: 4
+ .value_kind: by_value
+ - .name: a23
+ .offset: 92
+ .size: 4
+ .value_kind: by_value
+ - .name: a24
+ .offset: 96
+ .size: 4
+ .value_kind: by_value
+ - .name: a25
+ .offset: 100
+ .size: 4
+ .value_kind: by_value
+ - .name: a26
+ .offset: 104
+ .size: 4
+ .value_kind: by_value
+ - .name: a27
+ .offset: 108
+ .size: 4
+ .value_kind: by_value
+ - .name: a28
+ .offset: 112
+ .size: 4
+ .value_kind: by_value
+ - .name: a29
+ .offset: 116
+ .size: 4
+ .value_kind: by_value
+ - .name: a30
+ .offset: 120
+ .size: 4
+ .value_kind: by_value
+ - .name: a31
+ .offset: 124
+ .size: 4
+ .value_kind: by_value
+ - .name: a32
+ .offset: 128
+ .size: 4
+ .value_kind: by_value
+ - .name: a33
+ .offset: 132
+ .size: 4
+ .value_kind: by_value
+ - .name: a34
+ .offset: 136
+ .size: 4
+ .value_kind: by_value
+ - .name: a35
+ .offset: 140
+ .size: 4
+ .value_kind: by_value
+ - .offset: 144
+ .size: 4
+ .value_kind: hidden_block_count_x
+ - .offset: 148
+ .size: 4
+ .value_kind: hidden_block_count_y
+ - .offset: 152
+ .size: 4
+ .value_kind: hidden_block_count_z
+ - .offset: 156
+ .size: 2
+ .value_kind: hidden_group_size_x
+ - .offset: 158
+ .size: 2
+ .value_kind: hidden_group_size_y
+ - .offset: 160
+ .size: 2
+ .value_kind: hidden_group_size_z
+ - .offset: 162
+ .size: 2
+ .value_kind: hidden_remainder_x
+ - .offset: 164
+ .size: 2
+ .value_kind: hidden_remainder_y
+ - .offset: 166
+ .size: 2
+ .value_kind: hidden_remainder_z
+ - .offset: 184
+ .size: 8
+ .value_kind: hidden_global_offset_x
+ - .offset: 192
+ .size: 8
+ .value_kind: hidden_global_offset_y
+ - .offset: 200
+ .size: 8
+ .value_kind: hidden_global_offset_z
+ - .offset: 208
+ .size: 2
+ .value_kind: hidden_grid_dims
+ - .offset: 224
+ .size: 8
+ .value_kind: hidden_hostcall_buffer
+ - .offset: 232
+ .size: 8
+ .value_kind: hidden_multigrid_sync_arg
+ - .offset: 240
+ .size: 8
+ .value_kind: hidden_heap_v1
+ - .offset: 248
+ .size: 8
+ .value_kind: hidden_default_queue
+ - .offset: 256
+ .size: 8
+ .value_kind: hidden_completion_action
+ - .offset: 344
+ .size: 8
+ .value_kind: hidden_queue_ptr
+ .group_segment_fixed_size: 0
+ .kernarg_segment_align: 8
+ .kernarg_segment_size: 400
+ .max_flat_workgroup_size: 1024
+ .name: many_inreg_i32
+ .private_segment_fixed_size: 0
+ .sgpr_count: 32
+ .sgpr_spill_count: 0
+ .symbol: many_inreg_i32.kd
+ .uses_dynamic_stack: false
+ .vgpr_count: 0
+ .vgpr_spill_count: 0
+ .wavefront_size: 32
+amdhsa.target: amdgcn-amd-amdhsa--gfx1250
+amdhsa.version:
+ - 1
+ - 2
+...
+
+ .end_amdgpu_metadata
>From bb44a7e1d2c8576d51e60fb5a5f0cfa0c072cdbe Mon Sep 17 00:00:00 2001
From: zwu-2025 <Zeng.Wu2 at amd.com>
Date: Fri, 17 Apr 2026 20:39:08 +0000
Subject: [PATCH 3/7] comments
---
llvm/test/MC/AMDGPU/too_many_user_sgpr.s | 13 +++++++++++++
1 file changed, 13 insertions(+)
create mode 100644 llvm/test/MC/AMDGPU/too_many_user_sgpr.s
diff --git a/llvm/test/MC/AMDGPU/too_many_user_sgpr.s b/llvm/test/MC/AMDGPU/too_many_user_sgpr.s
new file mode 100644
index 0000000000000..55bd8668e012a
--- /dev/null
+++ b/llvm/test/MC/AMDGPU/too_many_user_sgpr.s
@@ -0,0 +1,13 @@
+// RUN: not llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx1250 -filetype=null %s 2>&1 | FileCheck --check-prefix=GFX1250-ERR %s
+
+.text
+.amdgcn_target "amdgcn-amd-amdhsa--gfx1250"
+.amdhsa_code_object_version 6
+
+.amdhsa_kernel many_inreg_i32
+ .amdhsa_next_free_vgpr 0
+ .amdhsa_next_free_sgpr 0
+
+ .amdhsa_user_sgpr_count 65
+// GFX1250-ERR: too many user SGPRs enabled
+.end_amdhsa_kernel
>From 9acff209e598efd31abc0129f7f035b8e36ace01 Mon Sep 17 00:00:00 2001
From: zwu-2025 <Zeng.Wu2 at amd.com>
Date: Fri, 17 Apr 2026 20:39:34 +0000
Subject: [PATCH 4/7] comments
---
.../amdhsa-kernarg-preload-num-sgprs.ll | 16 +
.../llvm.amdgcn.gfx125.user.sgpr.encoding.ll | 15 -
.../llvm.amdgcn.gfx125.user.sgpr.encoding.s | 326 ------------------
3 files changed, 16 insertions(+), 341 deletions(-)
delete mode 100644 llvm/test/CodeGen/AMDGPU/llvm.amdgcn.gfx125.user.sgpr.encoding.ll
delete mode 100644 llvm/test/CodeGen/AMDGPU/llvm.amdgcn.gfx125.user.sgpr.encoding.s
diff --git a/llvm/test/CodeGen/AMDGPU/amdhsa-kernarg-preload-num-sgprs.ll b/llvm/test/CodeGen/AMDGPU/amdhsa-kernarg-preload-num-sgprs.ll
index 4054d4c99c2b8..9c3baf422436e 100644
--- a/llvm/test/CodeGen/AMDGPU/amdhsa-kernarg-preload-num-sgprs.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdhsa-kernarg-preload-num-sgprs.ll
@@ -1,5 +1,6 @@
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx942 -filetype=obj < %s | llvm-objdump -s -j .rodata - | FileCheck --check-prefix=OBJDUMP %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx942 < %s | FileCheck --check-prefix=ASM %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 -filetype=obj < %s | llvm-readobj --notes - | FileCheck --check-prefix=GFX1250-OBJDUMP %s
; OBJDUMP: Contents of section .rodata:
; OBJDUMP-NEXT: 0000 00000000 00000000 10010000 00000000 ................
@@ -70,4 +71,19 @@ define amdgpu_kernel void @amdhsa_kernarg_preload_1_implicit_2(i32 inreg) #0 { r
define amdgpu_kernel void @amdhsa_kernarg_preload_0_implicit_2(i32) #0 { ret void }
+; GFX1250-OBJDUMP: .sgpr_count: 32
+define amdgpu_kernel void @many_inreg_i32(
+ i32 inreg %a0, i32 inreg %a1, i32 inreg %a2, i32 inreg %a3,
+ i32 inreg %a4, i32 inreg %a5, i32 inreg %a6, i32 inreg %a7,
+ i32 inreg %a8, i32 inreg %a9, i32 inreg %a10, i32 inreg %a11,
+ i32 inreg %a12, i32 inreg %a13, i32 inreg %a14, i32 inreg %a15,
+ i32 inreg %a16, i32 inreg %a17, i32 inreg %a18, i32 inreg %a19,
+ i32 inreg %a20, i32 inreg %a21, i32 inreg %a22, i32 inreg %a23,
+ i32 inreg %a24, i32 inreg %a25, i32 inreg %a26, i32 inreg %a27,
+ i32 inreg %a28, i32 inreg %a29, i32 inreg %a30, i32 inreg %a31,
+ i32 inreg %a32, i32 inreg %a33, i32 inreg %a34, i32 inreg %a35) {
+ 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-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.gfx125.user.sgpr.encoding.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.gfx125.user.sgpr.encoding.ll
deleted file mode 100644
index ac0578b883b49..0000000000000
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.gfx125.user.sgpr.encoding.ll
+++ /dev/null
@@ -1,15 +0,0 @@
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 < %s -o - | llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx1250 -filetype=obj -o - | llvm-readobj --notes -
-
-; CHECK: .sgpr_count: 32
-define amdgpu_kernel void @many_inreg_i32(
- i32 inreg %a0, i32 inreg %a1, i32 inreg %a2, i32 inreg %a3,
- i32 inreg %a4, i32 inreg %a5, i32 inreg %a6, i32 inreg %a7,
- i32 inreg %a8, i32 inreg %a9, i32 inreg %a10, i32 inreg %a11,
- i32 inreg %a12, i32 inreg %a13, i32 inreg %a14, i32 inreg %a15,
- i32 inreg %a16, i32 inreg %a17, i32 inreg %a18, i32 inreg %a19,
- i32 inreg %a20, i32 inreg %a21, i32 inreg %a22, i32 inreg %a23,
- i32 inreg %a24, i32 inreg %a25, i32 inreg %a26, i32 inreg %a27,
- i32 inreg %a28, i32 inreg %a29, i32 inreg %a30, i32 inreg %a31,
- i32 inreg %a32, i32 inreg %a33, i32 inreg %a34, i32 inreg %a35) {
- ret void
-}
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.gfx125.user.sgpr.encoding.s b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.gfx125.user.sgpr.encoding.s
deleted file mode 100644
index 81b4584ae785a..0000000000000
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.gfx125.user.sgpr.encoding.s
+++ /dev/null
@@ -1,326 +0,0 @@
- .amdgcn_target "amdgcn-amd-amdhsa--gfx1250"
- .amdhsa_code_object_version 6
- .text
- .globl many_inreg_i32 ; -- Begin function many_inreg_i32
- .p2align 8
- .type many_inreg_i32, at function
-many_inreg_i32: ; @many_inreg_i32
-; %bb.0:
- s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
- s_endpgm
- .section .rodata,"a", at progbits
- .p2align 6, 0x0
- .amdhsa_kernel many_inreg_i32
- .amdhsa_group_segment_fixed_size 0
- .amdhsa_private_segment_fixed_size 0
- .amdhsa_kernarg_size 400
- .amdhsa_user_sgpr_count 32
- .amdhsa_user_sgpr_dispatch_ptr 1
- .amdhsa_user_sgpr_queue_ptr 1
- .amdhsa_user_sgpr_kernarg_segment_ptr 1
- .amdhsa_user_sgpr_dispatch_id 1
- .amdhsa_user_sgpr_kernarg_preload_length 24
- .amdhsa_user_sgpr_kernarg_preload_offset 0
- .amdhsa_user_sgpr_private_segment_size 0
- .amdhsa_wavefront_size32 1
- .amdhsa_uses_dynamic_stack 0
- .amdhsa_enable_private_segment 0
- .amdhsa_system_sgpr_workgroup_id_x 1
- .amdhsa_system_sgpr_workgroup_id_y 1
- .amdhsa_system_sgpr_workgroup_id_z 1
- .amdhsa_system_sgpr_workgroup_info 0
- .amdhsa_system_vgpr_workitem_id 2
- .amdhsa_next_free_vgpr 1
- .amdhsa_next_free_sgpr 32
- .amdhsa_named_barrier_count 0
- .amdhsa_reserve_vcc 0
- .amdhsa_float_round_mode_32 0
- .amdhsa_float_round_mode_16_64 0
- .amdhsa_float_denorm_mode_32 3
- .amdhsa_float_denorm_mode_16_64 3
- .amdhsa_fp16_overflow 0
- .amdhsa_memory_ordered 1
- .amdhsa_forward_progress 1
- .amdhsa_inst_pref_size 1
- .amdhsa_round_robin_scheduling 0
- .amdhsa_exception_fp_ieee_invalid_op 0
- .amdhsa_exception_fp_denorm_src 0
- .amdhsa_exception_fp_ieee_div_zero 0
- .amdhsa_exception_fp_ieee_overflow 0
- .amdhsa_exception_fp_ieee_underflow 0
- .amdhsa_exception_fp_ieee_inexact 0
- .amdhsa_exception_int_div_zero 0
- .end_amdhsa_kernel
- .text
-.Lfunc_end0:
- .size many_inreg_i32, .Lfunc_end0-many_inreg_i32
- ; -- End function
- .set many_inreg_i32.num_vgpr, 0
- .set many_inreg_i32.num_agpr, 0
- .set many_inreg_i32.numbered_sgpr, 0
- .set many_inreg_i32.num_named_barrier, 0
- .set many_inreg_i32.private_seg_size, 0
- .set many_inreg_i32.uses_vcc, 0
- .set many_inreg_i32.uses_flat_scratch, 0
- .set many_inreg_i32.has_dyn_sized_stack, 0
- .set many_inreg_i32.has_recursion, 0
- .set many_inreg_i32.has_indirect_call, 0
- .section .AMDGPU.csdata,"", at progbits
-; Kernel info:
-; codeLenInByte = 12
-; TotalNumSgprs: 32
-; NumVgprs: 0
-; ScratchSize: 0
-; MemoryBound: 0
-; FloatMode: 240
-; IeeeMode: 1
-; LDSByteSize: 0 bytes/workgroup (compile time only)
-; SGPRBlocks: 0
-; VGPRBlocks: 0
-; NumSGPRsForWavesPerEU: 32
-; NumVGPRsForWavesPerEU: 1
-; NamedBarCnt: 0
-; Occupancy: 16
-; WaveLimiterHint : 0
-; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
-; COMPUTE_PGM_RSRC2:USER_SGPR: 32
-; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
-; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
-; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
-; COMPUTE_PGM_RSRC2:TGID_Z_EN: 1
-; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 2
- .text
- .p2alignl 7, 3214868480
- .fill 96, 4, 3214868480
- .section .AMDGPU.gpr_maximums,"", at progbits
- .set amdgpu.max_num_vgpr, 0
- .set amdgpu.max_num_agpr, 0
- .set amdgpu.max_num_sgpr, 0
- .set amdgpu.max_num_named_barrier, 0
- .text
- .section ".note.GNU-stack","", at progbits
- .amdgpu_metadata
----
-amdhsa.kernels:
- - .args:
- - .name: a0
- .offset: 0
- .size: 4
- .value_kind: by_value
- - .name: a1
- .offset: 4
- .size: 4
- .value_kind: by_value
- - .name: a2
- .offset: 8
- .size: 4
- .value_kind: by_value
- - .name: a3
- .offset: 12
- .size: 4
- .value_kind: by_value
- - .name: a4
- .offset: 16
- .size: 4
- .value_kind: by_value
- - .name: a5
- .offset: 20
- .size: 4
- .value_kind: by_value
- - .name: a6
- .offset: 24
- .size: 4
- .value_kind: by_value
- - .name: a7
- .offset: 28
- .size: 4
- .value_kind: by_value
- - .name: a8
- .offset: 32
- .size: 4
- .value_kind: by_value
- - .name: a9
- .offset: 36
- .size: 4
- .value_kind: by_value
- - .name: a10
- .offset: 40
- .size: 4
- .value_kind: by_value
- - .name: a11
- .offset: 44
- .size: 4
- .value_kind: by_value
- - .name: a12
- .offset: 48
- .size: 4
- .value_kind: by_value
- - .name: a13
- .offset: 52
- .size: 4
- .value_kind: by_value
- - .name: a14
- .offset: 56
- .size: 4
- .value_kind: by_value
- - .name: a15
- .offset: 60
- .size: 4
- .value_kind: by_value
- - .name: a16
- .offset: 64
- .size: 4
- .value_kind: by_value
- - .name: a17
- .offset: 68
- .size: 4
- .value_kind: by_value
- - .name: a18
- .offset: 72
- .size: 4
- .value_kind: by_value
- - .name: a19
- .offset: 76
- .size: 4
- .value_kind: by_value
- - .name: a20
- .offset: 80
- .size: 4
- .value_kind: by_value
- - .name: a21
- .offset: 84
- .size: 4
- .value_kind: by_value
- - .name: a22
- .offset: 88
- .size: 4
- .value_kind: by_value
- - .name: a23
- .offset: 92
- .size: 4
- .value_kind: by_value
- - .name: a24
- .offset: 96
- .size: 4
- .value_kind: by_value
- - .name: a25
- .offset: 100
- .size: 4
- .value_kind: by_value
- - .name: a26
- .offset: 104
- .size: 4
- .value_kind: by_value
- - .name: a27
- .offset: 108
- .size: 4
- .value_kind: by_value
- - .name: a28
- .offset: 112
- .size: 4
- .value_kind: by_value
- - .name: a29
- .offset: 116
- .size: 4
- .value_kind: by_value
- - .name: a30
- .offset: 120
- .size: 4
- .value_kind: by_value
- - .name: a31
- .offset: 124
- .size: 4
- .value_kind: by_value
- - .name: a32
- .offset: 128
- .size: 4
- .value_kind: by_value
- - .name: a33
- .offset: 132
- .size: 4
- .value_kind: by_value
- - .name: a34
- .offset: 136
- .size: 4
- .value_kind: by_value
- - .name: a35
- .offset: 140
- .size: 4
- .value_kind: by_value
- - .offset: 144
- .size: 4
- .value_kind: hidden_block_count_x
- - .offset: 148
- .size: 4
- .value_kind: hidden_block_count_y
- - .offset: 152
- .size: 4
- .value_kind: hidden_block_count_z
- - .offset: 156
- .size: 2
- .value_kind: hidden_group_size_x
- - .offset: 158
- .size: 2
- .value_kind: hidden_group_size_y
- - .offset: 160
- .size: 2
- .value_kind: hidden_group_size_z
- - .offset: 162
- .size: 2
- .value_kind: hidden_remainder_x
- - .offset: 164
- .size: 2
- .value_kind: hidden_remainder_y
- - .offset: 166
- .size: 2
- .value_kind: hidden_remainder_z
- - .offset: 184
- .size: 8
- .value_kind: hidden_global_offset_x
- - .offset: 192
- .size: 8
- .value_kind: hidden_global_offset_y
- - .offset: 200
- .size: 8
- .value_kind: hidden_global_offset_z
- - .offset: 208
- .size: 2
- .value_kind: hidden_grid_dims
- - .offset: 224
- .size: 8
- .value_kind: hidden_hostcall_buffer
- - .offset: 232
- .size: 8
- .value_kind: hidden_multigrid_sync_arg
- - .offset: 240
- .size: 8
- .value_kind: hidden_heap_v1
- - .offset: 248
- .size: 8
- .value_kind: hidden_default_queue
- - .offset: 256
- .size: 8
- .value_kind: hidden_completion_action
- - .offset: 344
- .size: 8
- .value_kind: hidden_queue_ptr
- .group_segment_fixed_size: 0
- .kernarg_segment_align: 8
- .kernarg_segment_size: 400
- .max_flat_workgroup_size: 1024
- .name: many_inreg_i32
- .private_segment_fixed_size: 0
- .sgpr_count: 32
- .sgpr_spill_count: 0
- .symbol: many_inreg_i32.kd
- .uses_dynamic_stack: false
- .vgpr_count: 0
- .vgpr_spill_count: 0
- .wavefront_size: 32
-amdhsa.target: amdgcn-amd-amdhsa--gfx1250
-amdhsa.version:
- - 1
- - 2
-...
-
- .end_amdgpu_metadata
>From d99c719bb307501e23a18b39f77e701da78c7e2d Mon Sep 17 00:00:00 2001
From: zwu-2025 <Zeng.Wu2 at amd.com>
Date: Fri, 17 Apr 2026 22:38:45 +0000
Subject: [PATCH 5/7] mc test
---
llvm/test/MC/AMDGPU/too_many_user_sgpr.s | 2 --
1 file changed, 2 deletions(-)
diff --git a/llvm/test/MC/AMDGPU/too_many_user_sgpr.s b/llvm/test/MC/AMDGPU/too_many_user_sgpr.s
index 55bd8668e012a..87f2625a0bcab 100644
--- a/llvm/test/MC/AMDGPU/too_many_user_sgpr.s
+++ b/llvm/test/MC/AMDGPU/too_many_user_sgpr.s
@@ -1,7 +1,5 @@
// RUN: not llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx1250 -filetype=null %s 2>&1 | FileCheck --check-prefix=GFX1250-ERR %s
-.text
-.amdgcn_target "amdgcn-amd-amdhsa--gfx1250"
.amdhsa_code_object_version 6
.amdhsa_kernel many_inreg_i32
>From ca64a267cd50675496d7feee0faf953ae559115c Mon Sep 17 00:00:00 2001
From: zwu-2025 <Zeng.Wu2 at amd.com>
Date: Wed, 22 Apr 2026 05:19:45 +0000
Subject: [PATCH 6/7] mc tests
---
...ny_user_sgpr.s => user-sgpr-count-gfx1250-diag.s} | 0
llvm/test/MC/AMDGPU/user-sgpr-count-gfx1250.s | 12 ++++++++++++
2 files changed, 12 insertions(+)
rename llvm/test/MC/AMDGPU/{too_many_user_sgpr.s => user-sgpr-count-gfx1250-diag.s} (100%)
create mode 100644 llvm/test/MC/AMDGPU/user-sgpr-count-gfx1250.s
diff --git a/llvm/test/MC/AMDGPU/too_many_user_sgpr.s b/llvm/test/MC/AMDGPU/user-sgpr-count-gfx1250-diag.s
similarity index 100%
rename from llvm/test/MC/AMDGPU/too_many_user_sgpr.s
rename to llvm/test/MC/AMDGPU/user-sgpr-count-gfx1250-diag.s
diff --git a/llvm/test/MC/AMDGPU/user-sgpr-count-gfx1250.s b/llvm/test/MC/AMDGPU/user-sgpr-count-gfx1250.s
new file mode 100644
index 0000000000000..2c0ca15340643
--- /dev/null
+++ b/llvm/test/MC/AMDGPU/user-sgpr-count-gfx1250.s
@@ -0,0 +1,12 @@
+// RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx1250 -filetype=asm %s 2>&1 | FileCheck %s
+
+.amdhsa_code_object_version 6
+
+ // CHECK:.amdhsa_user_sgpr_count 33
+.amdhsa_kernel many_inreg_i32
+ .amdhsa_next_free_vgpr 0
+ .amdhsa_next_free_sgpr 0
+
+ .amdhsa_user_sgpr_count 33
+
+.end_amdhsa_kernel
>From b1a3152e4a1f75500fae19374294f2b6cccc4f09 Mon Sep 17 00:00:00 2001
From: zwu-2025 <Zeng.Wu2 at amd.com>
Date: Thu, 23 Apr 2026 05:48:55 +0000
Subject: [PATCH 7/7] add more valid tests
---
.../MC/AMDGPU/user-sgpr-count-gfx1250-diag.s | 2 +-
llvm/test/MC/AMDGPU/user-sgpr-count-gfx1250.s | 36 +++++++++++++++++--
2 files changed, 35 insertions(+), 3 deletions(-)
diff --git a/llvm/test/MC/AMDGPU/user-sgpr-count-gfx1250-diag.s b/llvm/test/MC/AMDGPU/user-sgpr-count-gfx1250-diag.s
index 87f2625a0bcab..67b1c037d0f9c 100644
--- a/llvm/test/MC/AMDGPU/user-sgpr-count-gfx1250-diag.s
+++ b/llvm/test/MC/AMDGPU/user-sgpr-count-gfx1250-diag.s
@@ -6,6 +6,6 @@
.amdhsa_next_free_vgpr 0
.amdhsa_next_free_sgpr 0
- .amdhsa_user_sgpr_count 65
+ .amdhsa_user_sgpr_count 64
// GFX1250-ERR: too many user SGPRs enabled
.end_amdhsa_kernel
diff --git a/llvm/test/MC/AMDGPU/user-sgpr-count-gfx1250.s b/llvm/test/MC/AMDGPU/user-sgpr-count-gfx1250.s
index 2c0ca15340643..fbe29a95d8b79 100644
--- a/llvm/test/MC/AMDGPU/user-sgpr-count-gfx1250.s
+++ b/llvm/test/MC/AMDGPU/user-sgpr-count-gfx1250.s
@@ -2,11 +2,43 @@
.amdhsa_code_object_version 6
- // CHECK:.amdhsa_user_sgpr_count 33
-.amdhsa_kernel many_inreg_i32
+// CHECK:.amdhsa_user_sgpr_count 8
+ .amdhsa_kernel user_sgpr_0
.amdhsa_next_free_vgpr 0
.amdhsa_next_free_sgpr 0
+ .amdhsa_user_sgpr_count 8
+.end_amdhsa_kernel
+
+// CHECK:.amdhsa_user_sgpr_count 31
+.amdhsa_kernel user_sgpr_1
+ .amdhsa_next_free_vgpr 1
+ .amdhsa_next_free_sgpr 0
+
+ .amdhsa_user_sgpr_count 31
+.end_amdhsa_kernel
+
+// CHECK:.amdhsa_user_sgpr_count 32
+.amdhsa_kernel user_sgpr_2
+ .amdhsa_next_free_vgpr 1
+ .amdhsa_next_free_sgpr 0
+
+ .amdhsa_user_sgpr_count 32
+ .end_amdhsa_kernel
+
+// CHECK:.amdhsa_user_sgpr_count 33
+.amdhsa_kernel user_sgpr_3
+ .amdhsa_next_free_vgpr 2
+ .amdhsa_next_free_sgpr 0
+
.amdhsa_user_sgpr_count 33
+.end_amdhsa_kernel
+
+// CHECK:.amdhsa_user_sgpr_count 63
+.amdhsa_kernel user_sgpr_4
+ .amdhsa_next_free_vgpr 0
+ .amdhsa_next_free_sgpr 0
+
+ .amdhsa_user_sgpr_count 63
.end_amdhsa_kernel
More information about the llvm-commits
mailing list