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 <[email protected]> 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 <[email protected]> --- 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 <[email protected]> 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 <[email protected]> --- 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 <[email protected]> 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 <[email protected]> --- 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< _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
