llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-llvm-ir Author: jli (jli-melchior) <details> <summary>Changes</summary> --- Patch is 34.64 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/227953.diff 23 Files Affected: - (modified) clang/include/clang/Basic/BuiltinsAMDGPU.td (+1) - (modified) clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp (+16) - (modified) clang/lib/Sema/SemaAMDGPU.cpp (+19) - (added) clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-schedule-bank.cl (+71) - (modified) clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-param.cl (+9) - (modified) clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250.cl (+1) - (modified) llvm/include/llvm/IR/IntrinsicsAMDGPU.td (+12) - (modified) llvm/lib/Target/AMDGPU/AMDGPU.h (+3) - (modified) llvm/lib/Target/AMDGPU/AMDGPU.td (+2) - (modified) llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp (+25) - (modified) llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def (+1) - (added) llvm/lib/Target/AMDGPU/AMDGPUScheduleBank.cpp (+190) - (added) llvm/lib/Target/AMDGPU/AMDGPUScheduleBank.h (+23) - (modified) llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp (+4) - (modified) llvm/lib/Target/AMDGPU/CMakeLists.txt (+1) - (modified) llvm/lib/Target/AMDGPU/SIInstructions.td (+16) - (modified) llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp (+58) - (modified) llvm/lib/Target/AMDGPU/SIRegisterInfo.h (+1-1) - (added) llvm/test/CodeGen/AMDGPU/schedule-bank-hints.mir (+35) - (added) llvm/test/CodeGen/AMDGPU/schedule-bank.ll (+32) - (modified) mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td (+36) - (modified) mlir/test/Dialect/LLVMIR/rocdl.mlir (+8) - (modified) mlir/test/Target/LLVMIR/rocdl.mlir (+11) ``````````diff diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td index 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/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< 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; +... [truncated] `````````` </details> https://github.com/llvm/llvm-project/pull/227953 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
