[Mlir-commits] [clang] [llvm] [mlir] [AMDGPU] Add sched_vgpr/sgpr/agpr intrinsics for GPR register hints (PR #233773)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Tue Oct 20 03:21:28 PDT 2026
github-actions[bot] wrote:
<!--LLVM CODE FORMAT COMMENT: {clang-format}-->
:warning: C/C++ code formatter, clang-format found issues in your code. :warning:
<details>
<summary>
You can test this locally with the following command:
</summary>
``````````bash
git-clang-format --diff origin/main HEAD --extensions h,cl,cpp -- clang/test/CodeGenOpenCL/builtins-amdgcn-sched-gpr.cl llvm/lib/Target/AMDGPU/AMDGPUSchedGPR.cpp llvm/lib/Target/AMDGPU/AMDGPUSchedGPR.h clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp clang/lib/Sema/SemaAMDGPU.cpp llvm/lib/Target/AMDGPU/AMDGPU.h llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp llvm/lib/Target/AMDGPU/SIRegisterInfo.h --diff_from_common_commit
``````````
:warning:
The reproduction instructions above might return results for more than one PR
in a stack if you are using a stacked PR workflow. You can limit the results by
changing `origin/main` to the base branch/commit you want to compare against.
:warning:
</details>
<details>
<summary>
View the diff from clang-format here.
</summary>
``````````diff
diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
index c27a8e8f2..ab57c9b11 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 556526394..95d5842d3 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 a6c20d61b..072725db5 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 85b4776db..8fbcafd21 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 0bc47414f..ad83a89d7 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;
``````````
</details>
https://github.com/llvm/llvm-project/pull/233773
More information about the Mlir-commits
mailing list