[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:55:20 PDT 2026
https://github.com/jli-melchior created https://github.com/llvm/llvm-project/pull/227953
None
>From 355ed0ae73eb196fc1a04f0c1db02143bb491ef4 Mon Sep 17 00:00:00 2001
From: jli-melchior <Jun.Li at amd.com>
Date: Wed, 30 Sep 2026 02:53:55 +0000
Subject: [PATCH 1/3] [AMDGPU] Add llvm.amdgcn.schedule.bank intrinsic and
__builtin_amdgcn_schedule_bank
Add the schedule.bank VGPR bank hint infrastructure for gfx1250:
LLVM backend:
- Define the llvm.amdgcn.schedule.bank intrinsic in IntrinsicsAMDGPU.td
- Add V_SCHEDULE_BANK_B32/64/128/256 pseudo instructions for ISel
- Add AMDGPUScheduleBank pass (legacy + new-PM) that lowers the pseudos
to COPYs and attaches BankHint/StrictBankHint RA hints
- Extend SIRegisterInfo::getRegAllocationHints with BankHint and
StrictBankHint support for steering allocation into 256-register banks
Clang frontend:
- Add __builtin_amdgcn_schedule_bank(value, bank [, strict]) builtin
- Sema validation for bank range and type constraints
- CodeGen lowering that encodes bank+strict into a single i32
Co-Authored-By: Claude <noreply at anthropic.com>
---
clang/include/clang/Basic/BuiltinsAMDGPU.td | 1 +
clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp | 16 ++
clang/lib/Sema/SemaAMDGPU.cpp | 19 ++
.../builtins-amdgcn-gfx1250-schedule-bank.cl | 71 +++++++
.../builtins-amdgcn-error-gfx1250-param.cl | 9 +
.../builtins-amdgcn-error-gfx1250.cl | 1 +
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 12 ++
llvm/lib/Target/AMDGPU/AMDGPU.h | 3 +
llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp | 25 +++
llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def | 1 +
llvm/lib/Target/AMDGPU/AMDGPUScheduleBank.cpp | 190 ++++++++++++++++++
llvm/lib/Target/AMDGPU/AMDGPUScheduleBank.h | 23 +++
.../lib/Target/AMDGPU/AMDGPUTargetMachine.cpp | 4 +
llvm/lib/Target/AMDGPU/CMakeLists.txt | 1 +
llvm/lib/Target/AMDGPU/SIInstructions.td | 16 ++
llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp | 60 ++++++
llvm/lib/Target/AMDGPU/SIRegisterInfo.h | 2 +-
.../CodeGen/AMDGPU/schedule-bank-hints.mir | 35 ++++
llvm/test/CodeGen/AMDGPU/schedule-bank.ll | 32 +++
19 files changed, 520 insertions(+), 1 deletion(-)
create mode 100644 clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-schedule-bank.cl
create mode 100644 llvm/lib/Target/AMDGPU/AMDGPUScheduleBank.cpp
create mode 100644 llvm/lib/Target/AMDGPU/AMDGPUScheduleBank.h
create mode 100644 llvm/test/CodeGen/AMDGPU/schedule-bank-hints.mir
create mode 100644 llvm/test/CodeGen/AMDGPU/schedule-bank.ll
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td
index 4fd604390d3ce..b750be8a0e9ae 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 f3bf71dddc341..1a1f01c9c4518 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 ee6d989d2b107..de86ecc749e32 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 0000000000000..17fc1f9c37f66
--- /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 8c60b567ddc21..37bb036957459 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 87110d4d977db..01bb27927895b 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 e7829eb29ba34..13365c13e4618 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 c87e9adeaa7d1..391fdb6ce4e14 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/AMDGPUISelDAGToDAG.cpp b/llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp
index ff58f560314ab..199277c113724 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 29e9046e4f8ca..bda6e66153e48 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 0000000000000..fb2cc1651ea49
--- /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 = AMDGPUScheduleBank::ID;
+
+PreservedAnalyses
+AMDGPUScheduleBankPass::run(MachineFunction &MF,
+ MachineFunctionAnalysisManager &MFAM) {
+ if (!runScheduleBank(MF))
+ return PreservedAnalyses::all();
+
+ return getMachineFunctionPassPreservedAnalyses().preserveSet<CFGAnalyses>();
+}
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUScheduleBank.h b/llvm/lib/Target/AMDGPU/AMDGPUScheduleBank.h
new file mode 100644
index 0000000000000..ba8afdc9256f9
--- /dev/null
+++ b/llvm/lib/Target/AMDGPU/AMDGPUScheduleBank.h
@@ -0,0 +1,23 @@
+//===- AMDGPUScheduleBank.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_AMDGPUSCHEDULEBANK_H
+#define LLVM_LIB_TARGET_AMDGPU_AMDGPUSCHEDULEBANK_H
+
+#include "llvm/CodeGen/MachinePassManager.h"
+
+namespace llvm {
+class AMDGPUScheduleBankPass
+ : public PassInfoMixin<AMDGPUScheduleBankPass> {
+public:
+ PreservedAnalyses run(MachineFunction &MF,
+ MachineFunctionAnalysisManager &MFAM);
+};
+} // namespace llvm
+
+#endif // LLVM_LIB_TARGET_AMDGPU_AMDGPUSCHEDULEBANK_H
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp b/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
index 946d25400e17d..cc060ef41dfc0 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
@@ -32,6 +32,7 @@
#include "AMDGPUPerfHintAnalysis.h"
#include "AMDGPUPreloadKernArgProlog.h"
#include "AMDGPUPrepareAGPRAlloc.h"
+#include "AMDGPUScheduleBank.h"
#include "AMDGPURemoveIncompatibleFunctions.h"
#include "AMDGPUReserveWWMRegs.h"
#include "AMDGPUResourceUsageAnalysis.h"
@@ -695,6 +696,7 @@ extern "C" LLVM_ABI LLVM_EXTERNAL_VISIBILITY void LLVMInitializeAMDGPUTarget() {
initializeAMDGPUAsmPrinterPass(*PR);
initializeAMDGPUDAGToDAGISelLegacyPass(*PR);
initializeAMDGPUPrepareAGPRAllocLegacyPass(*PR);
+ initializeAMDGPUScheduleBankPass(*PR);
initializeGCNDPPCombineLegacyPass(*PR);
initializeSILowerI1CopiesLegacyPass(*PR);
initializeAMDGPUGlobalISelDivergenceLoweringLegacyPass(*PR);
@@ -1844,6 +1846,7 @@ void GCNPassConfig::addFastRegAlloc() {
}
void GCNPassConfig::addPreRegAlloc() {
+ addPass(&AMDGPUScheduleBankID);
if (getOptLevel() != CodeGenOptLevel::None)
addPass(&AMDGPUPrepareAGPRAllocLegacyID);
if (getOptLevel() >= CodeGenOptLevel::Default && EnableMachinePipeliner)
@@ -2693,6 +2696,7 @@ Error AMDGPUCodeGenPassBuilder::addOptimizedRegAlloc(PassManagerWrapper &PMW) {
}
void AMDGPUCodeGenPassBuilder::addPreRegAlloc(PassManagerWrapper &PMW) {
+ addMachineFunctionPass(AMDGPUScheduleBankPass(), 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 b7e679a69a80d..71c8611e7a726 100644
--- a/llvm/lib/Target/AMDGPU/CMakeLists.txt
+++ b/llvm/lib/Target/AMDGPU/CMakeLists.txt
@@ -80,6 +80,7 @@ add_llvm_target(AMDGPUCodeGen
AMDGPULowerKernelAttributes.cpp
AMDGPULowerModuleLDSPass.cpp
AMDGPUPrepareAGPRAlloc.cpp
+ AMDGPUScheduleBank.cpp
AMDGPULowerExecSync.cpp
AMDGPUSwLowerLDS.cpp
AMDGPUMachineFunctionInfo.cpp
diff --git a/llvm/lib/Target/AMDGPU/SIInstructions.td b/llvm/lib/Target/AMDGPU/SIInstructions.td
index eaece88e98525..51abd76751387 100644
--- a/llvm/lib/Target/AMDGPU/SIInstructions.td
+++ b/llvm/lib/Target/AMDGPU/SIInstructions.td
@@ -129,6 +129,22 @@ def V_CNDMASK_B64_PSEUDO : VOP3Common <(outs VReg_64:$vdst),
let usesCustomInserter = 1;
}
+// VGPR bank hint pseudos carrying a target bank number in [0, 3]. These are
+// identity copies produced by lowering llvm.amdgcn.schedule.bank. A pre-RA pass
+// (AMDGPUScheduleBank) reads the bank operand, sets a register allocation hint
+// on the destination vreg, and replaces the pseudo with a COPY.
+let isAsCheapAsAMove = 1, isReMaterializable = 1, hasSideEffects = 0,
+ mayLoad = 0, mayStore = 0, Size = 0 in {
+ def V_SCHEDULE_BANK_B32 : VPseudoInstSI<(outs VGPR_32:$vdst),
+ (ins VGPR_32:$src, i32imm:$bank)>;
+ def V_SCHEDULE_BANK_B64 : VPseudoInstSI<(outs VReg_64:$vdst),
+ (ins VReg_64:$src, i32imm:$bank)>;
+ def V_SCHEDULE_BANK_B128 : VPseudoInstSI<(outs VReg_128:$vdst),
+ (ins VReg_128:$src, i32imm:$bank)>;
+ def V_SCHEDULE_BANK_B256 : VPseudoInstSI<(outs VReg_256:$vdst),
+ (ins VReg_256:$src, i32imm:$bank)>;
+}
+
// 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),
diff --git a/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp b/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp
index 7cad384c98a8d..49b9354eed8b3 100644
--- a/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp
@@ -49,6 +49,11 @@ static cl::opt<unsigned> StressSGPRLimit(
"amdgpu-stress-sgpr", cl::Hidden, cl::init(0),
cl::desc("Limit SGPRs to N registers by reserving the rest"));
+static cl::opt<bool> StrictVGPRBankHints(
+ "amdgpu-strict-vgpr-bank-hints",
+ cl::desc("Restrict schedule.bank values to the requested VGPR bank"),
+ cl::Hidden, cl::init(false));
+
std::array<std::vector<int16_t>, 32> SIRegisterInfo::RegSplitParts;
std::array<std::array<uint16_t, 32>, 9> SIRegisterInfo::SubRegFromChannelTable;
@@ -4292,6 +4297,61 @@ bool SIRegisterInfo::getRegAllocationHints(Register VirtReg,
}
return false;
}
+ case AMDGPURI::SameBank: {
+ const auto *HintInfo = MRI.getRegAllocationHints(VirtReg);
+ if (!HintInfo || !VRM)
+ return TargetRegisterInfo::getRegAllocationHints(VirtReg, Order, Hints,
+ MF, VRM);
+
+ // Determine the target bank from any already-assigned hint register.
+ int TargetBank = -1;
+ for (Register HintReg : HintInfo->second) {
+ MCPhysReg PhysHint = 0;
+ if (HintReg.isPhysical())
+ PhysHint = HintReg;
+ else if (VRM->hasPhys(HintReg))
+ PhysHint = VRM->getPhys(HintReg);
+ if (PhysHint) {
+ TargetBank = static_cast<int>(getHWRegIndex(PhysHint) >> 8);
+ break;
+ }
+ }
+
+ if (TargetBank < 0)
+ return TargetRegisterInfo::getRegAllocationHints(VirtReg, Order, Hints,
+ MF, VRM);
+
+ for (MCPhysReg PhysReg : Order) {
+ if (MRI.isReserved(PhysReg))
+ continue;
+ unsigned RegBank = getHWRegIndex(PhysReg) >> 8;
+ if (static_cast<int>(RegBank) == TargetBank)
+ Hints.push_back(PhysReg);
+ }
+ return false;
+ }
+ case AMDGPURI::BankHint:
+ case AMDGPURI::StrictBankHint: {
+ // Absolute bank preference from llvm.amdgcn.schedule.bank. The bank number
+ // (0-3) is carried in Hint.second. Prefer physregs whose HW index falls in
+ // the requested 256-register bank. Advisory only: we return false so the
+ // default order still applies when the bank is full.
+ unsigned Bank = Hint.second;
+ if (Bank > 3)
+ return false;
+ unsigned BankStart = Bank * 256;
+ unsigned BankEnd = BankStart + 256;
+ for (MCPhysReg PhysReg : Order) {
+ if (MRI.isReserved(PhysReg))
+ continue;
+ unsigned HWIdx = getHWRegIndex(PhysReg);
+ if (HWIdx >= BankStart && HWIdx < BankEnd)
+ Hints.push_back(PhysReg);
+ }
+ LLVM_DEBUG(dbgs() << "BankHint: " << printReg(VirtReg, this) << " bank="
+ << Bank << " -> " << Hints.size() << " candidates\n");
+ return StrictVGPRBankHints || Hint.first == AMDGPURI::StrictBankHint;
+ }
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 1eb9a88b819ed..c2cb5d897a10d 100644
--- a/llvm/lib/Target/AMDGPU/SIRegisterInfo.h
+++ b/llvm/lib/Target/AMDGPU/SIRegisterInfo.h
@@ -33,7 +33,7 @@ struct SGPRSpillBuilder;
/// Register allocation hint types. Helps eliminate unneeded COPY with True16
namespace AMDGPURI {
-enum { Size16 = 1, Size32 = 2 };
+enum { Size16 = 1, Size32 = 2, SameBank = 3, BankHint = 4, StrictBankHint = 5 };
} // end namespace AMDGPURI
diff --git a/llvm/test/CodeGen/AMDGPU/schedule-bank-hints.mir b/llvm/test/CodeGen/AMDGPU/schedule-bank-hints.mir
new file mode 100644
index 0000000000000..2ba66444e22df
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/schedule-bank-hints.mir
@@ -0,0 +1,35 @@
+# RUN: llc -mtriple=amdgcn -mcpu=gfx1250 -run-pass=amdgpu-schedule-bank -o - %s | FileCheck %s
+
+# The AMDGPUScheduleBank pass attaches a bank register-allocation hint (not
+# printed in MIR) and lowers each V_SCHEDULE_BANK_B* pseudo to a plain COPY.
+
+---
+# CHECK-LABEL: name: lower_schedule_bank_b32
+# CHECK: %1:vgpr_32 = COPY %0
+# CHECK-NOT: V_SCHEDULE_BANK_B32
+name: lower_schedule_bank_b32
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0
+ %0:vgpr_32 = COPY $vgpr0
+ %1:vgpr_32 = V_SCHEDULE_BANK_B32 %0, 2, 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_schedule_bank_b64
+# CHECK: %1:vreg_64_align2 = COPY %0
+# CHECK-NOT: V_SCHEDULE_BANK_B64
+name: lower_schedule_bank_b64
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0_vgpr1
+ %0:vreg_64_align2 = COPY $vgpr0_vgpr1
+ %1:vreg_64_align2 = V_SCHEDULE_BANK_B64 %0, 3, implicit $exec
+ $vgpr0_vgpr1 = COPY %1
+ SI_RETURN_TO_EPILOG $vgpr0_vgpr1
+...
diff --git a/llvm/test/CodeGen/AMDGPU/schedule-bank.ll b/llvm/test/CodeGen/AMDGPU/schedule-bank.ll
new file mode 100644
index 0000000000000..f40681b1c6f2b
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/schedule-bank.ll
@@ -0,0 +1,32 @@
+; RUN: llc -mtriple=amdgcn -mcpu=gfx1250 < %s | FileCheck -check-prefix=GCN %s
+
+; Verify that llvm.amdgcn.schedule.bank steers register allocation into the
+; requested 256-register bank. The kernel is pinned to occupancy 1
+; (waves-per-eu=1,1 with a single-wave workgroup) so the whole 1024-VGPR file is
+; addressable and the high banks are actually available to the allocator.
+
+declare i32 @llvm.amdgcn.schedule.bank.i32(i32, i32)
+
+; A bank-2 hint places the value in v512-v767 and forces an s_set_vgpr_msb that
+; latches bank 2 (src1=2) around its use.
+; GCN-LABEL: {{^}}bank2:
+; GCN: v_mov_b32_e32 v0 /*v512*/, s2
+; GCN: s_set_vgpr_msb 0x8008
+define amdgpu_kernel void @bank2(ptr addrspace(1) %out, i32 %x) #0 {
+ %h = call i32 @llvm.amdgcn.schedule.bank.i32(i32 %x, i32 2)
+ %y = add i32 %h, 1
+ store i32 %y, ptr addrspace(1) %out
+ ret void
+}
+
+; Without the hint the value stays in bank 0 and no bank switch is emitted.
+; GCN-LABEL: {{^}}nohint:
+; GCN-NOT: s_set_vgpr_msb
+; 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" }
>From 80b9dbfb3ba6c5b825b81dabd35719f7294adc97 Mon Sep 17 00:00:00 2001
From: jli-melchior <Jun.Li at amd.com>
Date: Wed, 30 Sep 2026 03:35:47 +0000
Subject: [PATCH 2/3] [MLIR][ROCDL] Add rocdl.schedule.bank op for VGPR bank
hints
Add ROCDL_ScheduleBankOp that maps to the llvm.amdgcn.schedule.bank
intrinsic, enabling MLIR-based codegen pipelines to emit VGPR bank
scheduling hints for gfx1250.
The op defaults to strict mode (allocator must obey). An optional
`soft` keyword makes it advisory, consistent with the HIP builtin API:
%0 = rocdl.schedule.bank %val, 2 : f32 // strict bank 2
%1 = rocdl.schedule.bank %val, 1 soft : f32 // soft bank 1
Co-Authored-By: Claude <noreply at anthropic.com>
---
llvm/lib/Target/AMDGPU/AMDGPUScheduleBank.h | 2 +-
llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp | 2 --
mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td | 36 ++++++++++++++++++++
mlir/test/Dialect/LLVMIR/rocdl.mlir | 8 +++++
mlir/test/Target/LLVMIR/rocdl.mlir | 11 ++++++
5 files changed, 56 insertions(+), 3 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUScheduleBank.h b/llvm/lib/Target/AMDGPU/AMDGPUScheduleBank.h
index ba8afdc9256f9..79d6bab6e0556 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUScheduleBank.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPUScheduleBank.h
@@ -13,7 +13,7 @@
namespace llvm {
class AMDGPUScheduleBankPass
- : public PassInfoMixin<AMDGPUScheduleBankPass> {
+ : public OptionalPassInfoMixin<AMDGPUScheduleBankPass> {
public:
PreservedAnalyses run(MachineFunction &MF,
MachineFunctionAnalysisManager &MFAM);
diff --git a/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp b/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp
index 49b9354eed8b3..51ef57a1072d9 100644
--- a/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp
@@ -4348,8 +4348,6 @@ bool SIRegisterInfo::getRegAllocationHints(Register VirtReg,
if (HWIdx >= BankStart && HWIdx < BankEnd)
Hints.push_back(PhysReg);
}
- LLVM_DEBUG(dbgs() << "BankHint: " << printReg(VirtReg, this) << " bank="
- << Bank << " -> " << Hints.size() << " candidates\n");
return StrictVGPRBankHints || Hint.first == AMDGPURI::StrictBankHint;
}
default:
diff --git a/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td b/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td
index 716fd92e810ac..4b2194f1f7a22 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td
@@ -813,6 +813,42 @@ def ROCDL_IglpOpt : ROCDL_ConcreteNonMemIntrOp<"iglp.opt", [], 0, [0], ["variant
}];
}
+def ROCDL_ScheduleBankOp : ROCDL_IntrOp<"schedule.bank", [], [0],
+ [AllTypesMatch<["res", "value"]>], 1, 0, 0, 0, [], []>,
+ Arguments<(ins LLVM_Type:$value, I32Attr:$bank,
+ OptionalAttr<UnitAttr>:$soft)> {
+ let results = (outs LLVM_Type:$res);
+ let assemblyFormat = [{
+ $value `,` $bank (`soft` $soft^)? attr-dict `:` type($value)
+ }];
+ string llvmBuilder = [{
+ auto schedOp = cast<ROCDL::ScheduleBankOp>(opInst);
+ auto *value = moduleTranslation.lookupValue(schedOp.getValue());
+ int32_t bank = schedOp.getBank();
+ int32_t encodedBank = schedOp.getSoftAttr() ? bank : (bank | 4);
+ auto *fn = llvm::Intrinsic::getOrInsertDeclaration(
+ moduleTranslation.getLLVMModule(),
+ llvm::Intrinsic::amdgcn_schedule_bank,
+ {value->getType()});
+ auto *inst = builder.CreateCall(fn,
+ {value, builder.getInt32(encodedBank)});
+ moduleTranslation.mapValue(opInst.getResult(0), inst);
+ }];
+ let description = [{
+ VGPR bank scheduling hint (gfx1250). Identity at runtime; guides the
+ register allocator to place `value` in the requested 256-register bank.
+
+ Bank 0-3. Default is strict (allocator must obey). Pass `soft` for an
+ advisory hint that the allocator may ignore.
+
+ Example:
+ ```mlir
+ %0 = rocdl.schedule.bank %val, 2 : f32 // strict bank 2 (default)
+ %1 = rocdl.schedule.bank %val, 1 soft : f32 // soft bank 1
+ ```
+ }];
+}
+
//===---------------------------------------------------------------------===//
// Xdlops intrinsics
diff --git a/mlir/test/Dialect/LLVMIR/rocdl.mlir b/mlir/test/Dialect/LLVMIR/rocdl.mlir
index 1b5b641936881..ae77810a369e4 100644
--- a/mlir/test/Dialect/LLVMIR/rocdl.mlir
+++ b/mlir/test/Dialect/LLVMIR/rocdl.mlir
@@ -168,6 +168,14 @@ func.func @rocdl_iglp_opt() {
llvm.return
}
+func.func @rocdl_schedule_bank(%val : f32, %ival : i32) -> f32 {
+ // CHECK: rocdl.schedule.bank %{{.*}}, 2 : f32
+ %0 = rocdl.schedule.bank %val, 2 : f32
+ // CHECK: rocdl.schedule.bank %{{.*}}, 1 soft : i32
+ %1 = rocdl.schedule.bank %ival, 1 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 b4a94cef770fa..ffa584532e6d2 100644
--- a/mlir/test/Target/LLVMIR/rocdl.mlir
+++ b/mlir/test/Target/LLVMIR/rocdl.mlir
@@ -444,6 +444,17 @@ llvm.func @rocdl.iglp.opt() {
llvm.return
}
+llvm.func @rocdl.schedule.bank(%val : f32, %ival : i32) -> f32 {
+ // CHECK-LABEL: rocdl.schedule.bank
+ // strict (default): bank 2 → encoded as 6 (2 | 4)
+ // CHECK: call float @llvm.amdgcn.schedule.bank.f32(float %{{.*}}, i32 6)
+ %0 = rocdl.schedule.bank %val, 2 : f32
+ // soft: bank 1 → encoded as 1
+ // CHECK: call i32 @llvm.amdgcn.schedule.bank.i32(i32 %{{.*}}, i32 1)
+ %1 = rocdl.schedule.bank %ival, 1 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 5746eb18bb61c5623df248b377deb63b999c5741 Mon Sep 17 00:00:00 2001
From: jli-melchior <Jun.Li at amd.com>
Date: Thu, 1 Oct 2026 01:15:32 +0000
Subject: [PATCH 3/3] [AMDGPU] Add FeatureXNACKOnOffModes to gfx1250 feature
sets
gfx1250 had FeatureSupportsXNACK but was missing FeatureXNACKOnOffModes,
causing the target parser to treat XNACK as hardwired-on. This produced
ELF code objects with e_flags XNACK=on (0x300), which the HIP runtime
rejects when the GPU is not in XNACK-on mode.
Add FeatureXNACKOnOffModes to both FeatureISAVersion12_50_STRICT and
FeatureISAVersion12_50 so XNACK defaults to "any" (0x100), matching the
system compiler (hipcc) behavior.
Co-Authored-By: Claude <noreply at anthropic.com>
---
llvm/lib/Target/AMDGPU/AMDGPU.td | 2 ++
1 file changed, 2 insertions(+)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.td b/llvm/lib/Target/AMDGPU/AMDGPU.td
index c3b4d53a7effa..a418786ee2f8e 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<
More information about the llvm-commits
mailing list