[clang] [llvm] [mlir] [AMDGPU] Add schedule bank hint for gfx1250 to get high perf (PR #227953)
via llvm-commits
llvm-commits at lists.llvm.org
Wed Sep 30 21:56:13 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clang-codegen
Author: jli (jli-melchior)
<details>
<summary>Changes</summary>
---
Patch is 34.69 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/227953.diff
23 Files Affected:
- (modified) clang/include/clang/Basic/BuiltinsAMDGPU.td (+1)
- (modified) clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp (+16)
- (modified) clang/lib/Sema/SemaAMDGPU.cpp (+19)
- (added) clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-schedule-bank.cl (+71)
- (modified) clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-param.cl (+9)
- (modified) clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250.cl (+1)
- (modified) llvm/include/llvm/IR/IntrinsicsAMDGPU.td (+12)
- (modified) llvm/lib/Target/AMDGPU/AMDGPU.h (+3)
- (modified) llvm/lib/Target/AMDGPU/AMDGPU.td (+2)
- (modified) llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp (+25)
- (modified) llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def (+1)
- (added) llvm/lib/Target/AMDGPU/AMDGPUScheduleBank.cpp (+190)
- (added) llvm/lib/Target/AMDGPU/AMDGPUScheduleBank.h (+23)
- (modified) llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp (+4)
- (modified) llvm/lib/Target/AMDGPU/CMakeLists.txt (+1)
- (modified) llvm/lib/Target/AMDGPU/SIInstructions.td (+16)
- (modified) llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp (+58)
- (modified) llvm/lib/Target/AMDGPU/SIRegisterInfo.h (+1-1)
- (added) llvm/test/CodeGen/AMDGPU/schedule-bank-hints.mir (+35)
- (added) llvm/test/CodeGen/AMDGPU/schedule-bank.ll (+32)
- (modified) mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td (+36)
- (modified) mlir/test/Dialect/LLVMIR/rocdl.mlir (+8)
- (modified) mlir/test/Target/LLVMIR/rocdl.mlir (+11)
``````````diff
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td
index 4fd604390d3ce3..b750be8a0e9ae6 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.td
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td
@@ -986,6 +986,7 @@ def __builtin_amdgcn_cvt_sr_f16_f32 : AMDGPUBuiltin<"_ExtVector<2, _Float16>(_Ex
// GFX1250+ only builtins.
//===----------------------------------------------------------------------===//
def __builtin_amdgcn_s_cluster_barrier : AMDGPUBuiltin<"void()", [], "gfx1250-insts">;
+def __builtin_amdgcn_schedule_bank : AMDGPUBuiltin<"int(int, _Constant int, _Constant bool)", [Const, CustomTypeChecking], "gfx1250-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 f3bf71dddc341f..1a1f01c9c45180 100644
--- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
+++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
@@ -2286,6 +2286,22 @@ 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_schedule_bank: {
+ llvm::Value *Val = EmitScalarExpr(E->getArg(0));
+ llvm::Value *Bank = EmitScalarExpr(E->getArg(1));
+ llvm::Value *EncodedBank;
+ if (E->getNumArgs() > 2) {
+ llvm::Value *Strict = EmitScalarExpr(E->getArg(2));
+ Strict = Builder.CreateZExt(Strict, Bank->getType());
+ EncodedBank = Builder.CreateOr(Bank, Builder.CreateShl(Strict, 2));
+ } else {
+ // Default: strict (bank | 4)
+ EncodedBank = Builder.CreateOr(Bank, Builder.getInt32(4));
+ }
+ llvm::Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_schedule_bank,
+ Val->getType());
+ return Builder.CreateCall(F, {Val, EncodedBank});
+ }
default:
return nullptr;
}
diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp
index ee6d989d2b107a..de86ecc749e32c 100644
--- a/clang/lib/Sema/SemaAMDGPU.cpp
+++ b/clang/lib/Sema/SemaAMDGPU.cpp
@@ -459,6 +459,25 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(const TargetInfo &TI,
/*High=*/0) ||
SemaRef.BuiltinConstantArgRange(TheCall, /*ArgNum=*/2, /*Low=*/0,
/*High=*/0);
+ case AMDGPU::BI__builtin_amdgcn_schedule_bank: {
+ 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);
+ return SemaRef.BuiltinConstantArgRange(TheCall, 1, 0, 3);
+ }
default:
return false;
}
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-schedule-bank.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-schedule-bank.cl
new file mode 100644
index 00000000000000..17fc1f9c37f666
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-schedule-bank.cl
@@ -0,0 +1,71 @@
+// 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;
+
+// CHECK-LABEL: @test_schedule_bank_i32
+// CHECK: call i32 @llvm.amdgcn.schedule.bank.i32(i32 %{{.*}}, i32 2)
+void test_schedule_bank_i32(global int *out, int x) {
+ *out = __builtin_amdgcn_schedule_bank(x, 2, false);
+}
+
+// CHECK-LABEL: @test_schedule_bank_f32
+// CHECK: call float @llvm.amdgcn.schedule.bank.f32(float %{{.*}}, i32 3)
+void test_schedule_bank_f32(global float *out, float x) {
+ *out = __builtin_amdgcn_schedule_bank(x, 3, false);
+}
+
+// CHECK-LABEL: @test_schedule_bank_i64
+// CHECK: call i64 @llvm.amdgcn.schedule.bank.i64(i64 %{{.*}}, i32 0)
+void test_schedule_bank_i64(global long *out, long x) {
+ *out = __builtin_amdgcn_schedule_bank(x, 0, false);
+}
+
+// CHECK-LABEL: @test_schedule_bank_f64_strict
+// CHECK: call double @llvm.amdgcn.schedule.bank.f64(double %{{.*}}, i32 7)
+void test_schedule_bank_f64_strict(global double *out, double x) {
+ *out = __builtin_amdgcn_schedule_bank(x, 3, true);
+}
+
+// CHECK-LABEL: @test_schedule_bank_v2i32
+// CHECK: call <2 x i32> @llvm.amdgcn.schedule.bank.v2i32(<2 x i32> %{{.*}}, i32 1)
+void test_schedule_bank_v2i32(global int2 *out, int2 x) {
+ *out = __builtin_amdgcn_schedule_bank(x, 1, false);
+}
+
+// CHECK-LABEL: @test_schedule_bank_v4i32_strict
+// CHECK: call <4 x i32> @llvm.amdgcn.schedule.bank.v4i32(<4 x i32> %{{.*}}, i32 5)
+void test_schedule_bank_v4i32_strict(global int4 *out, int4 x) {
+ *out = __builtin_amdgcn_schedule_bank(x, 1, true);
+}
+
+// CHECK-LABEL: @test_schedule_bank_v8i32_strict
+// CHECK: call <8 x i32> @llvm.amdgcn.schedule.bank.v8i32(<8 x i32> %{{.*}}, i32 6)
+void test_schedule_bank_v8i32_strict(global int8 *out, int8 x) {
+ *out = __builtin_amdgcn_schedule_bank(x, 2, true);
+}
+
+// CHECK-LABEL: @test_schedule_bank_v4f32_strict
+// CHECK: call <4 x float> @llvm.amdgcn.schedule.bank.v4f32(<4 x float> %{{.*}}, i32 4)
+void test_schedule_bank_v4f32_strict(global float4 *out, float4 x) {
+ *out = __builtin_amdgcn_schedule_bank(x, 0, true);
+}
+
+// CHECK-LABEL: @test_schedule_bank_default_strict
+// CHECK: call i32 @llvm.amdgcn.schedule.bank.i32(i32 %{{.*}}, i32 6)
+void test_schedule_bank_default_strict(global int *out, int x) {
+ // 2 args: strict=true by default, bank 2 → encoded as 6
+ *out = __builtin_amdgcn_schedule_bank(x, 2);
+}
+
+// CHECK-LABEL: @test_schedule_bank_explicit_soft
+// CHECK: call i32 @llvm.amdgcn.schedule.bank.i32(i32 %{{.*}}, i32 2)
+void test_schedule_bank_explicit_soft(global int *out, int x) {
+ // 3 args: explicit soft, bank 2 → encoded as 2
+ *out = __builtin_amdgcn_schedule_bank(x, 2, false);
+}
diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-param.cl b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-param.cl
index 8c60b567ddc21c..37bb036957459f 100644
--- a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-param.cl
+++ b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-param.cl
@@ -216,3 +216,12 @@ void test_pk_add_min_max(global short2 *out, global ushort2 *uout, short2 a, sho
*out = __builtin_amdgcn_pk_add_min_i16(a, b, c, clamp); // expected-error {{'__builtin_amdgcn_pk_add_min_i16' must be a constant integer}}
*uout = __builtin_amdgcn_pk_add_min_u16(ua, ub, uc, clamp); // expected-error {{'__builtin_amdgcn_pk_add_min_u16' must be a constant integer}}
}
+
+void test_schedule_bank_non_const(int x, int b) {
+ __builtin_amdgcn_schedule_bank(x, b, false); // expected-error {{'__builtin_amdgcn_schedule_bank' must be a constant integer}}
+}
+
+void test_schedule_bank_out_of_range(int x) {
+ __builtin_amdgcn_schedule_bank(x, 4, false); // expected-error {{argument value 4 is outside the valid range [0, 3]}}
+ __builtin_amdgcn_schedule_bank(x, -1, true); // expected-error {{argument value -1 is outside the valid range [0, 3]}}
+}
diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250.cl b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250.cl
index 87110d4d977dba..01bb27927895b4 100644
--- a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250.cl
+++ b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250.cl
@@ -20,4 +20,5 @@ void test(global int* out, global short2 *s2out, global ushort2 *us2out,
*us2out = __builtin_amdgcn_pk_add_max_u16(us2a, us2b, us2c, true); // expected-error {{'__builtin_amdgcn_pk_add_max_u16' needs target feature pk-add-min-max-insts}}
*s2out = __builtin_amdgcn_pk_add_min_i16(s2a, s2b, s2c, false); // expected-error {{'__builtin_amdgcn_pk_add_min_i16' needs target feature pk-add-min-max-insts}}
*us2out = __builtin_amdgcn_pk_add_min_u16(us2a, us2b, us2c, true); // expected-error {{'__builtin_amdgcn_pk_add_min_u16' needs target feature pk-add-min-max-insts}}
+ *out = __builtin_amdgcn_schedule_bank(a, 2, false); // expected-error {{'__builtin_amdgcn_schedule_bank' needs target feature gfx1250-insts}}
}
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index e7829eb29ba34f..13365c13e4618c 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -4380,4 +4380,16 @@ 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 bank on
+// targets with 1024 addressable VGPRs (gfx1250). The bank operand is a
+// compile-time constant in [0, 3]:
+// bank 0 = v0-v255, bank 1 = v256-v511, bank 2 = v512-v767, bank 3 = v768-v1023.
+// This is an identity function at the IR level; the bank argument is consumed by
+// the pre-RA AMDGPUScheduleBank pass which attaches a register allocation hint.
+// The hint is advisory: the allocator honors it only when the bank has free
+// registers, and never spills to satisfy it.
+def int_amdgcn_schedule_bank : DefaultAttrsIntrinsic<
+ [llvm_any_ty], [LLVMMatchType<0>, llvm_i32_ty],
+ [IntrNoMem, IntrSpeculatable, IntrWillReturn, ImmArg<ArgIndex<1>>]>;
}
diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.h b/llvm/lib/Target/AMDGPU/AMDGPU.h
index c87e9adeaa7d10..391fdb6ce4e140 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPU.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPU.h
@@ -239,6 +239,9 @@ struct AMDGPULowerIntrinsicsPass
void initializeAMDGPUPrepareAGPRAllocLegacyPass(PassRegistry &);
extern char &AMDGPUPrepareAGPRAllocLegacyID;
+void initializeAMDGPUScheduleBankPass(PassRegistry &);
+extern char &AMDGPUScheduleBankID;
+
void initializeAMDGPUReserveWWMRegsLegacyPass(PassRegistry &);
extern char &AMDGPUReserveWWMRegsLegacyID;
diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.td b/llvm/lib/Target/AMDGPU/AMDGPU.td
index c3b4d53a7effa1..a418786ee2f8e2 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPU.td
+++ b/llvm/lib/Target/AMDGPU/AMDGPU.td
@@ -2548,6 +2548,7 @@ def FeatureISAVersion12_50_STRICT : FeatureSet<
FeatureNoSleepForever,
FeatureSlowMaxMinMulI64Insts,
FeatureSupportsXNACK,
+ FeatureXNACKOnOffModes,
])>;
def FeatureISAVersion12_50 : FeatureSet<
@@ -2581,6 +2582,7 @@ def FeatureISAVersion12_50 : FeatureSet<
FeatureNoSleepForever,
FeatureSlowMaxMinMulI64Insts,
FeatureSupportsXNACK,
+ FeatureXNACKOnOffModes,
])>;
def FeatureISAVersion12_51 : FeatureSet<
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp b/llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp
index ff58f560314ab7..199277c113724c 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp
@@ -3317,6 +3317,31 @@ void AMDGPUDAGToDAGISel::SelectINTRINSIC_WO_CHAIN(SDNode *N) {
case Intrinsic::amdgcn_interp_p1_f16:
SelectInterpP1F16(N);
return;
+ case Intrinsic::amdgcn_schedule_bank: {
+ SDValue Src = N->getOperand(1);
+ auto *BankC = dyn_cast<ConstantSDNode>(N->getOperand(2));
+ // Bank must be a constant; fall back (and error) otherwise.
+ if (!BankC) {
+ SelectCode(N);
+ return;
+ }
+ unsigned PseudoOpc;
+ switch (N->getValueType(0).getSizeInBits()) {
+ case 32: PseudoOpc = AMDGPU::V_SCHEDULE_BANK_B32; break;
+ case 64: PseudoOpc = AMDGPU::V_SCHEDULE_BANK_B64; break;
+ case 128: PseudoOpc = AMDGPU::V_SCHEDULE_BANK_B128; break;
+ case 256: PseudoOpc = AMDGPU::V_SCHEDULE_BANK_B256; break;
+ default:
+ // Unsupported width: drop the hint, forward the value unchanged.
+ ReplaceUses(SDValue(N, 0), Src);
+ CurDAG->RemoveDeadNode(N);
+ return;
+ }
+ SDValue BankImm = CurDAG->getTargetConstant(BankC->getZExtValue(), SDLoc(N),
+ MVT::i32);
+ CurDAG->SelectNodeTo(N, PseudoOpc, N->getVTList(), {Src, BankImm});
+ return;
+ }
case Intrinsic::amdgcn_permlane16_swap:
case Intrinsic::amdgcn_permlane32_swap: {
if ((IntrID == Intrinsic::amdgcn_permlane16_swap &&
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def b/llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def
index 29e9046e4f8caf..bda6e66153e48e 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def
+++ b/llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def
@@ -137,6 +137,7 @@ MACHINE_FUNCTION_PASS("amdgpu-reg-bank-legalize", AMDGPURegBankLegalizePass())
MACHINE_FUNCTION_PASS("amdgpu-regbank-combiner", AMDGPURegBankCombinerPass())
MACHINE_FUNCTION_PASS("amdgpu-preload-kern-arg-prolog", AMDGPUPreloadKernArgPrologPass())
MACHINE_FUNCTION_PASS("amdgpu-prepare-agpr-alloc", AMDGPUPrepareAGPRAllocPass())
+MACHINE_FUNCTION_PASS("amdgpu-schedule-bank", AMDGPUScheduleBankPass())
MACHINE_FUNCTION_PASS("amdgpu-nsa-reassign", GCNNSAReassignPass())
MACHINE_FUNCTION_PASS("amdgpu-wait-sgpr-hazards", AMDGPUWaitSGPRHazardsPass())
MACHINE_FUNCTION_PASS("gcn-create-vopd", GCNCreateVOPDPass())
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUScheduleBank.cpp b/llvm/lib/Target/AMDGPU/AMDGPUScheduleBank.cpp
new file mode 100644
index 00000000000000..fb2cc1651ea499
--- /dev/null
+++ b/llvm/lib/Target/AMDGPU/AMDGPUScheduleBank.cpp
@@ -0,0 +1,190 @@
+//===-- AMDGPUScheduleBank.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_BANK_B* pseudos produced from llvm.amdgcn.schedule.bank.
+/// For each pseudo this pass:
+/// 1. reads the requested bank (0-3) from the immediate operand,
+/// 2. attaches an AMDGPURI::BankHint register allocation hint to the
+/// destination (and source) vreg so the allocator prefers that 256-register
+/// bank,
+/// 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: SIRegisterInfo::getRegAllocationHints only reorders the
+/// allocation candidates, so a full bank falls back to the default order and the
+/// pass never forces a spill.
+//
+//===----------------------------------------------------------------------===//
+
+#include "AMDGPUScheduleBank.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-schedule-bank"
+
+static bool isScheduleBankPseudo(unsigned Opcode) {
+ switch (Opcode) {
+ case AMDGPU::V_SCHEDULE_BANK_B32:
+ case AMDGPU::V_SCHEDULE_BANK_B64:
+ case AMDGPU::V_SCHEDULE_BANK_B128:
+ case AMDGPU::V_SCHEDULE_BANK_B256:
+ return true;
+ default:
+ return false;
+ }
+}
+
+/// Walk the SSA use chain of \p Reg and propagate the bank hint through
+/// value-preserving pseudos so the preference reaches the vregs that survive
+/// coalescing.
+static void propagateBankHint(MachineRegisterInfo &MRI, Register Reg,
+ unsigned Bank, 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 bank hint.
+ std::pair<unsigned, Register> Existing = MRI.getRegAllocationHint(DefReg);
+ if (Existing.first != 0 &&
+ (Existing.first != HintKind || Existing.second != Bank))
+ continue;
+
+ MRI.setRegAllocationHint(DefReg, HintKind, Bank);
+ propagateBankHint(MRI, DefReg, Bank, HintKind, Visited);
+ }
+}
+
+// Shared implementation used by both the legacy and new-PM passes.
+static bool runScheduleBank(MachineFunction &MF) {
+ const GCNSubtarget &ST = MF.getSubtarget<GCNSubtarget>();
+ if (!ST.has1024AddressableVGPRs())
+ return false;
+
+ MachineRegisterInfo &MRI = MF.getRegInfo();
+ const SIInstrInfo *TII = ST.getInstrInfo();
+ bool Changed = false;
+
+ SmallVector<MachineInstr *, 16> ToErase;
+
+ for (MachineBasicBlock &MBB : MF) {
+ for (MachineInstr &MI : MBB) {
+ if (!isScheduleBankPseudo(MI.getOpcode()))
+ continue;
+
+ Register DstReg = MI.getOperand(0).getReg();
+ const MachineOperand &SrcMO = MI.getOperand(1);
+ Register SrcReg = SrcMO.getReg();
+ unsigned EncodedBank = MI.getOperand(2).getImm();
+ bool Strict = EncodedBank & 0x4;
+ unsigned Bank = EncodedBank & 0x3;
+ unsigned HintKind =
+ Strict ? AMDGPURI::StrictBankHint : AMDGPURI::BankHint;
+
+ // Bits 0..1 select bank; bit 2 requests strict allocation.
+ if (EncodedBank > 7) {
+ Bank = 0;
+ Strict = false;
+ }
+
+ LLVM_DEBUG(dbgs() << " schedule.bank: " << printReg(DstReg) << " <- "
+ << printReg(SrcReg) << " bank " << Bank
+ << " strict " << Strict << '\n');
+
+ // Hint both the destination and the source so the preference survives
+ // whichever side coalescing keeps.
+ if (DstReg.isVirtual() &&
+ (Strict || MRI.getRegAllocationHint(DstReg).first !=
+ AMDGPURI::StrictBankHint))
+ MRI.setRegAllocationHint(DstReg, HintKind, Bank);
+ if (SrcReg.isVirtual() &&
+ (Strict || MRI.getRegAllocationHint(SrcReg).first !=
+ AMDGPURI::StrictBankHint))
+ MRI.setRegAllocationHint(SrcReg, HintKind, Bank);
+
+ SmallDenseSet<unsigned, 32> Visited;
+ if (DstReg.isVirtual())
+ propagateBankHint(MRI, DstReg, Bank, 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 AMDGPUScheduleBank : public MachineFunctionPass {
+public:
+ static char ID;
+
+ AMDGPUScheduleBank() : MachineFunctionPass(ID) {}
+
+ bool runOnMachineFunction(MachineFunction &MF) override {
+ return runScheduleBank(MF);
+ }
+
+ StringRef getPassName() const override { return "AMDGPU Schedule Bank"; }
+
+ void getAnalysisUsage(AnalysisUsage &AU) const override {
+ AU.setPreservesCFG();
+ MachineFunctionPass::getAnalysisUsage(AU);
+ }
+};
+
+} // end anonymous namespace
+
+INITIALIZE_PASS(AMDGPUScheduleBank, DEBUG_TYPE, "AMDGPU Schedule Bank",
+ false, false)
+
+char AMDGPUScheduleBank::ID = 0;
+
+char &llvm::AMDGPUScheduleBankID = A...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/227953
More information about the llvm-commits
mailing list