https://github.com/jli-melchior updated https://github.com/llvm/llvm-project/pull/233773
>From 41303030ff50c6962839df03277ae784cb3adf8b Mon Sep 17 00:00:00 2001 From: Jun Li <[email protected]> Date: Tue, 20 Oct 2026 09:47:44 +0000 Subject: [PATCH 1/4] [AMDGPU] Add sched_vgpr/sgpr/agpr intrinsics, builtins, and MLIR ops Add exact GPR register scheduling hints for VGPR, SGPR, and AGPR: - Add llvm.amdgcn.sched.{vgpr,sgpr,agpr} intrinsics with (value, regnum, strict) signature - Add __builtin_amdgcn_sched_{vgpr,sgpr,agpr} clang builtins with Sema checks - Add ISel lowering to V_SCHEDULE_VGPR_B*, S_SCHEDULE_SGPR_B*, A_SCHEDULE_AGPR_B* pseudos (32-512 bit) - Add VGPRRegHint/SGPRRegHint/AGPRRegHint register allocation hint types in SIRegisterInfo - Add MLIR ROCDL rocdl.sched.{vgpr,sgpr,agpr} ops with strict/soft modes - Add CodeGen and MIR tests Depends on #227953 (sched_bank) which provides the AMDGPUSchedBank pass infrastructure that lowers these pseudos into register allocation hints. Co-Authored-By: Claude Opus 4 <[email protected]> --- clang/include/clang/Basic/BuiltinsAMDGPU.td | 5 + clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp | 26 +++++ clang/lib/Sema/SemaAMDGPU.cpp | 28 +++++ .../builtins-amdgcn-sched-gpr.cl | 68 ++++++++++++ llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 22 ++++ llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp | 57 ++++++++++ llvm/lib/Target/AMDGPU/SIInstructions.td | 48 +++++++++ llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp | 21 ++++ llvm/lib/Target/AMDGPU/SIRegisterInfo.h | 11 +- llvm/test/CodeGen/AMDGPU/sched-gpr-hints.mir | 66 ++++++++++++ llvm/test/CodeGen/AMDGPU/sched-gpr.ll | 101 ++++++++++++++++++ mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td | 73 +++++++++++++ mlir/test/Dialect/LLVMIR/rocdl.mlir | 24 +++++ mlir/test/Target/LLVMIR/rocdl.mlir | 29 +++++ 14 files changed, 578 insertions(+), 1 deletion(-) create mode 100644 clang/test/CodeGenOpenCL/builtins-amdgcn-sched-gpr.cl create mode 100644 llvm/test/CodeGen/AMDGPU/sched-gpr-hints.mir create mode 100644 llvm/test/CodeGen/AMDGPU/sched-gpr.ll diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td index 34a1bd2d301a37..466dfc8e99e758 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.td +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td @@ -1012,6 +1012,11 @@ def __builtin_amdgcn_cvt_sr_f16_f32 : AMDGPUBuiltin<"_ExtVector<2, _Float16>(_Ex //===----------------------------------------------------------------------===// def __builtin_amdgcn_s_cluster_barrier : AMDGPUBuiltin<"void()", [], "gfx1250-insts">; +// Exact GPR register scheduling hints. +def __builtin_amdgcn_sched_vgpr : AMDGPUBuiltin<"int(int, _Constant int, _Constant bool)", [Const, CustomTypeChecking]>; +def __builtin_amdgcn_sched_sgpr : AMDGPUBuiltin<"int(int, _Constant int, _Constant bool)", [Const, CustomTypeChecking]>; +def __builtin_amdgcn_sched_agpr : AMDGPUBuiltin<"int(int, _Constant int, _Constant bool)", [Const, CustomTypeChecking], "mai-insts">; + def __builtin_amdgcn_flat_prefetch : AMDGPUBuiltin<"void(void const address_space<0> *, _Constant int)", [Const], "vmem-pref-insts">; def __builtin_amdgcn_global_prefetch : AMDGPUBuiltin<"void(void const address_space<1> *, _Constant int)", [Const], "vmem-pref-insts">; diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp index 2f29e9557ba56f..c27a8e8f2a9fcf 100644 --- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp +++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp @@ -2283,6 +2283,32 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID, case AMDGPU::BI__builtin_amdgcn_permlane_xor: return emitBuiltinWithOneOverloadedType<3>(*this, E, Intrinsic::amdgcn_permlane_xor); + case AMDGPU::BI__builtin_amdgcn_sched_vgpr: + case AMDGPU::BI__builtin_amdgcn_sched_sgpr: + case AMDGPU::BI__builtin_amdgcn_sched_agpr: { + llvm::Value *Val = EmitScalarExpr(E->getArg(0)); + llvm::Value *RegNum = EmitScalarExpr(E->getArg(1)); + llvm::Value *Strict; + if (E->getNumArgs() > 2) { + Strict = EmitScalarExpr(E->getArg(2)); + Strict = Builder.CreateZExt(Strict, RegNum->getType()); + } else { + // Default: strict + Strict = Builder.getInt32(1); + } + Intrinsic::ID IID; + switch (BuiltinID) { + case AMDGPU::BI__builtin_amdgcn_sched_vgpr: + IID = Intrinsic::amdgcn_sched_vgpr; break; + case AMDGPU::BI__builtin_amdgcn_sched_sgpr: + IID = Intrinsic::amdgcn_sched_sgpr; break; + case AMDGPU::BI__builtin_amdgcn_sched_agpr: + IID = Intrinsic::amdgcn_sched_agpr; break; + default: llvm_unreachable("unexpected builtin"); + } + llvm::Function *F = CGM.getIntrinsic(IID, Val->getType()); + return Builder.CreateCall(F, {Val, RegNum, Strict}); + } default: return nullptr; } diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp index b1f6fa33649bd1..556526394d2872 100644 --- a/clang/lib/Sema/SemaAMDGPU.cpp +++ b/clang/lib/Sema/SemaAMDGPU.cpp @@ -465,6 +465,34 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(const TargetInfo &TI, /*High=*/0) || SemaRef.BuiltinConstantArgRange(TheCall, /*ArgNum=*/2, /*Low=*/0, /*High=*/0); + case AMDGPU::BI__builtin_amdgcn_sched_vgpr: + case AMDGPU::BI__builtin_amdgcn_sched_sgpr: + case AMDGPU::BI__builtin_amdgcn_sched_agpr: { + if (SemaRef.checkArgCountRange(TheCall, 2, 3)) + return true; + Expr *DataArg = TheCall->getArg(0); + QualType DataTy = DataArg->getType(); + if (DataTy->isAnyComplexType() || + !(DataTy->isArithmeticType() || + (DataTy->isVectorType() && + DataTy->castAs<VectorType>() + ->getElementType() + ->isArithmeticType()))) { + SemaRef.Diag(DataArg->getBeginLoc(), + diag::err_typecheck_cond_expect_int_float) + << DataTy << DataArg->getSourceRange(); + return true; + } + TheCall->setType(DataTy); + int MaxReg; + switch (BuiltinID) { + case AMDGPU::BI__builtin_amdgcn_sched_vgpr: MaxReg = 1023; break; + case AMDGPU::BI__builtin_amdgcn_sched_sgpr: MaxReg = 105; break; + case AMDGPU::BI__builtin_amdgcn_sched_agpr: MaxReg = 255; break; + default: llvm_unreachable("unexpected builtin"); + } + return SemaRef.BuiltinConstantArgRange(TheCall, 1, 0, MaxReg); + } default: return false; } diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-sched-gpr.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-sched-gpr.cl new file mode 100644 index 00000000000000..2d889f27f38a4e --- /dev/null +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-sched-gpr.cl @@ -0,0 +1,68 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py +// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu gfx1250 -emit-llvm -o - %s | FileCheck %s +// REQUIRES: amdgpu-registered-target + +typedef int __attribute__((ext_vector_type(2))) int2; +typedef int __attribute__((ext_vector_type(4))) int4; +typedef int __attribute__((ext_vector_type(8))) int8; +typedef float __attribute__((ext_vector_type(2))) float2; +typedef float __attribute__((ext_vector_type(4))) float4; + +// === sched_vgpr tests === + +// CHECK-LABEL: @test_sched_vgpr_i32 +// CHECK: call i32 @llvm.amdgcn.sched.vgpr.i32(i32 %{{.*}}, i32 42, i32 0) +void test_sched_vgpr_i32(global int *out, int x) { + *out = __builtin_amdgcn_sched_vgpr(x, 42, false); +} + +// CHECK-LABEL: @test_sched_vgpr_f32_strict +// CHECK: call float @llvm.amdgcn.sched.vgpr.f32(float %{{.*}}, i32 100, i32 1) +void test_sched_vgpr_f32_strict(global float *out, float x) { + *out = __builtin_amdgcn_sched_vgpr(x, 100, true); +} + +// CHECK-LABEL: @test_sched_vgpr_i64 +// CHECK: call i64 @llvm.amdgcn.sched.vgpr.i64(i64 %{{.*}}, i32 200, i32 0) +void test_sched_vgpr_i64(global long *out, long x) { + *out = __builtin_amdgcn_sched_vgpr(x, 200, false); +} + +// CHECK-LABEL: @test_sched_vgpr_v2i32 +// CHECK: call <2 x i32> @llvm.amdgcn.sched.vgpr.v2i32(<2 x i32> %{{.*}}, i32 50, i32 1) +void test_sched_vgpr_v2i32(global int2 *out, int2 x) { + *out = __builtin_amdgcn_sched_vgpr(x, 50, true); +} + +// CHECK-LABEL: @test_sched_vgpr_v4i32 +// CHECK: call <4 x i32> @llvm.amdgcn.sched.vgpr.v4i32(<4 x i32> %{{.*}}, i32 512, i32 1) +void test_sched_vgpr_v4i32(global int4 *out, int4 x) { + *out = __builtin_amdgcn_sched_vgpr(x, 512, true); +} + +// CHECK-LABEL: @test_sched_vgpr_default_strict +// CHECK: call i32 @llvm.amdgcn.sched.vgpr.i32(i32 %{{.*}}, i32 42, i32 1) +void test_sched_vgpr_default_strict(global int *out, int x) { + // 2 args: strict=true by default + *out = __builtin_amdgcn_sched_vgpr(x, 42); +} + +// === sched_sgpr tests === + +// CHECK-LABEL: @test_sched_sgpr_i32 +// CHECK: call i32 @llvm.amdgcn.sched.sgpr.i32(i32 %{{.*}}, i32 10, i32 0) +void test_sched_sgpr_i32(global int *out, int x) { + *out = __builtin_amdgcn_sched_sgpr(x, 10, false); +} + +// CHECK-LABEL: @test_sched_sgpr_i32_strict +// CHECK: call i32 @llvm.amdgcn.sched.sgpr.i32(i32 %{{.*}}, i32 50, i32 1) +void test_sched_sgpr_i32_strict(global int *out, int x) { + *out = __builtin_amdgcn_sched_sgpr(x, 50, true); +} + +// CHECK-LABEL: @test_sched_sgpr_default_strict +// CHECK: call i32 @llvm.amdgcn.sched.sgpr.i32(i32 %{{.*}}, i32 20, i32 1) +void test_sched_sgpr_default_strict(global int *out, int x) { + *out = __builtin_amdgcn_sched_sgpr(x, 20); +} diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td index ed80bc4628b6fa..2400bc1d6b5e4b 100644 --- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td +++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td @@ -4478,4 +4478,26 @@ def int_amdgcn_global_load_monitor_b128 : AMDGPULoadMonitor<global_ptr_ty>; /// incoming block. def int_amdgcn_dead: DefaultAttrsIntrinsic<[llvm_any_ty], [], [IntrNoMem]>; + +// Hint the register allocator to place the result in a specific VGPR/SGPR/AGPR +// register number. These are identity functions at the IR level; the register +// number and strict flag arguments are consumed by the pre-RA AMDGPUSchedBank +// pass which attaches a register allocation hint. +// regnum: the target HW register index (e.g., 42 means v42/s42/a42). +// strict: 1 = allocator must obey if possible, 0 = advisory hint. +// For wide values (64/128/256 bit), regnum is the starting register index. +def int_amdgcn_sched_vgpr : DefaultAttrsIntrinsic< + [llvm_any_ty], [LLVMMatchType<0>, llvm_i32_ty, llvm_i32_ty], + [IntrNoMem, IntrSpeculatable, IntrWillReturn, + ImmArg<ArgIndex<1>>, ImmArg<ArgIndex<2>>]>; + +def int_amdgcn_sched_sgpr : DefaultAttrsIntrinsic< + [llvm_any_ty], [LLVMMatchType<0>, llvm_i32_ty, llvm_i32_ty], + [IntrNoMem, IntrSpeculatable, IntrWillReturn, + ImmArg<ArgIndex<1>>, ImmArg<ArgIndex<2>>]>; + +def int_amdgcn_sched_agpr : DefaultAttrsIntrinsic< + [llvm_any_ty], [LLVMMatchType<0>, llvm_i32_ty, llvm_i32_ty], + [IntrNoMem, IntrSpeculatable, IntrWillReturn, + ImmArg<ArgIndex<1>>, ImmArg<ArgIndex<2>>]>; } diff --git a/llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp b/llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp index dc2abb60567ffa..a6c20d61be838f 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp @@ -3353,6 +3353,63 @@ void AMDGPUDAGToDAGISel::SelectINTRINSIC_WO_CHAIN(SDNode *N) { case Intrinsic::amdgcn_interp_p1_f16: SelectInterpP1F16(N); return; + case Intrinsic::amdgcn_sched_vgpr: + case Intrinsic::amdgcn_sched_sgpr: + case Intrinsic::amdgcn_sched_agpr: { + SDValue Src = N->getOperand(1); + auto *RegNumC = dyn_cast<ConstantSDNode>(N->getOperand(2)); + auto *StrictC = dyn_cast<ConstantSDNode>(N->getOperand(3)); + if (!RegNumC || !StrictC) { + SelectCode(N); + return; + } + unsigned Bits = N->getValueType(0).getSizeInBits(); + unsigned PseudoOpc; + if (IntrID == Intrinsic::amdgcn_sched_vgpr) { + switch (Bits) { + case 32: PseudoOpc = AMDGPU::V_SCHEDULE_VGPR_B32; break; + case 64: PseudoOpc = AMDGPU::V_SCHEDULE_VGPR_B64; break; + case 128: PseudoOpc = AMDGPU::V_SCHEDULE_VGPR_B128; break; + case 256: PseudoOpc = AMDGPU::V_SCHEDULE_VGPR_B256; break; + case 512: PseudoOpc = AMDGPU::V_SCHEDULE_VGPR_B512; break; + default: + ReplaceUses(SDValue(N, 0), Src); + CurDAG->RemoveDeadNode(N); + return; + } + } else if (IntrID == Intrinsic::amdgcn_sched_sgpr) { + switch (Bits) { + case 32: PseudoOpc = AMDGPU::S_SCHEDULE_SGPR_B32; break; + case 64: PseudoOpc = AMDGPU::S_SCHEDULE_SGPR_B64; break; + case 128: PseudoOpc = AMDGPU::S_SCHEDULE_SGPR_B128; break; + case 256: PseudoOpc = AMDGPU::S_SCHEDULE_SGPR_B256; break; + case 512: PseudoOpc = AMDGPU::S_SCHEDULE_SGPR_B512; break; + default: + ReplaceUses(SDValue(N, 0), Src); + CurDAG->RemoveDeadNode(N); + return; + } + } else { + switch (Bits) { + case 32: PseudoOpc = AMDGPU::A_SCHEDULE_AGPR_B32; break; + case 64: PseudoOpc = AMDGPU::A_SCHEDULE_AGPR_B64; break; + case 128: PseudoOpc = AMDGPU::A_SCHEDULE_AGPR_B128; break; + case 256: PseudoOpc = AMDGPU::A_SCHEDULE_AGPR_B256; break; + case 512: PseudoOpc = AMDGPU::A_SCHEDULE_AGPR_B512; break; + default: + ReplaceUses(SDValue(N, 0), Src); + CurDAG->RemoveDeadNode(N); + return; + } + } + SDValue RegNumImm = CurDAG->getTargetConstant(RegNumC->getZExtValue(), + SDLoc(N), MVT::i32); + SDValue StrictImm = CurDAG->getTargetConstant(StrictC->getZExtValue(), + SDLoc(N), MVT::i32); + CurDAG->SelectNodeTo(N, PseudoOpc, N->getVTList(), + {Src, RegNumImm, StrictImm}); + return; + } case Intrinsic::amdgcn_permlane16_swap: case Intrinsic::amdgcn_permlane32_swap: { if ((IntrID == Intrinsic::amdgcn_permlane16_swap && diff --git a/llvm/lib/Target/AMDGPU/SIInstructions.td b/llvm/lib/Target/AMDGPU/SIInstructions.td index c9e0e4c79f8ccd..af21dd0a6483e6 100644 --- a/llvm/lib/Target/AMDGPU/SIInstructions.td +++ b/llvm/lib/Target/AMDGPU/SIInstructions.td @@ -129,6 +129,23 @@ def V_CNDMASK_B64_PSEUDO : VOP3Common <(outs VReg_64:$vdst), let usesCustomInserter = 1; } +// Exact VGPR register hint pseudos. Identity copies carrying a target VGPR +// register number and strict flag. The pre-RA AMDGPUSchedBank pass reads +// these operands, sets a register allocation hint, and replaces with COPY. +let isAsCheapAsAMove = 1, isReMaterializable = 1, hasSideEffects = 0, + mayLoad = 0, mayStore = 0, Size = 0 in { + def V_SCHEDULE_VGPR_B32 : VPseudoInstSI<(outs VGPR_32:$vdst), + (ins VGPR_32:$src, i32imm:$regnum, i32imm:$strict)>; + def V_SCHEDULE_VGPR_B64 : VPseudoInstSI<(outs VReg_64:$vdst), + (ins VReg_64:$src, i32imm:$regnum, i32imm:$strict)>; + def V_SCHEDULE_VGPR_B128 : VPseudoInstSI<(outs VReg_128:$vdst), + (ins VReg_128:$src, i32imm:$regnum, i32imm:$strict)>; + def V_SCHEDULE_VGPR_B256 : VPseudoInstSI<(outs VReg_256:$vdst), + (ins VReg_256:$src, i32imm:$regnum, i32imm:$strict)>; + def V_SCHEDULE_VGPR_B512 : VPseudoInstSI<(outs VReg_512:$vdst), + (ins VReg_512:$src, i32imm:$regnum, i32imm:$strict)>; +} + // 64-bit vector move instruction. This is mainly used by the // SIFoldOperands pass to enable folding of inline immediates. def V_MOV_B64_PSEUDO : VPseudoInstSI <(outs VReg_64:$vdst), @@ -217,6 +234,37 @@ def STRICT_WQM : PseudoInstSI <(outs unknown:$vdst), (ins unknown:$src0)>; } // End let hasSideEffects = 0, mayLoad = 0, mayStore = 0, Uses = [EXEC] +// Exact SGPR register hint pseudos. Placed outside the Uses = [EXEC] block +// because SGPR operations do not use EXEC. +let isAsCheapAsAMove = 1, isReMaterializable = 1, hasSideEffects = 0, + mayLoad = 0, mayStore = 0, Size = 0 in { + def S_SCHEDULE_SGPR_B32 : SPseudoInstSI<(outs SGPR_32:$sdst), + (ins SGPR_32:$src, i32imm:$regnum, i32imm:$strict)>; + def S_SCHEDULE_SGPR_B64 : SPseudoInstSI<(outs SGPR_64:$sdst), + (ins SGPR_64:$src, i32imm:$regnum, i32imm:$strict)>; + def S_SCHEDULE_SGPR_B128 : SPseudoInstSI<(outs SGPR_128:$sdst), + (ins SGPR_128:$src, i32imm:$regnum, i32imm:$strict)>; + def S_SCHEDULE_SGPR_B256 : SPseudoInstSI<(outs SGPR_256:$sdst), + (ins SGPR_256:$src, i32imm:$regnum, i32imm:$strict)>; + def S_SCHEDULE_SGPR_B512 : SPseudoInstSI<(outs SGPR_512:$sdst), + (ins SGPR_512:$src, i32imm:$regnum, i32imm:$strict)>; +} + +// Exact AGPR register hint pseudos. Only valid on targets with MAI instructions. +let isAsCheapAsAMove = 1, isReMaterializable = 1, hasSideEffects = 0, + mayLoad = 0, mayStore = 0, Size = 0 in { + def A_SCHEDULE_AGPR_B32 : VPseudoInstSI<(outs AGPR_32:$vdst), + (ins AGPR_32:$src, i32imm:$regnum, i32imm:$strict)>; + def A_SCHEDULE_AGPR_B64 : VPseudoInstSI<(outs AReg_64:$vdst), + (ins AReg_64:$src, i32imm:$regnum, i32imm:$strict)>; + def A_SCHEDULE_AGPR_B128 : VPseudoInstSI<(outs AReg_128:$vdst), + (ins AReg_128:$src, i32imm:$regnum, i32imm:$strict)>; + def A_SCHEDULE_AGPR_B256 : VPseudoInstSI<(outs AReg_256:$vdst), + (ins AReg_256:$src, i32imm:$regnum, i32imm:$strict)>; + def A_SCHEDULE_AGPR_B512 : VPseudoInstSI<(outs AReg_512:$vdst), + (ins AReg_512:$src, i32imm:$regnum, i32imm:$strict)>; +} + def WWM_COPY : SPseudoInstSI < (outs unknown:$dst), (ins unknown:$src)> { let hasSideEffects = 0; diff --git a/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp b/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp index 0989263433f41c..0bc47414f52232 100644 --- a/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp +++ b/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp @@ -55,6 +55,7 @@ static cl::opt<unsigned> StressSGPRLimit( "amdgpu-stress-sgpr", cl::Hidden, cl::init(0), cl::desc("Limit SGPRs to N registers by reserving the rest")); + std::array<std::vector<int16_t>, 32> SIRegisterInfo::RegSplitParts; std::array<std::array<uint16_t, 32>, 9> SIRegisterInfo::SubRegFromChannelTable; @@ -4329,6 +4330,26 @@ bool SIRegisterInfo::getRegAllocationHints(Register VirtReg, } return false; } + case AMDGPURI::VGPRRegHint: + case AMDGPURI::StrictVGPRRegHint: + case AMDGPURI::SGPRRegHint: + case AMDGPURI::StrictSGPRRegHint: + case AMDGPURI::AGPRRegHint: + case AMDGPURI::StrictAGPRRegHint: { + // Exact register preference from llvm.amdgcn.sched.{vgpr,sgpr,agpr}. + // The target HW register index is carried in Hint.second. Prefer the + // physreg whose HW index matches. + unsigned TargetIdx = Hint.second; + for (MCPhysReg PhysReg : Order) { + if (MRI.isReserved(PhysReg)) + continue; + if (getHWRegIndex(PhysReg) == TargetIdx) + Hints.insert(PhysReg); + } + return Hint.first == AMDGPURI::StrictVGPRRegHint || + Hint.first == AMDGPURI::StrictSGPRRegHint || + Hint.first == AMDGPURI::StrictAGPRRegHint; + } default: return TargetRegisterInfo::getRegAllocationHints(VirtReg, Order, Hints, MF, VRM); diff --git a/llvm/lib/Target/AMDGPU/SIRegisterInfo.h b/llvm/lib/Target/AMDGPU/SIRegisterInfo.h index b70adfdd921012..a1a49f23ad52d4 100644 --- a/llvm/lib/Target/AMDGPU/SIRegisterInfo.h +++ b/llvm/lib/Target/AMDGPU/SIRegisterInfo.h @@ -38,7 +38,16 @@ struct SGPRSpillBuilder; /// Register allocation hint types. Helps eliminate unneeded COPY with True16 namespace AMDGPURI { -enum { Size16 = 1, Size32 = 2 }; +enum { + Size16 = 1, + Size32 = 2, + VGPRRegHint = 3, + StrictVGPRRegHint = 4, + SGPRRegHint = 5, + StrictSGPRRegHint = 6, + AGPRRegHint = 7, + StrictAGPRRegHint = 8, +}; } // end namespace AMDGPURI diff --git a/llvm/test/CodeGen/AMDGPU/sched-gpr-hints.mir b/llvm/test/CodeGen/AMDGPU/sched-gpr-hints.mir new file mode 100644 index 00000000000000..d480eea0ce596b --- /dev/null +++ b/llvm/test/CodeGen/AMDGPU/sched-gpr-hints.mir @@ -0,0 +1,66 @@ +# RUN: llc -mtriple=amdgcn -mcpu=gfx1250 -run-pass=amdgpu-sched-bank -o - %s | FileCheck %s + +# The AMDGPUSchedBank pass attaches exact register allocation hints (not +# printed in MIR) and lowers each schedule pseudo to a plain COPY. + +--- +# CHECK-LABEL: name: lower_sched_vgpr_b32 +# CHECK: %1:vgpr_32 = COPY %0 +# CHECK-NOT: V_SCHEDULE_VGPR_B32 +name: lower_sched_vgpr_b32 +tracksRegLiveness: true +body: | + bb.0: + liveins: $vgpr0 + %0:vgpr_32 = COPY $vgpr0 + %1:vgpr_32 = V_SCHEDULE_VGPR_B32 %0, 42, 1, implicit $exec + %2:vgpr_32 = V_ADD_U32_e32 1, %1, implicit $exec + $vgpr0 = COPY %2 + SI_RETURN_TO_EPILOG $vgpr0 +... + +--- +# CHECK-LABEL: name: lower_sched_vgpr_b64 +# CHECK: %1:vreg_64_align2 = COPY %0 +# CHECK-NOT: V_SCHEDULE_VGPR_B64 +name: lower_sched_vgpr_b64 +tracksRegLiveness: true +body: | + bb.0: + liveins: $vgpr0_vgpr1 + %0:vreg_64_align2 = COPY $vgpr0_vgpr1 + %1:vreg_64_align2 = V_SCHEDULE_VGPR_B64 %0, 100, 0, implicit $exec + $vgpr0_vgpr1 = COPY %1 + SI_RETURN_TO_EPILOG $vgpr0_vgpr1 +... + +--- +# CHECK-LABEL: name: lower_sched_sgpr_b32 +# CHECK: %1:sgpr_32 = COPY %0 +# CHECK-NOT: S_SCHEDULE_SGPR_B32 +name: lower_sched_sgpr_b32 +tracksRegLiveness: true +body: | + bb.0: + liveins: $sgpr0 + %0:sgpr_32 = COPY $sgpr0 + %1:sgpr_32 = S_SCHEDULE_SGPR_B32 %0, 10, 1 + %2:sgpr_32 = S_ADD_I32 %1, 1, implicit-def $scc + $sgpr0 = COPY %2 + SI_RETURN_TO_EPILOG $sgpr0 +... + +--- +# CHECK-LABEL: name: lower_sched_sgpr_b64 +# CHECK: %1:sgpr_64 = COPY %0 +# CHECK-NOT: S_SCHEDULE_SGPR_B64 +name: lower_sched_sgpr_b64 +tracksRegLiveness: true +body: | + bb.0: + liveins: $sgpr0_sgpr1 + %0:sgpr_64 = COPY $sgpr0_sgpr1 + %1:sgpr_64 = S_SCHEDULE_SGPR_B64 %0, 20, 0 + $sgpr0_sgpr1 = COPY %1 + SI_RETURN_TO_EPILOG $sgpr0_sgpr1 +... diff --git a/llvm/test/CodeGen/AMDGPU/sched-gpr.ll b/llvm/test/CodeGen/AMDGPU/sched-gpr.ll new file mode 100644 index 00000000000000..846b642d0b0cde --- /dev/null +++ b/llvm/test/CodeGen/AMDGPU/sched-gpr.ll @@ -0,0 +1,101 @@ +; RUN: llc -mtriple=amdgcn -mcpu=gfx1250 < %s | FileCheck -check-prefix=GCN %s + +; Verify that llvm.amdgcn.schedule.{vgpr,sgpr} steer register allocation to +; the exact requested register number. Kernels are pinned to occupancy 1 +; (waves-per-eu=1,1 with a single-wave workgroup) so the whole 1024-VGPR file +; is addressable and high-numbered registers are available to the allocator. + +declare i32 @llvm.amdgcn.sched.vgpr.i32(i32, i32, i32) +declare i64 @llvm.amdgcn.sched.vgpr.i64(i64, i32, i32) +declare <16 x float> @llvm.amdgcn.sched.vgpr.v16f32(<16 x float>, i32, i32) + +; -------------------------------------------------------------------------- +; sched_vgpr: strict hint to v42 + +; GCN-LABEL: {{^}}vgpr42: +; GCN: v42 +define amdgpu_kernel void @vgpr42(ptr addrspace(1) %out, i32 %x) #0 { + %h = call i32 @llvm.amdgcn.sched.vgpr.i32(i32 %x, i32 42, i32 1) + %y = add i32 %h, 1 + store i32 %y, ptr addrspace(1) %out + ret void +} + +; -------------------------------------------------------------------------- +; sched_vgpr: strict hint to v100 + +; GCN-LABEL: {{^}}vgpr100: +; GCN: v100 +define amdgpu_kernel void @vgpr100(ptr addrspace(1) %out, i32 %x) #0 { + %h = call i32 @llvm.amdgcn.sched.vgpr.i32(i32 %x, i32 100, i32 1) + %y = add i32 %h, 1 + store i32 %y, ptr addrspace(1) %out + ret void +} + +; -------------------------------------------------------------------------- +; sched_vgpr: strict hint to v0 + +; GCN-LABEL: {{^}}vgpr0: +; GCN: v0, +define amdgpu_kernel void @vgpr0(ptr addrspace(1) %out, i32 %x) #0 { + %h = call i32 @llvm.amdgcn.sched.vgpr.i32(i32 %x, i32 0, i32 1) + %y = add i32 %h, 1 + store i32 %y, ptr addrspace(1) %out + ret void +} + +; -------------------------------------------------------------------------- +; sched_vgpr: strict hint to v512 (bank 2 territory) + +; GCN-LABEL: {{^}}vgpr512: +; GCN: v{{(512|0 /\*v512\*/)}} + +define amdgpu_kernel void @vgpr512(ptr addrspace(1) %out, i32 %x) #0 { + %h = call i32 @llvm.amdgcn.sched.vgpr.i32(i32 %x, i32 512, i32 1) + %y = add i32 %h, 1 + store i32 %y, ptr addrspace(1) %out + ret void +} + +; -------------------------------------------------------------------------- +; sched_vgpr: 64-bit value, strict hint to v[200:201] + +; GCN-LABEL: {{^}}vgpr200_64: +; GCN: v[200:201] +define amdgpu_kernel void @vgpr200_64(ptr addrspace(1) %out, i64 %x) #0 { + %h = call i64 @llvm.amdgcn.sched.vgpr.i64(i64 %x, i32 200, i32 1) + %y = add i64 %h, 1 + store i64 %y, ptr addrspace(1) %out + ret void +} + +; -------------------------------------------------------------------------- +; sched_vgpr: 512-bit value (v16f32), strict hint to v[512:527] + +; GCN-LABEL: {{^}}vgpr512_v16f32: +; GCN: v{{(\[512:513\]|\[0:1\] /\*v\[512:513\]\*/)}} + +define amdgpu_kernel void @vgpr512_v16f32(ptr addrspace(1) %out, ptr addrspace(1) %in) #0 { + %x = load <16 x float>, ptr addrspace(1) %in + %h = call <16 x float> @llvm.amdgcn.sched.vgpr.v16f32(<16 x float> %x, i32 512, i32 1) + store <16 x float> %h, ptr addrspace(1) %out + ret void +} + +; -------------------------------------------------------------------------- +; Without hint: register allocator uses default low-numbered registers. +; Verify no high-numbered VGPRs appear. + +; GCN-LABEL: {{^}}nohint: +; GCN-NOT: v42 +; GCN-NOT: v100 +; GCN-NOT: v512 +; GCN: s_endpgm +define amdgpu_kernel void @nohint(ptr addrspace(1) %out, i32 %x) #0 { + %y = add i32 %x, 1 + store i32 %y, ptr addrspace(1) %out + ret void +} + +attributes #0 = { "amdgpu-flat-work-group-size"="32,32" "amdgpu-waves-per-eu"="1,1" } diff --git a/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td b/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td index 716fd92e810ace..fe3dfc164264bd 100644 --- a/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td +++ b/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td @@ -813,6 +813,79 @@ def ROCDL_IglpOpt : ROCDL_ConcreteNonMemIntrOp<"iglp.opt", [], 0, [0], ["variant }]; } +class ROCDL_ScheduleGPROp<string mnemonic, string intrName, string gprKind> + : ROCDL_IntrOp<mnemonic, [], [0], + [AllTypesMatch<["res", "value"]>], 1, 0, 0, 0, [], []>, + Arguments<(ins LLVM_Type:$value, I32Attr:$regnum, + OptionalAttr<UnitAttr>:$soft)> { + let results = (outs LLVM_Type:$res); + let assemblyFormat = [{ + $value `,` $regnum (`soft` $soft^)? attr-dict `:` type($value) + }]; + string llvmBuilder = [{ + auto schedOp = cast<ROCDL::Schedule}] # gprKind # [{Op>(opInst); + auto *value = moduleTranslation.lookupValue(schedOp.getValue()); + int32_t regnum = schedOp.getRegnum(); + int32_t strict = schedOp.getSoftAttr() ? 0 : 1; + auto *fn = llvm::Intrinsic::getOrInsertDeclaration( + moduleTranslation.getLLVMModule(), + llvm::Intrinsic::amdgcn_sched_}] # intrName # [{, + {value->getType()}); + auto *inst = builder.CreateCall(fn, + {value, builder.getInt32(regnum), builder.getInt32(strict)}); + moduleTranslation.mapValue(opInst.getResult(0), inst); + }]; +} + +def ROCDL_ScheduleVGPROp : ROCDL_ScheduleGPROp<"sched.vgpr", "vgpr", "VGPR"> { + let description = [{ + VGPR register scheduling hint. Identity at runtime; guides the register + allocator to place `value` in the specified VGPR register number. + + For wide values (64/128/256 bit), `regnum` is the starting register index. + Default is strict (allocator must obey). Pass `soft` for an advisory hint. + + Example: + ```mlir + %0 = rocdl.sched.vgpr %val, 42 : f32 // strict: allocate to v42 + %1 = rocdl.sched.vgpr %val, 10 soft : f32 // soft: prefer v10 + ``` + }]; +} + +def ROCDL_ScheduleSGPROp : ROCDL_ScheduleGPROp<"sched.sgpr", "sgpr", "SGPR"> { + let description = [{ + SGPR register scheduling hint. Identity at runtime; guides the register + allocator to place `value` in the specified SGPR register number. + + For wide values (64/128/256 bit), `regnum` is the starting register index. + Default is strict (allocator must obey). Pass `soft` for an advisory hint. + + Example: + ```mlir + %0 = rocdl.sched.sgpr %val, 10 : i32 // strict: allocate to s10 + %1 = rocdl.sched.sgpr %val, 4 soft : i32 // soft: prefer s4 + ``` + }]; +} + +def ROCDL_ScheduleAGPROp : ROCDL_ScheduleGPROp<"sched.agpr", "agpr", "AGPR"> { + let description = [{ + AGPR register scheduling hint. Identity at runtime; guides the register + allocator to place `value` in the specified AGPR register number. + Only valid on targets with MAI instructions (e.g., gfx90a, gfx940). + + For wide values (64/128/256 bit), `regnum` is the starting register index. + Default is strict (allocator must obey). Pass `soft` for an advisory hint. + + Example: + ```mlir + %0 = rocdl.sched.agpr %val, 8 : f32 // strict: allocate to a8 + %1 = rocdl.sched.agpr %val, 0 soft : f32 // soft: prefer a0 + ``` + }]; +} + //===---------------------------------------------------------------------===// // Xdlops intrinsics diff --git a/mlir/test/Dialect/LLVMIR/rocdl.mlir b/mlir/test/Dialect/LLVMIR/rocdl.mlir index 1b5b6419368819..ff62db8d69a253 100644 --- a/mlir/test/Dialect/LLVMIR/rocdl.mlir +++ b/mlir/test/Dialect/LLVMIR/rocdl.mlir @@ -168,6 +168,30 @@ func.func @rocdl_iglp_opt() { llvm.return } +func.func @rocdl_sched_vgpr(%val : f32, %ival : i32) -> f32 { + // CHECK: rocdl.sched.vgpr %{{.*}}, 42 : f32 + %0 = rocdl.sched.vgpr %val, 42 : f32 + // CHECK: rocdl.sched.vgpr %{{.*}}, 10 soft : i32 + %1 = rocdl.sched.vgpr %ival, 10 soft : i32 + llvm.return %0 : f32 +} + +func.func @rocdl_sched_sgpr(%val : f32, %ival : i32) -> f32 { + // CHECK: rocdl.sched.sgpr %{{.*}}, 20 : f32 + %0 = rocdl.sched.sgpr %val, 20 : f32 + // CHECK: rocdl.sched.sgpr %{{.*}}, 5 soft : i32 + %1 = rocdl.sched.sgpr %ival, 5 soft : i32 + llvm.return %0 : f32 +} + +func.func @rocdl_sched_agpr(%val : f32, %ival : i32) -> f32 { + // CHECK: rocdl.sched.agpr %{{.*}}, 8 : f32 + %0 = rocdl.sched.agpr %val, 8 : f32 + // CHECK: rocdl.sched.agpr %{{.*}}, 0 soft : i32 + %1 = rocdl.sched.agpr %ival, 0 soft : i32 + llvm.return %0 : f32 +} + func.func @rocdl.setprio() { // CHECK: rocdl.s.setprio rocdl.s.setprio 0 diff --git a/mlir/test/Target/LLVMIR/rocdl.mlir b/mlir/test/Target/LLVMIR/rocdl.mlir index 2f067b8ba98fda..193f02bf674941 100644 --- a/mlir/test/Target/LLVMIR/rocdl.mlir +++ b/mlir/test/Target/LLVMIR/rocdl.mlir @@ -451,6 +451,35 @@ llvm.func @rocdl.iglp.opt() { llvm.return } +llvm.func @rocdl.sched.vgpr(%val : f32, %ival : i32) -> f32 { + // CHECK-LABEL: rocdl.sched.vgpr + // strict (default): regnum 42, strict 1 + // CHECK: call float @llvm.amdgcn.sched.vgpr.f32(float %{{.*}}, i32 42, i32 1) + %0 = rocdl.sched.vgpr %val, 42 : f32 + // soft: regnum 10, strict 0 + // CHECK: call i32 @llvm.amdgcn.sched.vgpr.i32(i32 %{{.*}}, i32 10, i32 0) + %1 = rocdl.sched.vgpr %ival, 10 soft : i32 + llvm.return %0 : f32 +} + +llvm.func @rocdl.sched.sgpr(%val : f32, %ival : i32) -> f32 { + // CHECK-LABEL: rocdl.sched.sgpr + // CHECK: call float @llvm.amdgcn.sched.sgpr.f32(float %{{.*}}, i32 20, i32 1) + %0 = rocdl.sched.sgpr %val, 20 : f32 + // CHECK: call i32 @llvm.amdgcn.sched.sgpr.i32(i32 %{{.*}}, i32 5, i32 0) + %1 = rocdl.sched.sgpr %ival, 5 soft : i32 + llvm.return %0 : f32 +} + +llvm.func @rocdl.sched.agpr(%val : f32, %ival : i32) -> f32 { + // CHECK-LABEL: rocdl.sched.agpr + // CHECK: call float @llvm.amdgcn.sched.agpr.f32(float %{{.*}}, i32 8, i32 1) + %0 = rocdl.sched.agpr %val, 8 : f32 + // CHECK: call i32 @llvm.amdgcn.sched.agpr.i32(i32 %{{.*}}, i32 0, i32 0) + %1 = rocdl.sched.agpr %ival, 0 soft : i32 + llvm.return %0 : f32 +} + llvm.func @rocdl.xdlops(%arg0 : f32, %arg1 : f32, %arg2 : vector<32 x f32>, %arg3: i32, %arg4 : vector<16 x f32>, %arg5 : vector<4xf32>, >From 0aed547bb5ab54ca745881f116a405b65e05eb26 Mon Sep 17 00:00:00 2001 From: Jun Li <[email protected]> Date: Tue, 20 Oct 2026 10:01:17 +0000 Subject: [PATCH 2/4] [AMDGPU] Add AMDGPUSchedGPR pass to lower sched_vgpr/sgpr/agpr pseudos Add a pre-RA machine function pass that lowers V_SCHEDULE_VGPR_B*, S_SCHEDULE_SGPR_B*, and A_SCHEDULE_AGPR_B* pseudos into plain COPYs while attaching AMDGPURI register allocation hints to guide the register allocator to the requested physical register numbers. The pass propagates hints through COPY/REG_SEQUENCE/INSERT_SUBREG/ SUBREG_TO_REG so the preference survives coalescing. Strict hints cause getRegAllocationHints to return true, making the allocator treat the hint as a hard constraint. Co-Authored-By: Claude Opus 4 <[email protected]> --- llvm/lib/Target/AMDGPU/AMDGPU.h | 3 + llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def | 1 + llvm/lib/Target/AMDGPU/AMDGPUSchedGPR.cpp | 212 ++++++++++++++++++ llvm/lib/Target/AMDGPU/AMDGPUSchedGPR.h | 22 ++ .../lib/Target/AMDGPU/AMDGPUTargetMachine.cpp | 4 + llvm/lib/Target/AMDGPU/CMakeLists.txt | 1 + llvm/test/CodeGen/AMDGPU/sched-gpr-hints.mir | 4 +- 7 files changed, 245 insertions(+), 2 deletions(-) create mode 100644 llvm/lib/Target/AMDGPU/AMDGPUSchedGPR.cpp create mode 100644 llvm/lib/Target/AMDGPU/AMDGPUSchedGPR.h diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.h b/llvm/lib/Target/AMDGPU/AMDGPU.h index 75a83c55db2bc9..12fb9095ceafc6 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPU.h +++ b/llvm/lib/Target/AMDGPU/AMDGPU.h @@ -651,6 +651,9 @@ class AMDGPURewriteAGPRCopyMFMAPass void initializeAMDGPURewriteAGPRCopyMFMALegacyPass(PassRegistry &); extern char &AMDGPURewriteAGPRCopyMFMALegacyID; +void initializeAMDGPUSchedGPRPass(PassRegistry &); +extern char &AMDGPUSchedGPRID; + void initializeAMDGPUUniformIntrinsicCombineLegacyPass(PassRegistry &); extern char &AMDGPUUniformIntrinsicCombineLegacyPassID; FunctionPass *createAMDGPUUniformIntrinsicCombineLegacyPass(); diff --git a/llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def b/llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def index 372d5f5acab21a..97111a43198666 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def +++ b/llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def @@ -139,6 +139,7 @@ MACHINE_FUNCTION_PASS("amdgpu-preload-kern-arg-prolog", AMDGPUPreloadKernArgProl MACHINE_FUNCTION_PASS("amdgpu-prepare-agpr-alloc", AMDGPUPrepareAGPRAllocPass()) MACHINE_FUNCTION_PASS("amdgpu-nsa-reassign", GCNNSAReassignPass()) MACHINE_FUNCTION_PASS("amdgpu-wait-sgpr-hazards", AMDGPUWaitSGPRHazardsPass()) +MACHINE_FUNCTION_PASS("amdgpu-sched-gpr", AMDGPUSchedGPRPass()) MACHINE_FUNCTION_PASS("gcn-create-vopd", GCNCreateVOPDPass()) MACHINE_FUNCTION_PASS("gcn-dpp-combine", GCNDPPCombinePass()) MACHINE_FUNCTION_PASS("print<amdgpu-next-use-analysis>", AMDGPUNextUseAnalysisPrinterPass(errs())) diff --git a/llvm/lib/Target/AMDGPU/AMDGPUSchedGPR.cpp b/llvm/lib/Target/AMDGPU/AMDGPUSchedGPR.cpp new file mode 100644 index 00000000000000..85b4776db6a58a --- /dev/null +++ b/llvm/lib/Target/AMDGPU/AMDGPUSchedGPR.cpp @@ -0,0 +1,212 @@ +//===-- AMDGPUSchedGPR.cpp ------------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// +/// \file +/// Lower the V_SCHEDULE_VGPR_B*, S_SCHEDULE_SGPR_B*, and A_SCHEDULE_AGPR_B* +/// pseudos produced from llvm.amdgcn.sched.{vgpr,sgpr,agpr}. For each pseudo +/// this pass: +/// 1. reads the requested register number and strict flag from the immediate +/// operands, +/// 2. attaches an AMDGPURI::{VGPR,SGPR,AGPR}RegHint register allocation hint +/// to the destination (and source) vreg so the allocator prefers that exact +/// physical register, +/// 3. propagates the hint through COPY / REG_SEQUENCE / INSERT_SUBREG / +/// SUBREG_TO_REG so it survives coalescing, and +/// 4. replaces the pseudo with a plain COPY. +/// +/// The hint is advisory unless strict: SIRegisterInfo::getRegAllocationHints +/// only reorders the allocation candidates, so a non-strict hint falls back to +/// the default order and the pass never forces a spill. +// +//===----------------------------------------------------------------------===// + +#include "AMDGPUSchedGPR.h" +#include "AMDGPU.h" +#include "GCNSubtarget.h" +#include "SIInstrInfo.h" +#include "SIRegisterInfo.h" +#include "llvm/ADT/DenseSet.h" +#include "llvm/CodeGen/MachineFunctionPass.h" +#include "llvm/CodeGen/MachineInstrBuilder.h" +#include "llvm/CodeGen/MachineRegisterInfo.h" +#include "llvm/InitializePasses.h" + +using namespace llvm; + +#define DEBUG_TYPE "amdgpu-sched-gpr" + +enum ScheduleGPRKind { SGK_None, SGK_VGPR, SGK_SGPR, SGK_AGPR }; + +static ScheduleGPRKind getScheduleGPRKind(unsigned Opcode) { + switch (Opcode) { + case AMDGPU::V_SCHEDULE_VGPR_B32: + case AMDGPU::V_SCHEDULE_VGPR_B64: + case AMDGPU::V_SCHEDULE_VGPR_B128: + case AMDGPU::V_SCHEDULE_VGPR_B256: + case AMDGPU::V_SCHEDULE_VGPR_B512: + return SGK_VGPR; + case AMDGPU::S_SCHEDULE_SGPR_B32: + case AMDGPU::S_SCHEDULE_SGPR_B64: + case AMDGPU::S_SCHEDULE_SGPR_B128: + case AMDGPU::S_SCHEDULE_SGPR_B256: + case AMDGPU::S_SCHEDULE_SGPR_B512: + return SGK_SGPR; + case AMDGPU::A_SCHEDULE_AGPR_B32: + case AMDGPU::A_SCHEDULE_AGPR_B64: + case AMDGPU::A_SCHEDULE_AGPR_B128: + case AMDGPU::A_SCHEDULE_AGPR_B256: + case AMDGPU::A_SCHEDULE_AGPR_B512: + return SGK_AGPR; + default: + return SGK_None; + } +} + +/// Walk the SSA use chain of \p Reg and propagate the hint through +/// value-preserving pseudos so the preference reaches the vregs that survive +/// coalescing. +static void propagateHint(MachineRegisterInfo &MRI, Register Reg, + unsigned HintValue, unsigned HintKind, + SmallDenseSet<unsigned, 32> &Visited) { + if (!Reg.isVirtual() || !Visited.insert(Reg.id()).second) + return; + + for (MachineInstr &UseMI : MRI.use_nodbg_instructions(Reg)) { + Register DefReg; + switch (UseMI.getOpcode()) { + case TargetOpcode::COPY: + case TargetOpcode::REG_SEQUENCE: + case TargetOpcode::INSERT_SUBREG: + case TargetOpcode::SUBREG_TO_REG: + DefReg = UseMI.getOperand(0).getReg(); + break; + default: + continue; + } + if (!DefReg.isVirtual()) + continue; + + // Do not clobber an existing, conflicting hint. + std::pair<unsigned, Register> Existing = MRI.getRegAllocationHint(DefReg); + if (Existing.first != 0 && + (Existing.first != HintKind || Existing.second != HintValue)) + continue; + + MRI.setRegAllocationHint(DefReg, HintKind, HintValue); + propagateHint(MRI, DefReg, HintValue, HintKind, Visited); + } +} + +static bool runSchedGPR(MachineFunction &MF) { + MachineRegisterInfo &MRI = MF.getRegInfo(); + const GCNSubtarget &ST = MF.getSubtarget<GCNSubtarget>(); + const SIInstrInfo *TII = ST.getInstrInfo(); + bool Changed = false; + + SmallVector<MachineInstr *, 16> ToErase; + + for (MachineBasicBlock &MBB : MF) { + for (MachineInstr &MI : MBB) { + ScheduleGPRKind Kind = getScheduleGPRKind(MI.getOpcode()); + if (Kind == SGK_None) + continue; + + Register DstReg = MI.getOperand(0).getReg(); + const MachineOperand &SrcMO = MI.getOperand(1); + Register SrcReg = SrcMO.getReg(); + unsigned RegNum = MI.getOperand(2).getImm(); + bool Strict = MI.getOperand(3).getImm() != 0; + + unsigned HintKind; + switch (Kind) { + case SGK_VGPR: + HintKind = Strict ? AMDGPURI::StrictVGPRRegHint + : AMDGPURI::VGPRRegHint; + break; + case SGK_SGPR: + HintKind = Strict ? AMDGPURI::StrictSGPRRegHint + : AMDGPURI::SGPRRegHint; + break; + case SGK_AGPR: + HintKind = Strict ? AMDGPURI::StrictAGPRRegHint + : AMDGPURI::AGPRRegHint; + break; + default: + llvm_unreachable("unexpected schedule GPR kind"); + } + + LLVM_DEBUG(dbgs() << " sched." + << (Kind == SGK_VGPR ? "vgpr" + : Kind == SGK_SGPR ? "sgpr" + : "agpr") + << ": " << printReg(DstReg) << " <- " + << printReg(SrcReg) << " reg " << RegNum << " strict " + << Strict << '\n'); + + // Hint both the destination and the source so the preference survives + // whichever side coalescing keeps. + if (DstReg.isVirtual()) + MRI.setRegAllocationHint(DstReg, HintKind, RegNum); + if (SrcReg.isVirtual()) + MRI.setRegAllocationHint(SrcReg, HintKind, RegNum); + + SmallDenseSet<unsigned, 32> Visited; + if (DstReg.isVirtual()) + propagateHint(MRI, DstReg, RegNum, HintKind, Visited); + + // Replace the pseudo with a plain COPY. + BuildMI(MBB, MI, MI.getDebugLoc(), TII->get(TargetOpcode::COPY), DstReg) + .addReg(SrcReg, getRegState(SrcMO), SrcMO.getSubReg()); + + ToErase.push_back(&MI); + Changed = true; + } + } + + for (MachineInstr *MI : ToErase) + MI->eraseFromParent(); + + return Changed; +} + +namespace { + +class AMDGPUSchedGPR : public MachineFunctionPass { +public: + static char ID; + + AMDGPUSchedGPR() : MachineFunctionPass(ID) {} + + bool runOnMachineFunction(MachineFunction &MF) override { + return runSchedGPR(MF); + } + + StringRef getPassName() const override { return "AMDGPU Sched GPR"; } + + void getAnalysisUsage(AnalysisUsage &AU) const override { + AU.setPreservesCFG(); + MachineFunctionPass::getAnalysisUsage(AU); + } +}; + +} // end anonymous namespace + +INITIALIZE_PASS(AMDGPUSchedGPR, DEBUG_TYPE, "AMDGPU Sched GPR", false, false) + +char AMDGPUSchedGPR::ID = 0; + +char &llvm::AMDGPUSchedGPRID = AMDGPUSchedGPR::ID; + +PreservedAnalyses +AMDGPUSchedGPRPass::run(MachineFunction &MF, + MachineFunctionAnalysisManager &MFAM) { + if (!runSchedGPR(MF)) + return PreservedAnalyses::all(); + + return getMachineFunctionPassPreservedAnalyses().preserveSet<CFGAnalyses>(); +} diff --git a/llvm/lib/Target/AMDGPU/AMDGPUSchedGPR.h b/llvm/lib/Target/AMDGPU/AMDGPUSchedGPR.h new file mode 100644 index 00000000000000..fd7c0a824c98ee --- /dev/null +++ b/llvm/lib/Target/AMDGPU/AMDGPUSchedGPR.h @@ -0,0 +1,22 @@ +//===- AMDGPUSchedGPR.h -----------------------------------------*- C++ -*-===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#ifndef LLVM_LIB_TARGET_AMDGPU_AMDGPUSCHEDGPR_H +#define LLVM_LIB_TARGET_AMDGPU_AMDGPUSCHEDGPR_H + +#include "llvm/CodeGen/MachinePassManager.h" + +namespace llvm { +class AMDGPUSchedGPRPass : public OptionalPassInfoMixin<AMDGPUSchedGPRPass> { +public: + PreservedAnalyses run(MachineFunction &MF, + MachineFunctionAnalysisManager &MFAM); +}; +} // namespace llvm + +#endif // LLVM_LIB_TARGET_AMDGPU_AMDGPUSCHEDGPR_H diff --git a/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp b/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp index 6a5b27e5775c62..137ff98b173e64 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp @@ -35,6 +35,7 @@ #include "AMDGPURemoveIncompatibleFunctions.h" #include "AMDGPUReserveWWMRegs.h" #include "AMDGPUResourceUsageAnalysis.h" +#include "AMDGPUSchedGPR.h" #include "AMDGPUSplitModule.h" #include "AMDGPUTargetObjectFile.h" #include "AMDGPUTargetTransformInfo.h" @@ -733,6 +734,7 @@ extern "C" LLVM_ABI LLVM_EXTERNAL_VISIBILITY void LLVMInitializeAMDGPUTarget() { initializeAMDGPULowerIntrinsicsLegacyPass(*PR); initializeAMDGPUReserveWWMRegsLegacyPass(*PR); initializeAMDGPURewriteAGPRCopyMFMALegacyPass(*PR); + initializeAMDGPUSchedGPRPass(*PR); initializeAMDGPURewriteOutArgumentsPass(*PR); initializeAMDGPURewriteUndefForPHILegacyPass(*PR); initializeSIAnnotateControlFlowLegacyPass(*PR); @@ -1846,6 +1848,7 @@ void GCNPassConfig::addFastRegAlloc() { } void GCNPassConfig::addPreRegAlloc() { + addPass(&AMDGPUSchedGPRID); if (getOptLevel() != CodeGenOptLevel::None) addPass(&AMDGPUPrepareAGPRAllocLegacyID); if (getOptLevel() >= CodeGenOptLevel::Default && EnableMachinePipeliner) @@ -2708,6 +2711,7 @@ Error AMDGPUCodeGenPassBuilder::addOptimizedRegAlloc(PassManagerWrapper &PMW) { } void AMDGPUCodeGenPassBuilder::addPreRegAlloc(PassManagerWrapper &PMW) { + addMachineFunctionPass(AMDGPUSchedGPRPass(), PMW); if (getOptLevel() != CodeGenOptLevel::None) addMachineFunctionPass(AMDGPUPrepareAGPRAllocPass(), PMW); if (getOptLevel() >= CodeGenOptLevel::Default && EnableMachinePipeliner) diff --git a/llvm/lib/Target/AMDGPU/CMakeLists.txt b/llvm/lib/Target/AMDGPU/CMakeLists.txt index b6c5f63379344c..7baf68c021ae1c 100644 --- a/llvm/lib/Target/AMDGPU/CMakeLists.txt +++ b/llvm/lib/Target/AMDGPU/CMakeLists.txt @@ -113,6 +113,7 @@ add_llvm_target(AMDGPUCodeGen AMDGPUResourceUsageAnalysis.cpp AMDGPURewriteAGPRCopyMFMA.cpp AMDGPURewriteOutArguments.cpp + AMDGPUSchedGPR.cpp AMDGPURewriteUndefForPHI.cpp AMDGPUSelectionDAGInfo.cpp AMDGPUSetWavePriority.cpp diff --git a/llvm/test/CodeGen/AMDGPU/sched-gpr-hints.mir b/llvm/test/CodeGen/AMDGPU/sched-gpr-hints.mir index d480eea0ce596b..a1e178ece02ed9 100644 --- a/llvm/test/CodeGen/AMDGPU/sched-gpr-hints.mir +++ b/llvm/test/CodeGen/AMDGPU/sched-gpr-hints.mir @@ -1,6 +1,6 @@ -# RUN: llc -mtriple=amdgcn -mcpu=gfx1250 -run-pass=amdgpu-sched-bank -o - %s | FileCheck %s +# RUN: llc -mtriple=amdgcn -mcpu=gfx1250 -run-pass=amdgpu-sched-gpr -o - %s | FileCheck %s -# The AMDGPUSchedBank pass attaches exact register allocation hints (not +# The AMDGPUSchedGPR pass attaches exact register allocation hints (not # printed in MIR) and lowers each schedule pseudo to a plain COPY. --- >From 72113c3a42c257ac5eda5ce13a533fd537f0969f Mon Sep 17 00:00:00 2001 From: Jun Li <[email protected]> Date: Tue, 20 Oct 2026 10:04:02 +0000 Subject: [PATCH 3/4] [AMDGPU] Add AMDGPUSchedGPR to llc pipeline tests Update llc-pipeline.ll and llc-pipeline-npm.ll to include the amdgpu-sched-gpr pass in the expected pass ordering for all optimization levels (O0-O3). Co-Authored-By: Claude Opus 4 <[email protected]> --- llvm/test/CodeGen/AMDGPU/llc-pipeline-npm.ll | 3 +++ llvm/test/CodeGen/AMDGPU/llc-pipeline.ll | 5 +++++ 2 files changed, 8 insertions(+) diff --git a/llvm/test/CodeGen/AMDGPU/llc-pipeline-npm.ll b/llvm/test/CodeGen/AMDGPU/llc-pipeline-npm.ll index 3e44687d1856ed..73d5e1ea7e9583 100644 --- a/llvm/test/CodeGen/AMDGPU/llc-pipeline-npm.ll +++ b/llvm/test/CodeGen/AMDGPU/llc-pipeline-npm.ll @@ -68,6 +68,7 @@ ; GCN-O0-NEXT: function ; GCN-O0-NEXT: machine-function ; GCN-O0-NEXT: reg-usage-propagation +; GCN-O0-NEXT: amdgpu-sched-gpr ; GCN-O0-NEXT: phi-node-elimination ; GCN-O0-NEXT: si-lower-control-flow ; GCN-O0-NEXT: two-address-instruction @@ -215,6 +216,7 @@ ; GCN-O2-NEXT: function ; GCN-O2-NEXT: machine-function ; GCN-O2-NEXT: reg-usage-propagation +; GCN-O2-NEXT: amdgpu-sched-gpr ; GCN-O2-NEXT: amdgpu-prepare-agpr-alloc ; GCN-O2-NEXT: detect-dead-lanes ; GCN-O2-NEXT: dead-mi-elimination @@ -404,6 +406,7 @@ ; GCN-O3-NEXT: function ; GCN-O3-NEXT: machine-function ; GCN-O3-NEXT: reg-usage-propagation +; GCN-O3-NEXT: amdgpu-sched-gpr ; GCN-O3-NEXT: amdgpu-prepare-agpr-alloc ; GCN-O3-NEXT: detect-dead-lanes ; GCN-O3-NEXT: dead-mi-elimination diff --git a/llvm/test/CodeGen/AMDGPU/llc-pipeline.ll b/llvm/test/CodeGen/AMDGPU/llc-pipeline.ll index a97c1955c94f5c..e9e4b8d364ec95 100644 --- a/llvm/test/CodeGen/AMDGPU/llc-pipeline.ll +++ b/llvm/test/CodeGen/AMDGPU/llc-pipeline.ll @@ -104,6 +104,7 @@ ; GCN-O0-NEXT: Finalize ISel and expand pseudo-instructions ; GCN-O0-NEXT: Local Stack Slot Allocation ; GCN-O0-NEXT: Register Usage Information Propagation +; GCN-O0-NEXT: AMDGPU Sched GPR ; GCN-O0-NEXT: Eliminate PHI nodes for register allocation ; GCN-O0-NEXT: SI Lower control flow pseudo instructions ; GCN-O0-NEXT: Two-Address instruction pass @@ -344,6 +345,7 @@ ; GCN-O1-NEXT: Remove dead machine instructions ; GCN-O1-NEXT: SI Shrink Instructions ; GCN-O1-NEXT: Register Usage Information Propagation +; GCN-O1-NEXT: AMDGPU Sched GPR ; GCN-O1-NEXT: AMDGPU Prepare AGPR Alloc ; GCN-O1-NEXT: Detect Dead Lanes ; GCN-O1-NEXT: Remove dead machine instructions @@ -673,6 +675,7 @@ ; GCN-O1-OPTS-NEXT: Remove dead machine instructions ; GCN-O1-OPTS-NEXT: SI Shrink Instructions ; GCN-O1-OPTS-NEXT: Register Usage Information Propagation +; GCN-O1-OPTS-NEXT: AMDGPU Sched GPR ; GCN-O1-OPTS-NEXT: AMDGPU Prepare AGPR Alloc ; GCN-O1-OPTS-NEXT: Detect Dead Lanes ; GCN-O1-OPTS-NEXT: Remove dead machine instructions @@ -1006,6 +1009,7 @@ ; GCN-O2-NEXT: Remove dead machine instructions ; GCN-O2-NEXT: SI Shrink Instructions ; GCN-O2-NEXT: Register Usage Information Propagation +; GCN-O2-NEXT: AMDGPU Sched GPR ; GCN-O2-NEXT: AMDGPU Prepare AGPR Alloc ; GCN-O2-NEXT: Detect Dead Lanes ; GCN-O2-NEXT: Remove dead machine instructions @@ -1355,6 +1359,7 @@ ; GCN-O3-NEXT: Remove dead machine instructions ; GCN-O3-NEXT: SI Shrink Instructions ; GCN-O3-NEXT: Register Usage Information Propagation +; GCN-O3-NEXT: AMDGPU Sched GPR ; GCN-O3-NEXT: AMDGPU Prepare AGPR Alloc ; GCN-O3-NEXT: Detect Dead Lanes ; GCN-O3-NEXT: Remove dead machine instructions >From b5758c40f5f4e3ef17d34fd06714a7966dfc3788 Mon Sep 17 00:00:00 2001 From: Jun Li <[email protected]> Date: Tue, 20 Oct 2026 10:27:18 +0000 Subject: [PATCH 4/4] [AMDGPU] Fix clang-format issues in sched_gpr code Co-Authored-By: Claude Opus 4 <[email protected]> --- clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp | 12 ++-- clang/lib/Sema/SemaAMDGPU.cpp | 22 +++--- llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp | 68 +++++++++++++------ llvm/lib/Target/AMDGPU/AMDGPUSchedGPR.cpp | 13 ++-- llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp | 1 - 5 files changed, 76 insertions(+), 40 deletions(-) diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp index c27a8e8f2a9fcf..ab57c9b1126ceb 100644 --- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp +++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp @@ -2299,12 +2299,16 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID, Intrinsic::ID IID; switch (BuiltinID) { case AMDGPU::BI__builtin_amdgcn_sched_vgpr: - IID = Intrinsic::amdgcn_sched_vgpr; break; + IID = Intrinsic::amdgcn_sched_vgpr; + break; case AMDGPU::BI__builtin_amdgcn_sched_sgpr: - IID = Intrinsic::amdgcn_sched_sgpr; break; + IID = Intrinsic::amdgcn_sched_sgpr; + break; case AMDGPU::BI__builtin_amdgcn_sched_agpr: - IID = Intrinsic::amdgcn_sched_agpr; break; - default: llvm_unreachable("unexpected builtin"); + IID = Intrinsic::amdgcn_sched_agpr; + break; + default: + llvm_unreachable("unexpected builtin"); } llvm::Function *F = CGM.getIntrinsic(IID, Val->getType()); return Builder.CreateCall(F, {Val, RegNum, Strict}); diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp index 556526394d2872..95d5842d3baf50 100644 --- a/clang/lib/Sema/SemaAMDGPU.cpp +++ b/clang/lib/Sema/SemaAMDGPU.cpp @@ -474,10 +474,9 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(const TargetInfo &TI, QualType DataTy = DataArg->getType(); if (DataTy->isAnyComplexType() || !(DataTy->isArithmeticType() || - (DataTy->isVectorType() && - DataTy->castAs<VectorType>() - ->getElementType() - ->isArithmeticType()))) { + (DataTy->isVectorType() && DataTy->castAs<VectorType>() + ->getElementType() + ->isArithmeticType()))) { SemaRef.Diag(DataArg->getBeginLoc(), diag::err_typecheck_cond_expect_int_float) << DataTy << DataArg->getSourceRange(); @@ -486,10 +485,17 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(const TargetInfo &TI, TheCall->setType(DataTy); int MaxReg; switch (BuiltinID) { - case AMDGPU::BI__builtin_amdgcn_sched_vgpr: MaxReg = 1023; break; - case AMDGPU::BI__builtin_amdgcn_sched_sgpr: MaxReg = 105; break; - case AMDGPU::BI__builtin_amdgcn_sched_agpr: MaxReg = 255; break; - default: llvm_unreachable("unexpected builtin"); + case AMDGPU::BI__builtin_amdgcn_sched_vgpr: + MaxReg = 1023; + break; + case AMDGPU::BI__builtin_amdgcn_sched_sgpr: + MaxReg = 105; + break; + case AMDGPU::BI__builtin_amdgcn_sched_agpr: + MaxReg = 255; + break; + default: + llvm_unreachable("unexpected builtin"); } return SemaRef.BuiltinConstantArgRange(TheCall, 1, 0, MaxReg); } diff --git a/llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp b/llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp index a6c20d61be838f..072725db57338f 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp @@ -3367,11 +3367,21 @@ void AMDGPUDAGToDAGISel::SelectINTRINSIC_WO_CHAIN(SDNode *N) { unsigned PseudoOpc; if (IntrID == Intrinsic::amdgcn_sched_vgpr) { switch (Bits) { - case 32: PseudoOpc = AMDGPU::V_SCHEDULE_VGPR_B32; break; - case 64: PseudoOpc = AMDGPU::V_SCHEDULE_VGPR_B64; break; - case 128: PseudoOpc = AMDGPU::V_SCHEDULE_VGPR_B128; break; - case 256: PseudoOpc = AMDGPU::V_SCHEDULE_VGPR_B256; break; - case 512: PseudoOpc = AMDGPU::V_SCHEDULE_VGPR_B512; break; + case 32: + PseudoOpc = AMDGPU::V_SCHEDULE_VGPR_B32; + break; + case 64: + PseudoOpc = AMDGPU::V_SCHEDULE_VGPR_B64; + break; + case 128: + PseudoOpc = AMDGPU::V_SCHEDULE_VGPR_B128; + break; + case 256: + PseudoOpc = AMDGPU::V_SCHEDULE_VGPR_B256; + break; + case 512: + PseudoOpc = AMDGPU::V_SCHEDULE_VGPR_B512; + break; default: ReplaceUses(SDValue(N, 0), Src); CurDAG->RemoveDeadNode(N); @@ -3379,11 +3389,21 @@ void AMDGPUDAGToDAGISel::SelectINTRINSIC_WO_CHAIN(SDNode *N) { } } else if (IntrID == Intrinsic::amdgcn_sched_sgpr) { switch (Bits) { - case 32: PseudoOpc = AMDGPU::S_SCHEDULE_SGPR_B32; break; - case 64: PseudoOpc = AMDGPU::S_SCHEDULE_SGPR_B64; break; - case 128: PseudoOpc = AMDGPU::S_SCHEDULE_SGPR_B128; break; - case 256: PseudoOpc = AMDGPU::S_SCHEDULE_SGPR_B256; break; - case 512: PseudoOpc = AMDGPU::S_SCHEDULE_SGPR_B512; break; + case 32: + PseudoOpc = AMDGPU::S_SCHEDULE_SGPR_B32; + break; + case 64: + PseudoOpc = AMDGPU::S_SCHEDULE_SGPR_B64; + break; + case 128: + PseudoOpc = AMDGPU::S_SCHEDULE_SGPR_B128; + break; + case 256: + PseudoOpc = AMDGPU::S_SCHEDULE_SGPR_B256; + break; + case 512: + PseudoOpc = AMDGPU::S_SCHEDULE_SGPR_B512; + break; default: ReplaceUses(SDValue(N, 0), Src); CurDAG->RemoveDeadNode(N); @@ -3391,21 +3411,31 @@ void AMDGPUDAGToDAGISel::SelectINTRINSIC_WO_CHAIN(SDNode *N) { } } else { switch (Bits) { - case 32: PseudoOpc = AMDGPU::A_SCHEDULE_AGPR_B32; break; - case 64: PseudoOpc = AMDGPU::A_SCHEDULE_AGPR_B64; break; - case 128: PseudoOpc = AMDGPU::A_SCHEDULE_AGPR_B128; break; - case 256: PseudoOpc = AMDGPU::A_SCHEDULE_AGPR_B256; break; - case 512: PseudoOpc = AMDGPU::A_SCHEDULE_AGPR_B512; break; + case 32: + PseudoOpc = AMDGPU::A_SCHEDULE_AGPR_B32; + break; + case 64: + PseudoOpc = AMDGPU::A_SCHEDULE_AGPR_B64; + break; + case 128: + PseudoOpc = AMDGPU::A_SCHEDULE_AGPR_B128; + break; + case 256: + PseudoOpc = AMDGPU::A_SCHEDULE_AGPR_B256; + break; + case 512: + PseudoOpc = AMDGPU::A_SCHEDULE_AGPR_B512; + break; default: ReplaceUses(SDValue(N, 0), Src); CurDAG->RemoveDeadNode(N); return; } } - SDValue RegNumImm = CurDAG->getTargetConstant(RegNumC->getZExtValue(), - SDLoc(N), MVT::i32); - SDValue StrictImm = CurDAG->getTargetConstant(StrictC->getZExtValue(), - SDLoc(N), MVT::i32); + SDValue RegNumImm = + CurDAG->getTargetConstant(RegNumC->getZExtValue(), SDLoc(N), MVT::i32); + SDValue StrictImm = + CurDAG->getTargetConstant(StrictC->getZExtValue(), SDLoc(N), MVT::i32); CurDAG->SelectNodeTo(N, PseudoOpc, N->getVTList(), {Src, RegNumImm, StrictImm}); return; diff --git a/llvm/lib/Target/AMDGPU/AMDGPUSchedGPR.cpp b/llvm/lib/Target/AMDGPU/AMDGPUSchedGPR.cpp index 85b4776db6a58a..8fbcafd21c8d6c 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUSchedGPR.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUSchedGPR.cpp @@ -13,8 +13,8 @@ /// 1. reads the requested register number and strict flag from the immediate /// operands, /// 2. attaches an AMDGPURI::{VGPR,SGPR,AGPR}RegHint register allocation hint -/// to the destination (and source) vreg so the allocator prefers that exact -/// physical register, +/// to the destination (and source) vreg so the allocator prefers that +/// exact physical register, /// 3. propagates the hint through COPY / REG_SEQUENCE / INSERT_SUBREG / /// SUBREG_TO_REG so it survives coalescing, and /// 4. replaces the pseudo with a plain COPY. @@ -125,16 +125,13 @@ static bool runSchedGPR(MachineFunction &MF) { unsigned HintKind; switch (Kind) { case SGK_VGPR: - HintKind = Strict ? AMDGPURI::StrictVGPRRegHint - : AMDGPURI::VGPRRegHint; + HintKind = Strict ? AMDGPURI::StrictVGPRRegHint : AMDGPURI::VGPRRegHint; break; case SGK_SGPR: - HintKind = Strict ? AMDGPURI::StrictSGPRRegHint - : AMDGPURI::SGPRRegHint; + HintKind = Strict ? AMDGPURI::StrictSGPRRegHint : AMDGPURI::SGPRRegHint; break; case SGK_AGPR: - HintKind = Strict ? AMDGPURI::StrictAGPRRegHint - : AMDGPURI::AGPRRegHint; + HintKind = Strict ? AMDGPURI::StrictAGPRRegHint : AMDGPURI::AGPRRegHint; break; default: llvm_unreachable("unexpected schedule GPR kind"); diff --git a/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp b/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp index 0bc47414f52232..ad83a89d758839 100644 --- a/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp +++ b/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp @@ -55,7 +55,6 @@ static cl::opt<unsigned> StressSGPRLimit( "amdgpu-stress-sgpr", cl::Hidden, cl::init(0), cl::desc("Limit SGPRs to N registers by reserving the rest")); - std::array<std::vector<int16_t>, 32> SIRegisterInfo::RegSplitParts; std::array<std::array<uint16_t, 32>, 9> SIRegisterInfo::SubRegFromChannelTable; _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
