[llvm] [AMDGPU] Add llvm.amdgcn.s.prefetch.inst intrinsic (PR #192440)
via llvm-commits
llvm-commits at lists.llvm.org
Thu Apr 16 05:50:43 PDT 2026
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-amdgpu
Author: Mirko BrkuĊĦanin (mbrkusanin)
<details>
<summary>Changes</summary>
Limited to constant addrspace unlike llvm.amdgcn.s.prefetch.data.
---
Full diff: https://github.com/llvm/llvm-project/pull/192440.diff
6 Files Affected:
- (modified) llvm/include/llvm/IR/IntrinsicsAMDGPU.td (+8)
- (modified) llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp (+1-1)
- (modified) llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp (+7-1)
- (modified) llvm/lib/Target/AMDGPU/SIISelLowering.cpp (+1)
- (modified) llvm/lib/Target/AMDGPU/SMInstructions.td (+20)
- (added) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.prefetch.inst.ll (+206)
``````````diff
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: {{.*}}
``````````
</details>
https://github.com/llvm/llvm-project/pull/192440
More information about the llvm-commits
mailing list