[llvm] [AMDGPU] Add llvm.amdgcn.s.prefetch.inst intrinsic (PR #192440)
Mirko BrkuĊĦanin via llvm-commits
llvm-commits at lists.llvm.org
Thu Apr 16 07:49:12 PDT 2026
https://github.com/mbrkusanin updated https://github.com/llvm/llvm-project/pull/192440
>From 88e46818a637361e6c15d7e4b1daaeffccd5603f Mon Sep 17 00:00:00 2001
From: Mirko Brkusanin <Mirko.Brkusanin at amd.com>
Date: Thu, 16 Apr 2026 14:46:06 +0200
Subject: [PATCH 1/2] [AMDGPU] Add llvm.amdgcn.s.prefetch.inst intrinsic
---
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 8 +
.../AMDGPU/AMDGPURegBankLegalizeRules.cpp | 2 +-
.../Target/AMDGPU/AMDGPURegisterBankInfo.cpp | 8 +-
llvm/lib/Target/AMDGPU/SIISelLowering.cpp | 1 +
llvm/lib/Target/AMDGPU/SMInstructions.td | 20 ++
.../AMDGPU/llvm.amdgcn.s.prefetch.inst.ll | 206 ++++++++++++++++++
6 files changed, 243 insertions(+), 2 deletions(-)
create mode 100644 llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.prefetch.inst.ll
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index 9528fb2b446bc..c7e566e8fd896 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -3195,6 +3195,14 @@ def int_amdgcn_s_prefetch_data :
"", [SDNPMemOperand]
>;
+def int_amdgcn_s_prefetch_inst :
+ Intrinsic<[],
+ [LLVMQualPointerType<4>, // Pointer to constant memory
+ llvm_i32_ty], // Length to prefetch 0-31 (1-32 chunks, units of 128 bytes)
+ [IntrInaccessibleMemOrArgMemOnly, IntrWillReturn, NoCapture<ArgIndex<0>>, IntrNoCallback, IntrNoFree],
+ "", [SDNPMemOperand]
+ >;
+
// llvm.amdgcn.ds.bpermute.fi.b32 <index> <src>
def int_amdgcn_ds_bpermute_fi_b32 :
ClangBuiltin<"__builtin_amdgcn_ds_bpermute_fi_b32">,
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
index 8028a4cb7c878..639a422b26267 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
@@ -1565,7 +1565,7 @@ RegBankLegalizeRules::RegBankLegalizeRules(const GCNSubtarget &_ST,
{amdgcn_s_get_named_barrier_state, amdgcn_s_get_barrier_state}, Standard)
.Uni(S32, {{Sgpr32}, {IntrId, SgprB32_M0}});
- addRulesForIOpcs({amdgcn_s_prefetch_data})
+ addRulesForIOpcs({amdgcn_s_prefetch_data, amdgcn_s_prefetch_inst})
.Any({{}, {{}, {IntrId, SgprB64_ReadFirstLane, SgprB32_ReadFirstLane}}});
addRulesForIOpcs({amdgcn_class})
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
index ebdd709c34f08..597f0f5fe281f 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
@@ -3381,6 +3381,11 @@ void AMDGPURegisterBankInfo::applyMappingImpl(
MI.eraseFromParent();
return;
}
+ case Intrinsic::amdgcn_s_prefetch_inst: {
+ constrainOpWithReadfirstlane(B, MI, 1);
+ constrainOpWithReadfirstlane(B, MI, 2);
+ return;
+ }
case Intrinsic::amdgcn_tensor_load_to_lds:
case Intrinsic::amdgcn_tensor_store_from_lds: {
constrainOpWithReadfirstlane(B, MI, 1);
@@ -5660,7 +5665,8 @@ AMDGPURegisterBankInfo::getInstrMapping(const MachineInstr &MI) const {
}
break;
}
- case Intrinsic::amdgcn_s_prefetch_data: {
+ case Intrinsic::amdgcn_s_prefetch_data:
+ case Intrinsic::amdgcn_s_prefetch_inst: {
OpdsMapping[1] = getSGPROpMapping(MI.getOperand(1).getReg(), MRI, *TRI);
OpdsMapping[2] = getSGPROpMapping(MI.getOperand(2).getReg(), MRI, *TRI);
break;
diff --git a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
index 752ea02119e03..5fc846474efa8 100644
--- a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
+++ b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
@@ -1768,6 +1768,7 @@ void SITargetLowering::getTgtMemIntrinsic(SmallVectorImpl<IntrinsicInfo> &Infos,
return;
}
case Intrinsic::amdgcn_s_prefetch_data:
+ case Intrinsic::amdgcn_s_prefetch_inst:
case Intrinsic::amdgcn_flat_prefetch:
case Intrinsic::amdgcn_global_prefetch: {
Info.opc = ISD::INTRINSIC_VOID;
diff --git a/llvm/lib/Target/AMDGPU/SMInstructions.td b/llvm/lib/Target/AMDGPU/SMInstructions.td
index ee8d29c77708b..045c8abdaaecd 100644
--- a/llvm/lib/Target/AMDGPU/SMInstructions.td
+++ b/llvm/lib/Target/AMDGPU/SMInstructions.td
@@ -1184,6 +1184,26 @@ let SubtargetPredicate = isGFX12Plus in {
(SIsbuffer_prefetch v4i32:$sbase, (SMRDBufferImm i32:$offset), imm:$len),
(S_BUFFER_PREFETCH_DATA SReg_128:$sbase, i32imm:$offset, (i32 SGPR_NULL), (as_i8timm $len))
>;
+
+ def : GCNPat <
+ (int_amdgcn_s_prefetch_inst (SMRDImm i64:$sbase, i32:$offset), (i32 SReg_32:$len)),
+ (S_PREFETCH_INST $sbase, $offset, $len, 0)
+ >;
+
+ def : GCNPat <
+ (int_amdgcn_s_prefetch_inst (i64 SReg_64:$sbase), (i32 SReg_32:$len)),
+ (S_PREFETCH_INST $sbase, 0, $len, 0)
+ >;
+
+ def : GCNPat <
+ (int_amdgcn_s_prefetch_inst (SMRDImm i64:$sbase, i32:$offset), imm:$len),
+ (S_PREFETCH_INST $sbase, $offset, (i32 SGPR_NULL), (as_i8timm $len))
+ >;
+
+ def : GCNPat <
+ (int_amdgcn_s_prefetch_inst (i64 SReg_64:$sbase), imm:$len),
+ (S_PREFETCH_INST $sbase, 0, (i32 SGPR_NULL), (as_i8timm $len))
+ >;
} // End let SubtargetPredicate = isGFX12Plus
//===----------------------------------------------------------------------===//
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.prefetch.inst.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.prefetch.inst.ll
new file mode 100644
index 0000000000000..4d752c5371e4d
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.prefetch.inst.ll
@@ -0,0 +1,206 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1200 < %s | FileCheck --check-prefixes=GCN,GFX12,GFX12-SDAG %s
+; RUN: llc -global-isel=1 -new-reg-bank-select -mtriple=amdgcn -mcpu=gfx1200 < %s | FileCheck --check-prefixes=GCN,GFX12,GFX12-GISEL %s
+; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1250 < %s | FileCheck --check-prefixes=GCN,GFX1250,GFX1250-SDAG %s
+; RUN: llc -global-isel=1 -new-reg-bank-select -mtriple=amdgcn -mcpu=gfx1250 < %s | FileCheck --check-prefixes=GCN,GFX1250,GFX1250-GISEL %s
+
+define amdgpu_ps void @prefetch_inst_sgpr_base_sgpr_len(ptr addrspace(4) inreg %ptr, i32 inreg %len) {
+; GFX12-LABEL: prefetch_inst_sgpr_base_sgpr_len:
+; GFX12: ; %bb.0:
+; GFX12-NEXT: s_prefetch_inst s[0:1], 0x0, s2, 0
+; GFX12-NEXT: s_endpgm
+;
+; GFX1250-LABEL: prefetch_inst_sgpr_base_sgpr_len:
+; GFX1250: ; %bb.0:
+; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX1250-NEXT: s_prefetch_inst s[0:1], 0x0, s2, 0
+; GFX1250-NEXT: s_endpgm
+ tail call void @llvm.amdgcn.s.prefetch.inst(ptr addrspace(4) %ptr, i32 %len)
+ ret void
+}
+
+define amdgpu_ps void @prefetch_inst_sgpr_imm_base_sgpr_len(ptr addrspace(4) inreg %ptr, i32 inreg %len) {
+; GFX12-LABEL: prefetch_inst_sgpr_imm_base_sgpr_len:
+; GFX12: ; %bb.0:
+; GFX12-NEXT: s_prefetch_inst s[0:1], 0x200, s2, 0
+; GFX12-NEXT: s_endpgm
+;
+; GFX1250-LABEL: prefetch_inst_sgpr_imm_base_sgpr_len:
+; GFX1250: ; %bb.0:
+; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX1250-NEXT: s_prefetch_inst s[0:1], 0x200, s2, 0
+; GFX1250-NEXT: s_endpgm
+ %gep = getelementptr i32, ptr addrspace(4) %ptr, i32 128
+ tail call void @llvm.amdgcn.s.prefetch.inst(ptr addrspace(4) %gep, i32 %len)
+ ret void
+}
+
+define amdgpu_ps void @prefetch_inst_sgpr_base_imm_len(ptr addrspace(4) inreg %ptr) {
+; GFX12-LABEL: prefetch_inst_sgpr_base_imm_len:
+; GFX12: ; %bb.0:
+; GFX12-NEXT: s_prefetch_inst s[0:1], 0x0, null, 31
+; GFX12-NEXT: s_endpgm
+;
+; GFX1250-LABEL: prefetch_inst_sgpr_base_imm_len:
+; GFX1250: ; %bb.0:
+; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX1250-NEXT: s_prefetch_inst s[0:1], 0x0, null, 31
+; GFX1250-NEXT: s_endpgm
+ tail call void @llvm.amdgcn.s.prefetch.inst(ptr addrspace(4) %ptr, i32 31)
+ ret void
+}
+
+define amdgpu_ps void @prefetch_inst_sgpr_imm_base_imm_len(ptr addrspace(4) inreg %ptr) {
+; GFX12-LABEL: prefetch_inst_sgpr_imm_base_imm_len:
+; GFX12: ; %bb.0:
+; GFX12-NEXT: s_prefetch_inst s[0:1], 0x200, null, 31
+; GFX12-NEXT: s_endpgm
+;
+; GFX1250-LABEL: prefetch_inst_sgpr_imm_base_imm_len:
+; GFX1250: ; %bb.0:
+; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX1250-NEXT: s_prefetch_inst s[0:1], 0x200, null, 31
+; GFX1250-NEXT: s_endpgm
+ %gep = getelementptr i32, ptr addrspace(4) %ptr, i32 128
+ tail call void @llvm.amdgcn.s.prefetch.inst(ptr addrspace(4) %gep, i32 31)
+ ret void
+}
+
+define amdgpu_ps void @prefetch_inst_vgpr_base_sgpr_len(ptr addrspace(4) %ptr, i32 inreg %len) {
+; GFX12-LABEL: prefetch_inst_vgpr_base_sgpr_len:
+; GFX12: ; %bb.0:
+; GFX12-NEXT: v_readfirstlane_b32 s2, v0
+; GFX12-NEXT: v_readfirstlane_b32 s3, v1
+; GFX12-NEXT: s_prefetch_inst s[2:3], 0x0, s0, 0
+; GFX12-NEXT: s_endpgm
+;
+; GFX1250-LABEL: prefetch_inst_vgpr_base_sgpr_len:
+; GFX1250: ; %bb.0:
+; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX1250-NEXT: v_readfirstlane_b32 s2, v0
+; GFX1250-NEXT: v_readfirstlane_b32 s3, v1
+; GFX1250-NEXT: s_prefetch_inst s[2:3], 0x0, s0, 0
+; GFX1250-NEXT: s_endpgm
+ tail call void @llvm.amdgcn.s.prefetch.inst(ptr addrspace(4) %ptr, i32 %len)
+ ret void
+}
+
+define amdgpu_ps void @prefetch_inst_vgpr_imm_base_sgpr_len(ptr addrspace(4) %ptr, i32 inreg %len) {
+; GFX12-SDAG-LABEL: prefetch_inst_vgpr_imm_base_sgpr_len:
+; GFX12-SDAG: ; %bb.0:
+; GFX12-SDAG-NEXT: v_readfirstlane_b32 s2, v0
+; GFX12-SDAG-NEXT: v_readfirstlane_b32 s3, v1
+; GFX12-SDAG-NEXT: s_prefetch_inst s[2:3], 0x200, s0, 0
+; GFX12-SDAG-NEXT: s_endpgm
+;
+; GFX12-GISEL-LABEL: prefetch_inst_vgpr_imm_base_sgpr_len:
+; GFX12-GISEL: ; %bb.0:
+; GFX12-GISEL-NEXT: v_add_co_u32 v0, vcc_lo, 0x200, v0
+; GFX12-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
+; GFX12-GISEL-NEXT: v_add_co_ci_u32_e64 v1, null, 0, v1, vcc_lo
+; GFX12-GISEL-NEXT: v_readfirstlane_b32 s2, v0
+; GFX12-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_2)
+; GFX12-GISEL-NEXT: v_readfirstlane_b32 s3, v1
+; GFX12-GISEL-NEXT: s_prefetch_inst s[2:3], 0x0, s0, 0
+; GFX12-GISEL-NEXT: s_endpgm
+;
+; GFX1250-SDAG-LABEL: prefetch_inst_vgpr_imm_base_sgpr_len:
+; GFX1250-SDAG: ; %bb.0:
+; GFX1250-SDAG-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX1250-SDAG-NEXT: v_readfirstlane_b32 s2, v0
+; GFX1250-SDAG-NEXT: v_readfirstlane_b32 s3, v1
+; GFX1250-SDAG-NEXT: s_prefetch_inst s[2:3], 0x200, s0, 0
+; GFX1250-SDAG-NEXT: s_endpgm
+;
+; GFX1250-GISEL-LABEL: prefetch_inst_vgpr_imm_base_sgpr_len:
+; GFX1250-GISEL: ; %bb.0:
+; GFX1250-GISEL-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX1250-GISEL-NEXT: v_add_co_u32 v0, vcc_lo, 0x200, v0
+; GFX1250-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
+; GFX1250-GISEL-NEXT: v_add_co_ci_u32_e64 v1, null, 0, v1, vcc_lo
+; GFX1250-GISEL-NEXT: v_readfirstlane_b32 s2, v0
+; GFX1250-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_2)
+; GFX1250-GISEL-NEXT: v_readfirstlane_b32 s3, v1
+; GFX1250-GISEL-NEXT: s_prefetch_inst s[2:3], 0x0, s0, 0
+; GFX1250-GISEL-NEXT: s_endpgm
+ %gep = getelementptr i32, ptr addrspace(4) %ptr, i32 128
+ tail call void @llvm.amdgcn.s.prefetch.inst(ptr addrspace(4) %gep, i32 %len)
+ ret void
+}
+
+define amdgpu_ps void @prefetch_inst_sgpr_base_vgpr_len(ptr addrspace(4) inreg %ptr, i32 %len) {
+; GFX12-LABEL: prefetch_inst_sgpr_base_vgpr_len:
+; GFX12: ; %bb.0:
+; GFX12-NEXT: v_readfirstlane_b32 s2, v0
+; GFX12-NEXT: s_prefetch_inst s[0:1], 0x0, s2, 0
+; GFX12-NEXT: s_endpgm
+;
+; GFX1250-LABEL: prefetch_inst_sgpr_base_vgpr_len:
+; GFX1250: ; %bb.0:
+; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX1250-NEXT: v_readfirstlane_b32 s2, v0
+; GFX1250-NEXT: s_prefetch_inst s[0:1], 0x0, s2, 0
+; GFX1250-NEXT: s_endpgm
+ tail call void @llvm.amdgcn.s.prefetch.inst(ptr addrspace(4) %ptr, i32 %len)
+ ret void
+}
+
+define amdgpu_ps void @prefetch_inst_vgpr_base_imm_len(ptr addrspace(4) %ptr) {
+; GFX12-LABEL: prefetch_inst_vgpr_base_imm_len:
+; GFX12: ; %bb.0:
+; GFX12-NEXT: v_readfirstlane_b32 s0, v0
+; GFX12-NEXT: v_readfirstlane_b32 s1, v1
+; GFX12-NEXT: s_prefetch_inst s[0:1], 0x0, null, 0
+; GFX12-NEXT: s_endpgm
+;
+; GFX1250-LABEL: prefetch_inst_vgpr_base_imm_len:
+; GFX1250: ; %bb.0:
+; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX1250-NEXT: v_readfirstlane_b32 s0, v0
+; GFX1250-NEXT: v_readfirstlane_b32 s1, v1
+; GFX1250-NEXT: s_prefetch_inst s[0:1], 0x0, null, 0
+; GFX1250-NEXT: s_endpgm
+ tail call void @llvm.amdgcn.s.prefetch.inst(ptr addrspace(4) %ptr, i32 0)
+ ret void
+}
+
+define amdgpu_ps void @prefetch_inst_vgpr_base_vgpr_len(ptr addrspace(4) %ptr, i32 %len) {
+; GFX12-SDAG-LABEL: prefetch_inst_vgpr_base_vgpr_len:
+; GFX12-SDAG: ; %bb.0:
+; GFX12-SDAG-NEXT: v_readfirstlane_b32 s2, v2
+; GFX12-SDAG-NEXT: v_readfirstlane_b32 s0, v0
+; GFX12-SDAG-NEXT: v_readfirstlane_b32 s1, v1
+; GFX12-SDAG-NEXT: s_prefetch_inst s[0:1], 0x0, s2, 0
+; GFX12-SDAG-NEXT: s_endpgm
+;
+; GFX12-GISEL-LABEL: prefetch_inst_vgpr_base_vgpr_len:
+; GFX12-GISEL: ; %bb.0:
+; GFX12-GISEL-NEXT: v_readfirstlane_b32 s0, v0
+; GFX12-GISEL-NEXT: v_readfirstlane_b32 s1, v1
+; GFX12-GISEL-NEXT: v_readfirstlane_b32 s2, v2
+; GFX12-GISEL-NEXT: s_prefetch_inst s[0:1], 0x0, s2, 0
+; GFX12-GISEL-NEXT: s_endpgm
+;
+; GFX1250-SDAG-LABEL: prefetch_inst_vgpr_base_vgpr_len:
+; GFX1250-SDAG: ; %bb.0:
+; GFX1250-SDAG-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX1250-SDAG-NEXT: v_readfirstlane_b32 s2, v2
+; GFX1250-SDAG-NEXT: v_readfirstlane_b32 s0, v0
+; GFX1250-SDAG-NEXT: v_readfirstlane_b32 s1, v1
+; GFX1250-SDAG-NEXT: s_prefetch_inst s[0:1], 0x0, s2, 0
+; GFX1250-SDAG-NEXT: s_endpgm
+;
+; GFX1250-GISEL-LABEL: prefetch_inst_vgpr_base_vgpr_len:
+; GFX1250-GISEL: ; %bb.0:
+; GFX1250-GISEL-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX1250-GISEL-NEXT: v_readfirstlane_b32 s0, v0
+; GFX1250-GISEL-NEXT: v_readfirstlane_b32 s1, v1
+; GFX1250-GISEL-NEXT: v_readfirstlane_b32 s2, v2
+; GFX1250-GISEL-NEXT: s_prefetch_inst s[0:1], 0x0, s2, 0
+; GFX1250-GISEL-NEXT: s_endpgm
+ tail call void @llvm.amdgcn.s.prefetch.inst(ptr addrspace(4) %ptr, i32 %len)
+ ret void
+}
+
+;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:
+; GCN: {{.*}}
>From 0097b17458be2b0da279b6d412d2c990c483f30e Mon Sep 17 00:00:00 2001
From: Mirko Brkusanin <Mirko.Brkusanin at amd.com>
Date: Thu, 16 Apr 2026 16:41:05 +0200
Subject: [PATCH 2/2] support flat/global
---
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 2 +-
.../Target/AMDGPU/AMDGPURegisterBankInfo.cpp | 8 +---
llvm/lib/Target/AMDGPU/SIISelLowering.cpp | 3 +-
.../AMDGPU/llvm.amdgcn.s.prefetch.inst.ll | 48 +++++++++++++++----
4 files changed, 44 insertions(+), 17 deletions(-)
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index c7e566e8fd896..db401ff19aba8 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -3197,7 +3197,7 @@ def int_amdgcn_s_prefetch_data :
def int_amdgcn_s_prefetch_inst :
Intrinsic<[],
- [LLVMQualPointerType<4>, // Pointer to constant memory
+ [llvm_anyptr_ty, // Pointer to instruction memory
llvm_i32_ty], // Length to prefetch 0-31 (1-32 chunks, units of 128 bytes)
[IntrInaccessibleMemOrArgMemOnly, IntrWillReturn, NoCapture<ArgIndex<0>>, IntrNoCallback, IntrNoFree],
"", [SDNPMemOperand]
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
index 597f0f5fe281f..af0e8b3506e46 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
@@ -3371,7 +3371,8 @@ void AMDGPURegisterBankInfo::applyMappingImpl(
constrainOpWithReadfirstlane(B, MI, 2);
return;
}
- case Intrinsic::amdgcn_s_prefetch_data: {
+ case Intrinsic::amdgcn_s_prefetch_data:
+ case Intrinsic::amdgcn_s_prefetch_inst: {
Register PtrReg = MI.getOperand(1).getReg();
unsigned AS = MRI.getType(PtrReg).getAddressSpace();
if (AMDGPU::isFlatGlobalAddrSpace(AS)) {
@@ -3381,11 +3382,6 @@ void AMDGPURegisterBankInfo::applyMappingImpl(
MI.eraseFromParent();
return;
}
- case Intrinsic::amdgcn_s_prefetch_inst: {
- constrainOpWithReadfirstlane(B, MI, 1);
- constrainOpWithReadfirstlane(B, MI, 2);
- return;
- }
case Intrinsic::amdgcn_tensor_load_to_lds:
case Intrinsic::amdgcn_tensor_store_from_lds: {
constrainOpWithReadfirstlane(B, MI, 1);
diff --git a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
index 5fc846474efa8..61fab7461a8c5 100644
--- a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
+++ b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
@@ -12386,7 +12386,8 @@ SDValue SITargetLowering::LowerINTRINSIC_VOID(SDValue Op,
auto *NewMI = DAG.getMachineNode(Opc, DL, Op->getVTList(), Ops);
return SDValue(NewMI, 0);
}
- case Intrinsic::amdgcn_s_prefetch_data: {
+ case Intrinsic::amdgcn_s_prefetch_data:
+ case Intrinsic::amdgcn_s_prefetch_inst: {
// For non-global address space preserve the chain and remove the call.
if (!AMDGPU::isFlatGlobalAddrSpace(cast<MemSDNode>(Op)->getAddressSpace()))
return Op.getOperand(0);
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.prefetch.inst.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.prefetch.inst.ll
index 4d752c5371e4d..9423518a47887 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.prefetch.inst.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.prefetch.inst.ll
@@ -15,7 +15,7 @@ define amdgpu_ps void @prefetch_inst_sgpr_base_sgpr_len(ptr addrspace(4) inreg %
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
; GFX1250-NEXT: s_prefetch_inst s[0:1], 0x0, s2, 0
; GFX1250-NEXT: s_endpgm
- tail call void @llvm.amdgcn.s.prefetch.inst(ptr addrspace(4) %ptr, i32 %len)
+ tail call void @llvm.amdgcn.s.prefetch.inst.p4(ptr addrspace(4) %ptr, i32 %len)
ret void
}
@@ -31,7 +31,7 @@ define amdgpu_ps void @prefetch_inst_sgpr_imm_base_sgpr_len(ptr addrspace(4) inr
; GFX1250-NEXT: s_prefetch_inst s[0:1], 0x200, s2, 0
; GFX1250-NEXT: s_endpgm
%gep = getelementptr i32, ptr addrspace(4) %ptr, i32 128
- tail call void @llvm.amdgcn.s.prefetch.inst(ptr addrspace(4) %gep, i32 %len)
+ tail call void @llvm.amdgcn.s.prefetch.inst.p4(ptr addrspace(4) %gep, i32 %len)
ret void
}
@@ -46,7 +46,7 @@ define amdgpu_ps void @prefetch_inst_sgpr_base_imm_len(ptr addrspace(4) inreg %p
; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
; GFX1250-NEXT: s_prefetch_inst s[0:1], 0x0, null, 31
; GFX1250-NEXT: s_endpgm
- tail call void @llvm.amdgcn.s.prefetch.inst(ptr addrspace(4) %ptr, i32 31)
+ tail call void @llvm.amdgcn.s.prefetch.inst.p4(ptr addrspace(4) %ptr, i32 31)
ret void
}
@@ -62,7 +62,7 @@ define amdgpu_ps void @prefetch_inst_sgpr_imm_base_imm_len(ptr addrspace(4) inre
; GFX1250-NEXT: s_prefetch_inst s[0:1], 0x200, null, 31
; GFX1250-NEXT: s_endpgm
%gep = getelementptr i32, ptr addrspace(4) %ptr, i32 128
- tail call void @llvm.amdgcn.s.prefetch.inst(ptr addrspace(4) %gep, i32 31)
+ tail call void @llvm.amdgcn.s.prefetch.inst.p4(ptr addrspace(4) %gep, i32 31)
ret void
}
@@ -81,7 +81,7 @@ define amdgpu_ps void @prefetch_inst_vgpr_base_sgpr_len(ptr addrspace(4) %ptr, i
; GFX1250-NEXT: v_readfirstlane_b32 s3, v1
; GFX1250-NEXT: s_prefetch_inst s[2:3], 0x0, s0, 0
; GFX1250-NEXT: s_endpgm
- tail call void @llvm.amdgcn.s.prefetch.inst(ptr addrspace(4) %ptr, i32 %len)
+ tail call void @llvm.amdgcn.s.prefetch.inst.p4(ptr addrspace(4) %ptr, i32 %len)
ret void
}
@@ -124,7 +124,7 @@ define amdgpu_ps void @prefetch_inst_vgpr_imm_base_sgpr_len(ptr addrspace(4) %pt
; GFX1250-GISEL-NEXT: s_prefetch_inst s[2:3], 0x0, s0, 0
; GFX1250-GISEL-NEXT: s_endpgm
%gep = getelementptr i32, ptr addrspace(4) %ptr, i32 128
- tail call void @llvm.amdgcn.s.prefetch.inst(ptr addrspace(4) %gep, i32 %len)
+ tail call void @llvm.amdgcn.s.prefetch.inst.p4(ptr addrspace(4) %gep, i32 %len)
ret void
}
@@ -141,7 +141,37 @@ define amdgpu_ps void @prefetch_inst_sgpr_base_vgpr_len(ptr addrspace(4) inreg %
; GFX1250-NEXT: v_readfirstlane_b32 s2, v0
; GFX1250-NEXT: s_prefetch_inst s[0:1], 0x0, s2, 0
; GFX1250-NEXT: s_endpgm
- tail call void @llvm.amdgcn.s.prefetch.inst(ptr addrspace(4) %ptr, i32 %len)
+ tail call void @llvm.amdgcn.s.prefetch.inst.p4(ptr addrspace(4) %ptr, i32 %len)
+ ret void
+}
+
+define amdgpu_ps void @prefetch_inst_sgpr_base_imm_len_global(ptr addrspace(1) inreg %ptr) {
+; GFX12-LABEL: prefetch_inst_sgpr_base_imm_len_global:
+; GFX12: ; %bb.0:
+; GFX12-NEXT: s_prefetch_inst s[0:1], 0x0, null, 31
+; GFX12-NEXT: s_endpgm
+;
+; GFX1250-LABEL: prefetch_inst_sgpr_base_imm_len_global:
+; GFX1250: ; %bb.0:
+; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX1250-NEXT: s_prefetch_inst s[0:1], 0x0, null, 31
+; GFX1250-NEXT: s_endpgm
+ tail call void @llvm.amdgcn.s.prefetch.inst.p1(ptr addrspace(1) %ptr, i32 31)
+ ret void
+}
+
+define amdgpu_ps void @prefetch_inst_sgpr_base_imm_len_flat(ptr inreg %ptr) {
+; GFX12-LABEL: prefetch_inst_sgpr_base_imm_len_flat:
+; GFX12: ; %bb.0:
+; GFX12-NEXT: s_prefetch_inst s[0:1], 0x0, null, 31
+; GFX12-NEXT: s_endpgm
+;
+; GFX1250-LABEL: prefetch_inst_sgpr_base_imm_len_flat:
+; GFX1250: ; %bb.0:
+; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX1250-NEXT: s_prefetch_inst s[0:1], 0x0, null, 31
+; GFX1250-NEXT: s_endpgm
+ tail call void @llvm.amdgcn.s.prefetch.inst.p0(ptr %ptr, i32 31)
ret void
}
@@ -160,7 +190,7 @@ define amdgpu_ps void @prefetch_inst_vgpr_base_imm_len(ptr addrspace(4) %ptr) {
; GFX1250-NEXT: v_readfirstlane_b32 s1, v1
; GFX1250-NEXT: s_prefetch_inst s[0:1], 0x0, null, 0
; GFX1250-NEXT: s_endpgm
- tail call void @llvm.amdgcn.s.prefetch.inst(ptr addrspace(4) %ptr, i32 0)
+ tail call void @llvm.amdgcn.s.prefetch.inst.p4(ptr addrspace(4) %ptr, i32 0)
ret void
}
@@ -198,7 +228,7 @@ define amdgpu_ps void @prefetch_inst_vgpr_base_vgpr_len(ptr addrspace(4) %ptr, i
; GFX1250-GISEL-NEXT: v_readfirstlane_b32 s2, v2
; GFX1250-GISEL-NEXT: s_prefetch_inst s[0:1], 0x0, s2, 0
; GFX1250-GISEL-NEXT: s_endpgm
- tail call void @llvm.amdgcn.s.prefetch.inst(ptr addrspace(4) %ptr, i32 %len)
+ tail call void @llvm.amdgcn.s.prefetch.inst.p4(ptr addrspace(4) %ptr, i32 %len)
ret void
}
More information about the llvm-commits
mailing list