[Mlir-commits] [clang] [llvm] [mlir] [AMDGPU] Add schedule bank hint for gfx1250 to get high perf (PR #227953)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Tue Oct 20 00:21:34 PDT 2026
https://github.com/jli-melchior updated https://github.com/llvm/llvm-project/pull/227953
>From e34a2863d02690ad2870fe4806596202939121bf 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/4] [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 | 18 ++
.../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 | 33 +++
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, 527 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 34a1bd2d301a3..c4daf33b8da29 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.td
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td
@@ -1011,6 +1011,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 2f29e9557ba56..0f73b36cc3519 100644
--- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
+++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
@@ -2283,6 +2283,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 b1f6fa33649bd..6a467f39fd756 100644
--- a/clang/lib/Sema/SemaAMDGPU.cpp
+++ b/clang/lib/Sema/SemaAMDGPU.cpp
@@ -465,6 +465,24 @@ 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 0fdf8df507bec..61db05c2a92a7 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 ed80bc4628b6f..6b7063f7904de 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -4478,4 +4478,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 75a83c55db2bc..92371186bb567 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPU.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPU.h
@@ -222,6 +222,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 dc2abb60567ff..45d9bad615d61 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp
@@ -3353,6 +3353,39 @@ 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 372d5f5acab21..d0d85faa1cf12 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..8427b751b704e
--- /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 6a5b27e5775c6..e66d5ecd42ab4 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 "AMDGPUScheduleBank.h"
#include "AMDGPUSplitModule.h"
#include "AMDGPUTargetObjectFile.h"
#include "AMDGPUTargetTransformInfo.h"
@@ -693,6 +694,7 @@ extern "C" LLVM_ABI LLVM_EXTERNAL_VISIBILITY void LLVMInitializeAMDGPUTarget() {
initializeAMDGPUAsmPrinterPass(*PR);
initializeAMDGPUDAGToDAGISelLegacyPass(*PR);
initializeAMDGPUPrepareAGPRAllocLegacyPass(*PR);
+ initializeAMDGPUScheduleBankPass(*PR);
initializeGCNDPPCombineLegacyPass(*PR);
initializeSILowerI1CopiesLegacyPass(*PR);
initializeAMDGPUGlobalISelDivergenceLoweringLegacyPass(*PR);
@@ -1846,6 +1848,7 @@ void GCNPassConfig::addFastRegAlloc() {
}
void GCNPassConfig::addPreRegAlloc() {
+ addPass(&AMDGPUScheduleBankID);
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(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 b6c5f63379344..73a0eb08696d8 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 c9e0e4c79f8cc..8d37a02019896 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 0989263433f41..26819d1fba235 100644
--- a/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp
@@ -55,6 +55,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;
@@ -4329,6 +4334,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 b70adfdd92101..27e1f5b3c01c1 100644
--- a/llvm/lib/Target/AMDGPU/SIRegisterInfo.h
+++ b/llvm/lib/Target/AMDGPU/SIRegisterInfo.h
@@ -38,7 +38,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 76bcf69ef24b26b64c37317d0cde011b6bee1b0e 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/4] [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 26819d1fba235..4723f7f07b59f 100644
--- a/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp
@@ -4385,8 +4385,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 2f067b8ba98fd..3bbc47cbdda87 100644
--- a/mlir/test/Target/LLVMIR/rocdl.mlir
+++ b/mlir/test/Target/LLVMIR/rocdl.mlir
@@ -451,6 +451,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 3708ed967cde38661880557cf1f67f4dcb7bd8c4 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/4] [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/GCNFeatures.td | 2 ++
1 file changed, 2 insertions(+)
diff --git a/llvm/lib/Target/AMDGPU/GCNFeatures.td b/llvm/lib/Target/AMDGPU/GCNFeatures.td
index 23c79f5949122..0f380c4745937 100644
--- a/llvm/lib/Target/AMDGPU/GCNFeatures.td
+++ b/llvm/lib/Target/AMDGPU/GCNFeatures.td
@@ -2600,6 +2600,7 @@ def FeatureISAVersion12_50_STRICT : FeatureSet<
FeatureSlowMaxMinMulI64Insts,
FeatureXNACKOnOffModes,
FeatureSupportsXNACK,
+ FeatureXNACKOnOffModes,
])>;
def FeatureISAVersion12_50 : FeatureSet<
@@ -2638,6 +2639,7 @@ def FeatureISAVersion12_50 : FeatureSet<
FeatureSlowMaxMinMulI64Insts,
FeatureXNACKOnOffModes,
FeatureSupportsXNACK,
+ FeatureXNACKOnOffModes,
])>;
def FeatureISAVersion12_51 : FeatureSet<
>From e42950022f6745f6721547c021af04b767a9964d Mon Sep 17 00:00:00 2001
From: Jun Li <jli10004 at amd.com>
Date: Tue, 20 Oct 2026 07:19:19 +0000
Subject: [PATCH 4/4] [AMDGPU] Rename schedule_bank to sched_bank and split
strict into separate ImmArg
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit
1. Rename schedule_bank → sched_bank across the full stack (intrinsic,
builtin, pseudo instructions, pass, ROCDL op, tests) to align with
existing LLVM naming conventions (sched_barrier, sched_group_barrier).
2. Split the strict flag from a bit-packed encoding (bank | (strict << 2))
into a separate ImmArg on the intrinsic:
Before: llvm.amdgcn.schedule.bank(value, encoded_bank)
After: llvm.amdgcn.sched.bank(value, bank, strict)
This removes manual bit-packing in Clang CodeGen and decoding in the
backend pass, making the IR more readable and extensible.
Co-Authored-By: Claude Opus 4 <noreply at anthropic.com>
---
clang/include/clang/Basic/BuiltinsAMDGPU.td | 2 +-
clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp | 16 ++---
clang/lib/Sema/SemaAMDGPU.cpp | 2 +-
.../builtins-amdgcn-gfx1250-sched-bank.cl | 71 +++++++++++++++++++
.../builtins-amdgcn-gfx1250-schedule-bank.cl | 71 -------------------
.../builtins-amdgcn-error-gfx1250-param.cl | 10 +--
.../builtins-amdgcn-error-gfx1250.cl | 2 +-
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 11 +--
llvm/lib/Target/AMDGPU/AMDGPU.h | 4 +-
llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp | 20 +++---
llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def | 2 +-
...PUScheduleBank.cpp => AMDGPUSchedBank.cpp} | 54 +++++++-------
...AMDGPUScheduleBank.h => AMDGPUSchedBank.h} | 12 ++--
.../lib/Target/AMDGPU/AMDGPUTargetMachine.cpp | 8 +--
llvm/lib/Target/AMDGPU/CMakeLists.txt | 2 +-
llvm/lib/Target/AMDGPU/SIInstructions.td | 20 +++---
llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp | 4 +-
llvm/test/CodeGen/AMDGPU/sched-bank-hints.mir | 35 +++++++++
.../{schedule-bank.ll => sched-bank.ll} | 6 +-
.../CodeGen/AMDGPU/schedule-bank-hints.mir | 35 ---------
mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td | 14 ++--
mlir/test/Dialect/LLVMIR/rocdl.mlir | 10 +--
mlir/test/Target/LLVMIR/rocdl.mlir | 16 ++---
23 files changed, 214 insertions(+), 213 deletions(-)
create mode 100644 clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-sched-bank.cl
delete mode 100644 clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-schedule-bank.cl
rename llvm/lib/Target/AMDGPU/{AMDGPUScheduleBank.cpp => AMDGPUSchedBank.cpp} (78%)
rename llvm/lib/Target/AMDGPU/{AMDGPUScheduleBank.h => AMDGPUSchedBank.h} (59%)
create mode 100644 llvm/test/CodeGen/AMDGPU/sched-bank-hints.mir
rename llvm/test/CodeGen/AMDGPU/{schedule-bank.ll => sched-bank.ll} (84%)
delete mode 100644 llvm/test/CodeGen/AMDGPU/schedule-bank-hints.mir
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td
index c4daf33b8da29..5047a81d0b2bc 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.td
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td
@@ -1011,7 +1011,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_sched_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 0f73b36cc3519..fc00d60b9b342 100644
--- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
+++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
@@ -2283,21 +2283,19 @@ 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: {
+ case AMDGPU::BI__builtin_amdgcn_sched_bank: {
llvm::Value *Val = EmitScalarExpr(E->getArg(0));
llvm::Value *Bank = EmitScalarExpr(E->getArg(1));
- llvm::Value *EncodedBank;
+ llvm::Value *Strict;
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));
+ Strict = EmitScalarExpr(E->getArg(2));
} else {
- // Default: strict (bank | 4)
- EncodedBank = Builder.CreateOr(Bank, Builder.getInt32(4));
+ // Default: strict
+ Strict = Builder.getInt1(true);
}
llvm::Function *F =
- CGM.getIntrinsic(Intrinsic::amdgcn_schedule_bank, Val->getType());
- return Builder.CreateCall(F, {Val, EncodedBank});
+ CGM.getIntrinsic(Intrinsic::amdgcn_sched_bank, Val->getType());
+ return Builder.CreateCall(F, {Val, Bank, Strict});
}
default:
return nullptr;
diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp
index 6a467f39fd756..464f2747fa7e8 100644
--- a/clang/lib/Sema/SemaAMDGPU.cpp
+++ b/clang/lib/Sema/SemaAMDGPU.cpp
@@ -465,7 +465,7 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(const TargetInfo &TI,
/*High=*/0) ||
SemaRef.BuiltinConstantArgRange(TheCall, /*ArgNum=*/2, /*Low=*/0,
/*High=*/0);
- case AMDGPU::BI__builtin_amdgcn_schedule_bank: {
+ case AMDGPU::BI__builtin_amdgcn_sched_bank: {
if (SemaRef.checkArgCountRange(TheCall, 2, 3))
return true;
Expr *DataArg = TheCall->getArg(0);
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-sched-bank.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-sched-bank.cl
new file mode 100644
index 0000000000000..ba5fb8d09fde2
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-sched-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_sched_bank_i32
+// CHECK: call i32 @llvm.amdgcn.sched.bank.i32(i32 %{{.*}}, i32 2, i1 false)
+void test_sched_bank_i32(global int *out, int x) {
+ *out = __builtin_amdgcn_sched_bank(x, 2, false);
+}
+
+// CHECK-LABEL: @test_sched_bank_f32
+// CHECK: call float @llvm.amdgcn.sched.bank.f32(float %{{.*}}, i32 3, i1 false)
+void test_sched_bank_f32(global float *out, float x) {
+ *out = __builtin_amdgcn_sched_bank(x, 3, false);
+}
+
+// CHECK-LABEL: @test_sched_bank_i64
+// CHECK: call i64 @llvm.amdgcn.sched.bank.i64(i64 %{{.*}}, i32 0, i1 false)
+void test_sched_bank_i64(global long *out, long x) {
+ *out = __builtin_amdgcn_sched_bank(x, 0, false);
+}
+
+// CHECK-LABEL: @test_sched_bank_f64_strict
+// CHECK: call double @llvm.amdgcn.sched.bank.f64(double %{{.*}}, i32 3, i1 true)
+void test_sched_bank_f64_strict(global double *out, double x) {
+ *out = __builtin_amdgcn_sched_bank(x, 3, true);
+}
+
+// CHECK-LABEL: @test_sched_bank_v2i32
+// CHECK: call <2 x i32> @llvm.amdgcn.sched.bank.v2i32(<2 x i32> %{{.*}}, i32 1, i1 false)
+void test_sched_bank_v2i32(global int2 *out, int2 x) {
+ *out = __builtin_amdgcn_sched_bank(x, 1, false);
+}
+
+// CHECK-LABEL: @test_sched_bank_v4i32_strict
+// CHECK: call <4 x i32> @llvm.amdgcn.sched.bank.v4i32(<4 x i32> %{{.*}}, i32 1, i1 true)
+void test_sched_bank_v4i32_strict(global int4 *out, int4 x) {
+ *out = __builtin_amdgcn_sched_bank(x, 1, true);
+}
+
+// CHECK-LABEL: @test_sched_bank_v8i32_strict
+// CHECK: call <8 x i32> @llvm.amdgcn.sched.bank.v8i32(<8 x i32> %{{.*}}, i32 2, i1 true)
+void test_sched_bank_v8i32_strict(global int8 *out, int8 x) {
+ *out = __builtin_amdgcn_sched_bank(x, 2, true);
+}
+
+// CHECK-LABEL: @test_sched_bank_v4f32_strict
+// CHECK: call <4 x float> @llvm.amdgcn.sched.bank.v4f32(<4 x float> %{{.*}}, i32 0, i1 true)
+void test_sched_bank_v4f32_strict(global float4 *out, float4 x) {
+ *out = __builtin_amdgcn_sched_bank(x, 0, true);
+}
+
+// CHECK-LABEL: @test_sched_bank_default_strict
+// CHECK: call i32 @llvm.amdgcn.sched.bank.i32(i32 %{{.*}}, i32 2, i1 true)
+void test_sched_bank_default_strict(global int *out, int x) {
+ // 2 args: strict=true by default
+ *out = __builtin_amdgcn_sched_bank(x, 2);
+}
+
+// CHECK-LABEL: @test_sched_bank_explicit_soft
+// CHECK: call i32 @llvm.amdgcn.sched.bank.i32(i32 %{{.*}}, i32 2, i1 false)
+void test_sched_bank_explicit_soft(global int *out, int x) {
+ // 3 args: explicit soft
+ *out = __builtin_amdgcn_sched_bank(x, 2, false);
+}
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-schedule-bank.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-schedule-bank.cl
deleted file mode 100644
index 17fc1f9c37f66..0000000000000
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-schedule-bank.cl
+++ /dev/null
@@ -1,71 +0,0 @@
-// 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 61db05c2a92a7..abf5e4491ddfa 100644
--- a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-param.cl
+++ b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-param.cl
@@ -217,11 +217,11 @@ void test_pk_add_min_max(global short2 *out, global ushort2 *uout, short2 a, sho
*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_sched_bank_non_const(int x, int b) {
+ __builtin_amdgcn_sched_bank(x, b, false); // expected-error {{'__builtin_amdgcn_sched_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]}}
+void test_sched_bank_out_of_range(int x) {
+ __builtin_amdgcn_sched_bank(x, 4, false); // expected-error {{argument value 4 is outside the valid range [0, 3]}}
+ __builtin_amdgcn_sched_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 01bb27927895b..4a973d215919c 100644
--- a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250.cl
+++ b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250.cl
@@ -20,5 +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}}
+ *out = __builtin_amdgcn_sched_bank(a, 2, false); // expected-error {{'__builtin_amdgcn_sched_bank' needs target feature gfx1250-insts}}
}
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index 6b7063f7904de..01e1043c69069 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -4484,10 +4484,13 @@ def int_amdgcn_dead: DefaultAttrsIntrinsic<[llvm_any_ty], [],
// 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 pre-RA AMDGPUSchedBank 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>>]>;
+// The strict operand (i1) controls whether the hint is mandatory (true) or
+// advisory (false).
+def int_amdgcn_sched_bank : DefaultAttrsIntrinsic<
+ [llvm_any_ty], [LLVMMatchType<0>, llvm_i32_ty, llvm_i1_ty],
+ [IntrNoMem, IntrSpeculatable, IntrWillReturn,
+ ImmArg<ArgIndex<1>>, ImmArg<ArgIndex<2>>]>;
}
diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.h b/llvm/lib/Target/AMDGPU/AMDGPU.h
index 92371186bb567..3c1ef6f9de2f1 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPU.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPU.h
@@ -222,8 +222,8 @@ struct AMDGPULowerIntrinsicsPass
void initializeAMDGPUPrepareAGPRAllocLegacyPass(PassRegistry &);
extern char &AMDGPUPrepareAGPRAllocLegacyID;
-void initializeAMDGPUScheduleBankPass(PassRegistry &);
-extern char &AMDGPUScheduleBankID;
+void initializeAMDGPUSchedBankPass(PassRegistry &);
+extern char &AMDGPUSchedBankID;
void initializeAMDGPUReserveWWMRegsLegacyPass(PassRegistry &);
extern char &AMDGPUReserveWWMRegsLegacyID;
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp b/llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp
index 45d9bad615d61..03fcae193c96f 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp
@@ -3353,27 +3353,28 @@ void AMDGPUDAGToDAGISel::SelectINTRINSIC_WO_CHAIN(SDNode *N) {
case Intrinsic::amdgcn_interp_p1_f16:
SelectInterpP1F16(N);
return;
- case Intrinsic::amdgcn_schedule_bank: {
+ case Intrinsic::amdgcn_sched_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) {
+ auto *StrictC = dyn_cast<ConstantSDNode>(N->getOperand(3));
+ // Bank and strict must be constants; fall back (and error) otherwise.
+ if (!BankC || !StrictC) {
SelectCode(N);
return;
}
unsigned PseudoOpc;
switch (N->getValueType(0).getSizeInBits()) {
case 32:
- PseudoOpc = AMDGPU::V_SCHEDULE_BANK_B32;
+ PseudoOpc = AMDGPU::V_SCHED_BANK_B32;
break;
case 64:
- PseudoOpc = AMDGPU::V_SCHEDULE_BANK_B64;
+ PseudoOpc = AMDGPU::V_SCHED_BANK_B64;
break;
case 128:
- PseudoOpc = AMDGPU::V_SCHEDULE_BANK_B128;
+ PseudoOpc = AMDGPU::V_SCHED_BANK_B128;
break;
case 256:
- PseudoOpc = AMDGPU::V_SCHEDULE_BANK_B256;
+ PseudoOpc = AMDGPU::V_SCHED_BANK_B256;
break;
default:
// Unsupported width: drop the hint, forward the value unchanged.
@@ -3383,7 +3384,10 @@ void AMDGPUDAGToDAGISel::SelectINTRINSIC_WO_CHAIN(SDNode *N) {
}
SDValue BankImm =
CurDAG->getTargetConstant(BankC->getZExtValue(), SDLoc(N), MVT::i32);
- CurDAG->SelectNodeTo(N, PseudoOpc, N->getVTList(), {Src, BankImm});
+ SDValue StrictImm =
+ CurDAG->getTargetConstant(StrictC->getZExtValue(), SDLoc(N), MVT::i1);
+ CurDAG->SelectNodeTo(N, PseudoOpc, N->getVTList(),
+ {Src, BankImm, StrictImm});
return;
}
case Intrinsic::amdgcn_permlane16_swap:
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def b/llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def
index d0d85faa1cf12..350002a4bd5de 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def
+++ b/llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def
@@ -137,7 +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-sched-bank", AMDGPUSchedBankPass())
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/AMDGPUSchedBank.cpp
similarity index 78%
rename from llvm/lib/Target/AMDGPU/AMDGPUScheduleBank.cpp
rename to llvm/lib/Target/AMDGPU/AMDGPUSchedBank.cpp
index 8427b751b704e..1da71dd668a2d 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUScheduleBank.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUSchedBank.cpp
@@ -1,4 +1,4 @@
-//===-- AMDGPUScheduleBank.cpp --------------------------------------------===//
+//===-- AMDGPUSchedBank.cpp --------------------------------------------===//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
@@ -7,8 +7,8 @@
//===----------------------------------------------------------------------===//
//
/// \file
-/// Lower the V_SCHEDULE_BANK_B* pseudos produced from
-/// llvm.amdgcn.schedule.bank. For each pseudo this pass:
+/// Lower the V_SCHED_BANK_B* pseudos produced from
+/// llvm.amdgcn.sched.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
@@ -23,7 +23,7 @@
//
//===----------------------------------------------------------------------===//
-#include "AMDGPUScheduleBank.h"
+#include "AMDGPUSchedBank.h"
#include "AMDGPU.h"
#include "GCNSubtarget.h"
#include "SIInstrInfo.h"
@@ -36,14 +36,14 @@
using namespace llvm;
-#define DEBUG_TYPE "amdgpu-schedule-bank"
+#define DEBUG_TYPE "amdgpu-sched-bank"
-static bool isScheduleBankPseudo(unsigned Opcode) {
+static bool isSchedBankPseudo(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:
+ case AMDGPU::V_SCHED_BANK_B32:
+ case AMDGPU::V_SCHED_BANK_B64:
+ case AMDGPU::V_SCHED_BANK_B128:
+ case AMDGPU::V_SCHED_BANK_B256:
return true;
default:
return false;
@@ -86,7 +86,7 @@ static void propagateBankHint(MachineRegisterInfo &MRI, Register Reg,
}
// Shared implementation used by both the legacy and new-PM passes.
-static bool runScheduleBank(MachineFunction &MF) {
+static bool runSchedBank(MachineFunction &MF) {
const GCNSubtarget &ST = MF.getSubtarget<GCNSubtarget>();
if (!ST.has1024AddressableVGPRs())
return false;
@@ -99,25 +99,21 @@ static bool runScheduleBank(MachineFunction &MF) {
for (MachineBasicBlock &MBB : MF) {
for (MachineInstr &MI : MBB) {
- if (!isScheduleBankPseudo(MI.getOpcode()))
+ if (!isSchedBankPseudo(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 Bank = MI.getOperand(2).getImm();
+ bool Strict = MI.getOperand(3).getImm() != 0;
unsigned HintKind =
Strict ? AMDGPURI::StrictBankHint : AMDGPURI::BankHint;
- // Bits 0..1 select bank; bit 2 requests strict allocation.
- if (EncodedBank > 7) {
+ if (Bank > 3)
Bank = 0;
- Strict = false;
- }
- LLVM_DEBUG(dbgs() << " schedule.bank: " << printReg(DstReg) << " <- "
+ LLVM_DEBUG(dbgs() << " sched.bank: " << printReg(DstReg) << " <- "
<< printReg(SrcReg) << " bank " << Bank << " strict "
<< Strict << '\n');
@@ -153,17 +149,17 @@ static bool runScheduleBank(MachineFunction &MF) {
namespace {
-class AMDGPUScheduleBank : public MachineFunctionPass {
+class AMDGPUSchedBank : public MachineFunctionPass {
public:
static char ID;
- AMDGPUScheduleBank() : MachineFunctionPass(ID) {}
+ AMDGPUSchedBank() : MachineFunctionPass(ID) {}
bool runOnMachineFunction(MachineFunction &MF) override {
- return runScheduleBank(MF);
+ return runSchedBank(MF);
}
- StringRef getPassName() const override { return "AMDGPU Schedule Bank"; }
+ StringRef getPassName() const override { return "AMDGPU Sched Bank"; }
void getAnalysisUsage(AnalysisUsage &AU) const override {
AU.setPreservesCFG();
@@ -173,17 +169,17 @@ class AMDGPUScheduleBank : public MachineFunctionPass {
} // end anonymous namespace
-INITIALIZE_PASS(AMDGPUScheduleBank, DEBUG_TYPE, "AMDGPU Schedule Bank", false,
+INITIALIZE_PASS(AMDGPUSchedBank, DEBUG_TYPE, "AMDGPU Sched Bank", false,
false)
-char AMDGPUScheduleBank::ID = 0;
+char AMDGPUSchedBank::ID = 0;
-char &llvm::AMDGPUScheduleBankID = AMDGPUScheduleBank::ID;
+char &llvm::AMDGPUSchedBankID = AMDGPUSchedBank::ID;
PreservedAnalyses
-AMDGPUScheduleBankPass::run(MachineFunction &MF,
+AMDGPUSchedBankPass::run(MachineFunction &MF,
MachineFunctionAnalysisManager &MFAM) {
- if (!runScheduleBank(MF))
+ if (!runSchedBank(MF))
return PreservedAnalyses::all();
return getMachineFunctionPassPreservedAnalyses().preserveSet<CFGAnalyses>();
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUScheduleBank.h b/llvm/lib/Target/AMDGPU/AMDGPUSchedBank.h
similarity index 59%
rename from llvm/lib/Target/AMDGPU/AMDGPUScheduleBank.h
rename to llvm/lib/Target/AMDGPU/AMDGPUSchedBank.h
index 79d6bab6e0556..3cff8cd765422 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUScheduleBank.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPUSchedBank.h
@@ -1,4 +1,4 @@
-//===- AMDGPUScheduleBank.h -------------------------------------*- C++- *-===//
+//===- AMDGPUSchedBank.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.
@@ -6,18 +6,18 @@
//
//===----------------------------------------------------------------------===//
-#ifndef LLVM_LIB_TARGET_AMDGPU_AMDGPUSCHEDULEBANK_H
-#define LLVM_LIB_TARGET_AMDGPU_AMDGPUSCHEDULEBANK_H
+#ifndef LLVM_LIB_TARGET_AMDGPU_AMDGPUSCHEDBANK_H
+#define LLVM_LIB_TARGET_AMDGPU_AMDGPUSCHEDBANK_H
#include "llvm/CodeGen/MachinePassManager.h"
namespace llvm {
-class AMDGPUScheduleBankPass
- : public OptionalPassInfoMixin<AMDGPUScheduleBankPass> {
+class AMDGPUSchedBankPass
+ : public OptionalPassInfoMixin<AMDGPUSchedBankPass> {
public:
PreservedAnalyses run(MachineFunction &MF,
MachineFunctionAnalysisManager &MFAM);
};
} // namespace llvm
-#endif // LLVM_LIB_TARGET_AMDGPU_AMDGPUSCHEDULEBANK_H
+#endif // LLVM_LIB_TARGET_AMDGPU_AMDGPUSCHEDBANK_H
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp b/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
index e66d5ecd42ab4..abdd1129116e3 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
@@ -35,7 +35,7 @@
#include "AMDGPURemoveIncompatibleFunctions.h"
#include "AMDGPUReserveWWMRegs.h"
#include "AMDGPUResourceUsageAnalysis.h"
-#include "AMDGPUScheduleBank.h"
+#include "AMDGPUSchedBank.h"
#include "AMDGPUSplitModule.h"
#include "AMDGPUTargetObjectFile.h"
#include "AMDGPUTargetTransformInfo.h"
@@ -694,7 +694,7 @@ extern "C" LLVM_ABI LLVM_EXTERNAL_VISIBILITY void LLVMInitializeAMDGPUTarget() {
initializeAMDGPUAsmPrinterPass(*PR);
initializeAMDGPUDAGToDAGISelLegacyPass(*PR);
initializeAMDGPUPrepareAGPRAllocLegacyPass(*PR);
- initializeAMDGPUScheduleBankPass(*PR);
+ initializeAMDGPUSchedBankPass(*PR);
initializeGCNDPPCombineLegacyPass(*PR);
initializeSILowerI1CopiesLegacyPass(*PR);
initializeAMDGPUGlobalISelDivergenceLoweringLegacyPass(*PR);
@@ -1848,7 +1848,7 @@ void GCNPassConfig::addFastRegAlloc() {
}
void GCNPassConfig::addPreRegAlloc() {
- addPass(&AMDGPUScheduleBankID);
+ addPass(&AMDGPUSchedBankID);
if (getOptLevel() != CodeGenOptLevel::None)
addPass(&AMDGPUPrepareAGPRAllocLegacyID);
if (getOptLevel() >= CodeGenOptLevel::Default && EnableMachinePipeliner)
@@ -2711,7 +2711,7 @@ Error AMDGPUCodeGenPassBuilder::addOptimizedRegAlloc(PassManagerWrapper &PMW) {
}
void AMDGPUCodeGenPassBuilder::addPreRegAlloc(PassManagerWrapper &PMW) {
- addMachineFunctionPass(AMDGPUScheduleBankPass(), PMW);
+ addMachineFunctionPass(AMDGPUSchedBankPass(), 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 73a0eb08696d8..5d2931f561d33 100644
--- a/llvm/lib/Target/AMDGPU/CMakeLists.txt
+++ b/llvm/lib/Target/AMDGPU/CMakeLists.txt
@@ -80,7 +80,7 @@ add_llvm_target(AMDGPUCodeGen
AMDGPULowerKernelAttributes.cpp
AMDGPULowerModuleLDSPass.cpp
AMDGPUPrepareAGPRAlloc.cpp
- AMDGPUScheduleBank.cpp
+ AMDGPUSchedBank.cpp
AMDGPULowerExecSync.cpp
AMDGPUSwLowerLDS.cpp
AMDGPUMachineFunctionInfo.cpp
diff --git a/llvm/lib/Target/AMDGPU/SIInstructions.td b/llvm/lib/Target/AMDGPU/SIInstructions.td
index 8d37a02019896..0b7f818497829 100644
--- a/llvm/lib/Target/AMDGPU/SIInstructions.td
+++ b/llvm/lib/Target/AMDGPU/SIInstructions.td
@@ -130,19 +130,19 @@ def V_CNDMASK_B64_PSEUDO : VOP3Common <(outs VReg_64:$vdst),
}
// 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
+// identity copies produced by lowering llvm.amdgcn.sched.bank. A pre-RA pass
+// (AMDGPUSchedBank) 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)>;
+ def V_SCHED_BANK_B32 : VPseudoInstSI<(outs VGPR_32:$vdst),
+ (ins VGPR_32:$src, i32imm:$bank, i1imm:$strict)>;
+ def V_SCHED_BANK_B64 : VPseudoInstSI<(outs VReg_64:$vdst),
+ (ins VReg_64:$src, i32imm:$bank, i1imm:$strict)>;
+ def V_SCHED_BANK_B128 : VPseudoInstSI<(outs VReg_128:$vdst),
+ (ins VReg_128:$src, i32imm:$bank, i1imm:$strict)>;
+ def V_SCHED_BANK_B256 : VPseudoInstSI<(outs VReg_256:$vdst),
+ (ins VReg_256:$src, i32imm:$bank, i1imm:$strict)>;
}
// 64-bit vector move instruction. This is mainly used by the
diff --git a/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp b/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp
index 4723f7f07b59f..09d8c18dc5ddf 100644
--- a/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp
@@ -57,7 +57,7 @@ static cl::opt<unsigned> StressSGPRLimit(
static cl::opt<bool> StrictVGPRBankHints(
"amdgpu-strict-vgpr-bank-hints",
- cl::desc("Restrict schedule.bank values to the requested VGPR bank"),
+ cl::desc("Restrict sched.bank values to the requested VGPR bank"),
cl::Hidden, cl::init(false));
std::array<std::vector<int16_t>, 32> SIRegisterInfo::RegSplitParts;
@@ -4369,7 +4369,7 @@ bool SIRegisterInfo::getRegAllocationHints(Register VirtReg,
}
case AMDGPURI::BankHint:
case AMDGPURI::StrictBankHint: {
- // Absolute bank preference from llvm.amdgcn.schedule.bank. The bank number
+ // Absolute bank preference from llvm.amdgcn.sched.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.
diff --git a/llvm/test/CodeGen/AMDGPU/sched-bank-hints.mir b/llvm/test/CodeGen/AMDGPU/sched-bank-hints.mir
new file mode 100644
index 0000000000000..ee15f45b8c74f
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/sched-bank-hints.mir
@@ -0,0 +1,35 @@
+# RUN: llc -mtriple=amdgcn -mcpu=gfx1250 -run-pass=amdgpu-sched-bank -o - %s | FileCheck %s
+
+# The AMDGPUSchedBank pass attaches a bank register-allocation hint (not
+# printed in MIR) and lowers each V_SCHED_BANK_B* pseudo to a plain COPY.
+
+---
+# CHECK-LABEL: name: lower_sched_bank_b32
+# CHECK: %1:vgpr_32 = COPY %0
+# CHECK-NOT: V_SCHED_BANK_B32
+name: lower_sched_bank_b32
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0
+ %0:vgpr_32 = COPY $vgpr0
+ %1:vgpr_32 = V_SCHED_BANK_B32 %0, 2, 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_bank_b64
+# CHECK: %1:vreg_64_align2 = COPY %0
+# CHECK-NOT: V_SCHED_BANK_B64
+name: lower_sched_bank_b64
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0_vgpr1
+ %0:vreg_64_align2 = COPY $vgpr0_vgpr1
+ %1:vreg_64_align2 = V_SCHED_BANK_B64 %0, 3, 1, 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/sched-bank.ll
similarity index 84%
rename from llvm/test/CodeGen/AMDGPU/schedule-bank.ll
rename to llvm/test/CodeGen/AMDGPU/sched-bank.ll
index f40681b1c6f2b..69751ee7134a8 100644
--- a/llvm/test/CodeGen/AMDGPU/schedule-bank.ll
+++ b/llvm/test/CodeGen/AMDGPU/sched-bank.ll
@@ -1,11 +1,11 @@
; RUN: llc -mtriple=amdgcn -mcpu=gfx1250 < %s | FileCheck -check-prefix=GCN %s
-; Verify that llvm.amdgcn.schedule.bank steers register allocation into the
+; Verify that llvm.amdgcn.sched.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)
+declare i32 @llvm.amdgcn.sched.bank.i32(i32, i32, i1)
; 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.
@@ -13,7 +13,7 @@ declare i32 @llvm.amdgcn.schedule.bank.i32(i32, i32)
; 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)
+ %h = call i32 @llvm.amdgcn.sched.bank.i32(i32 %x, i32 2, i1 true)
%y = add i32 %h, 1
store i32 %y, ptr addrspace(1) %out
ret void
diff --git a/llvm/test/CodeGen/AMDGPU/schedule-bank-hints.mir b/llvm/test/CodeGen/AMDGPU/schedule-bank-hints.mir
deleted file mode 100644
index 2ba66444e22df..0000000000000
--- a/llvm/test/CodeGen/AMDGPU/schedule-bank-hints.mir
+++ /dev/null
@@ -1,35 +0,0 @@
-# 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/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td b/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td
index 4b2194f1f7a22..0b63e028e00b9 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td
@@ -813,7 +813,7 @@ def ROCDL_IglpOpt : ROCDL_ConcreteNonMemIntrOp<"iglp.opt", [], 0, [0], ["variant
}];
}
-def ROCDL_ScheduleBankOp : ROCDL_IntrOp<"schedule.bank", [], [0],
+def ROCDL_SchedBankOp : ROCDL_IntrOp<"sched.bank", [], [0],
[AllTypesMatch<["res", "value"]>], 1, 0, 0, 0, [], []>,
Arguments<(ins LLVM_Type:$value, I32Attr:$bank,
OptionalAttr<UnitAttr>:$soft)> {
@@ -822,16 +822,16 @@ def ROCDL_ScheduleBankOp : ROCDL_IntrOp<"schedule.bank", [], [0],
$value `,` $bank (`soft` $soft^)? attr-dict `:` type($value)
}];
string llvmBuilder = [{
- auto schedOp = cast<ROCDL::ScheduleBankOp>(opInst);
+ auto schedOp = cast<ROCDL::SchedBankOp>(opInst);
auto *value = moduleTranslation.lookupValue(schedOp.getValue());
int32_t bank = schedOp.getBank();
- int32_t encodedBank = schedOp.getSoftAttr() ? bank : (bank | 4);
+ bool strict = !schedOp.getSoftAttr();
auto *fn = llvm::Intrinsic::getOrInsertDeclaration(
moduleTranslation.getLLVMModule(),
- llvm::Intrinsic::amdgcn_schedule_bank,
+ llvm::Intrinsic::amdgcn_sched_bank,
{value->getType()});
auto *inst = builder.CreateCall(fn,
- {value, builder.getInt32(encodedBank)});
+ {value, builder.getInt32(bank), builder.getInt1(strict)});
moduleTranslation.mapValue(opInst.getResult(0), inst);
}];
let description = [{
@@ -843,8 +843,8 @@ def ROCDL_ScheduleBankOp : ROCDL_IntrOp<"schedule.bank", [], [0],
Example:
```mlir
- %0 = rocdl.schedule.bank %val, 2 : f32 // strict bank 2 (default)
- %1 = rocdl.schedule.bank %val, 1 soft : f32 // soft bank 1
+ %0 = rocdl.sched.bank %val, 2 : f32 // strict bank 2 (default)
+ %1 = rocdl.sched.bank %val, 1 soft : f32 // soft bank 1
```
}];
}
diff --git a/mlir/test/Dialect/LLVMIR/rocdl.mlir b/mlir/test/Dialect/LLVMIR/rocdl.mlir
index ae77810a369e4..0ace8265cd563 100644
--- a/mlir/test/Dialect/LLVMIR/rocdl.mlir
+++ b/mlir/test/Dialect/LLVMIR/rocdl.mlir
@@ -168,11 +168,11 @@ 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
+func.func @rocdl_sched_bank(%val : f32, %ival : i32) -> f32 {
+ // CHECK: rocdl.sched.bank %{{.*}}, 2 : f32
+ %0 = rocdl.sched.bank %val, 2 : f32
+ // CHECK: rocdl.sched.bank %{{.*}}, 1 soft : i32
+ %1 = rocdl.sched.bank %ival, 1 soft : i32
llvm.return %0 : f32
}
diff --git a/mlir/test/Target/LLVMIR/rocdl.mlir b/mlir/test/Target/LLVMIR/rocdl.mlir
index 3bbc47cbdda87..12a3cedaafa09 100644
--- a/mlir/test/Target/LLVMIR/rocdl.mlir
+++ b/mlir/test/Target/LLVMIR/rocdl.mlir
@@ -451,14 +451,14 @@ 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.func @rocdl.sched.bank(%val : f32, %ival : i32) -> f32 {
+ // CHECK-LABEL: rocdl.sched.bank
+ // strict (default): bank 2, strict true
+ // CHECK: call float @llvm.amdgcn.sched.bank.f32(float %{{.*}}, i32 2, i1 true)
+ %0 = rocdl.sched.bank %val, 2 : f32
+ // soft: bank 1, strict false
+ // CHECK: call i32 @llvm.amdgcn.sched.bank.i32(i32 %{{.*}}, i32 1, i1 false)
+ %1 = rocdl.sched.bank %ival, 1 soft : i32
llvm.return %0 : f32
}
More information about the Mlir-commits
mailing list