[llvm] [AMDGPU] Add llvm.amdgcn.s.prefetch.inst intrinsic (PR #192440)
Mirko BrkuĊĦanin via llvm-commits
llvm-commits at lists.llvm.org
Mon Apr 27 07:13:32 PDT 2026
https://github.com/mbrkusanin updated https://github.com/llvm/llvm-project/pull/192440
>From eb3bf4821413a343a5431a222475e912f5aeae5e 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/4] [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 8631985de9a0a..f93be7c9387fa 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -3197,6 +3197,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 77802f5f4349a..3569e3945bb99 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
@@ -1613,7 +1613,7 @@ RegBankLegalizeRules::RegBankLegalizeRules(const GCNSubtarget &_ST,
addRulesForIOpcs({amdgcn_global_prefetch}).Any({{}, {{}, {IntrId, VgprP1}}});
- 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 f08e12a7fbf31..ed7dfe491cbff 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 ca757a993e61e0d454d03c31194df66bb6fe7e67 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/4] 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 f93be7c9387fa..10d02eb9142da 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -3199,7 +3199,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 ed7dfe491cbff..6b341b18eeebe 100644
--- a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
+++ b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
@@ -12398,7 +12398,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
}
>From 8d3e142fa0c11747141023ec143b030dfec365a0 Mon Sep 17 00:00:00 2001
From: Mirko Brkusanin <Mirko.Brkusanin at amd.com>
Date: Tue, 21 Apr 2026 14:55:00 +0200
Subject: [PATCH 3/4] more tests
---
.../AMDGPU/llvm.amdgcn.s.prefetch.inst.ll | 299 +++++++++++++++++-
1 file changed, 289 insertions(+), 10 deletions(-)
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 9423518a47887..ac93af78e2890 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.prefetch.inst.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.prefetch.inst.ll
@@ -145,6 +145,93 @@ define amdgpu_ps void @prefetch_inst_sgpr_base_vgpr_len(ptr addrspace(4) inreg %
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.p4(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.p4(ptr addrspace(4) %ptr, i32 %len)
+ ret void
+}
+define amdgpu_ps void @prefetch_inst_sgpr_base_sgpr_len_global(ptr addrspace(1) inreg %ptr, i32 inreg %len) {
+; GFX12-LABEL: prefetch_inst_sgpr_base_sgpr_len_global:
+; 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_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, s2, 0
+; GFX1250-NEXT: s_endpgm
+ tail call void @llvm.amdgcn.s.prefetch.inst.p1(ptr addrspace(1) %ptr, i32 %len)
+ ret void
+}
+
+define amdgpu_ps void @prefetch_inst_sgpr_imm_base_sgpr_len_global(ptr addrspace(1) inreg %ptr, i32 inreg %len) {
+; GFX12-LABEL: prefetch_inst_sgpr_imm_base_sgpr_len_global:
+; 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_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], 0x200, s2, 0
+; GFX1250-NEXT: s_endpgm
+ %gep = getelementptr i32, ptr addrspace(1) %ptr, i32 128
+ tail call void @llvm.amdgcn.s.prefetch.inst.p1(ptr addrspace(1) %gep, 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:
@@ -160,6 +247,146 @@ define amdgpu_ps void @prefetch_inst_sgpr_base_imm_len_global(ptr addrspace(1) i
ret void
}
+define amdgpu_ps void @prefetch_inst_sgpr_imm_base_imm_len_global(ptr addrspace(1) inreg %ptr) {
+; GFX12-LABEL: prefetch_inst_sgpr_imm_base_imm_len_global:
+; 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_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], 0x200, null, 31
+; GFX1250-NEXT: s_endpgm
+ %gep = getelementptr i32, ptr addrspace(1) %ptr, i32 128
+ tail call void @llvm.amdgcn.s.prefetch.inst.p1(ptr addrspace(1) %gep, i32 31)
+ ret void
+}
+
+define amdgpu_ps void @prefetch_inst_vgpr_base_sgpr_len_global(ptr addrspace(1) %ptr, i32 inreg %len) {
+; GFX12-LABEL: prefetch_inst_vgpr_base_sgpr_len_global:
+; 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_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: 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.p1(ptr addrspace(1) %ptr, i32 %len)
+ ret void
+}
+
+define amdgpu_ps void @prefetch_inst_vgpr_base_imm_len_global(ptr addrspace(1) %ptr) {
+; GFX12-LABEL: prefetch_inst_vgpr_base_imm_len_global:
+; 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_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: 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.p1(ptr addrspace(1) %ptr, i32 0)
+ ret void
+}
+
+define amdgpu_ps void @prefetch_inst_sgpr_base_vgpr_len_global(ptr addrspace(1) inreg %ptr, i32 %len) {
+; GFX12-LABEL: prefetch_inst_sgpr_base_vgpr_len_global:
+; 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_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: 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.p1(ptr addrspace(1) %ptr, i32 %len)
+ ret void
+}
+
+define amdgpu_ps void @prefetch_inst_vgpr_base_vgpr_len_global(ptr addrspace(1) %ptr, i32 %len) {
+; GFX12-SDAG-LABEL: prefetch_inst_vgpr_base_vgpr_len_global:
+; 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_global:
+; 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_global:
+; 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_global:
+; 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.p1(ptr addrspace(1) %ptr, i32 %len)
+ ret void
+}
+
+define amdgpu_ps void @prefetch_inst_sgpr_base_sgpr_len_flat(ptr inreg %ptr, i32 inreg %len) {
+; GFX12-LABEL: prefetch_inst_sgpr_base_sgpr_len_flat:
+; 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_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, s2, 0
+; GFX1250-NEXT: s_endpgm
+ tail call void @llvm.amdgcn.s.prefetch.inst.p0(ptr %ptr, i32 %len)
+ ret void
+}
+
+define amdgpu_ps void @prefetch_inst_sgpr_imm_base_sgpr_len_flat(ptr inreg %ptr, i32 inreg %len) {
+; GFX12-LABEL: prefetch_inst_sgpr_imm_base_sgpr_len_flat:
+; 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_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], 0x200, s2, 0
+; GFX1250-NEXT: s_endpgm
+ %gep = getelementptr i32, ptr %ptr, i32 128
+ tail call void @llvm.amdgcn.s.prefetch.inst.p0(ptr %gep, i32 %len)
+ 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:
@@ -175,27 +402,79 @@ define amdgpu_ps void @prefetch_inst_sgpr_base_imm_len_flat(ptr inreg %ptr) {
ret void
}
-define amdgpu_ps void @prefetch_inst_vgpr_base_imm_len(ptr addrspace(4) %ptr) {
-; GFX12-LABEL: prefetch_inst_vgpr_base_imm_len:
+define amdgpu_ps void @prefetch_inst_sgpr_imm_base_imm_len_flat(ptr inreg %ptr) {
+; GFX12-LABEL: prefetch_inst_sgpr_imm_base_imm_len_flat:
+; 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_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], 0x200, null, 31
+; GFX1250-NEXT: s_endpgm
+ %gep = getelementptr i32, ptr %ptr, i32 128
+ tail call void @llvm.amdgcn.s.prefetch.inst.p0(ptr %gep, i32 31)
+ ret void
+}
+
+define amdgpu_ps void @prefetch_inst_vgpr_base_sgpr_len_flat(ptr %ptr, i32 inreg %len) {
+; GFX12-LABEL: prefetch_inst_vgpr_base_sgpr_len_flat:
+; 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_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: 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.p0(ptr %ptr, i32 %len)
+ ret void
+}
+
+define amdgpu_ps void @prefetch_inst_vgpr_base_imm_len_flat(ptr %ptr) {
+; GFX12-LABEL: prefetch_inst_vgpr_base_imm_len_flat:
; 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-LABEL: prefetch_inst_vgpr_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: 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.p4(ptr addrspace(4) %ptr, i32 0)
+ tail call void @llvm.amdgcn.s.prefetch.inst.p0(ptr %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:
+define amdgpu_ps void @prefetch_inst_sgpr_base_vgpr_len_flat(ptr inreg %ptr, i32 %len) {
+; GFX12-LABEL: prefetch_inst_sgpr_base_vgpr_len_flat:
+; 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_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: 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.p0(ptr %ptr, i32 %len)
+ ret void
+}
+
+define amdgpu_ps void @prefetch_inst_vgpr_base_vgpr_len_flat(ptr %ptr, i32 %len) {
+; GFX12-SDAG-LABEL: prefetch_inst_vgpr_base_vgpr_len_flat:
; GFX12-SDAG: ; %bb.0:
; GFX12-SDAG-NEXT: v_readfirstlane_b32 s2, v2
; GFX12-SDAG-NEXT: v_readfirstlane_b32 s0, v0
@@ -203,7 +482,7 @@ define amdgpu_ps void @prefetch_inst_vgpr_base_vgpr_len(ptr addrspace(4) %ptr, i
; 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-LABEL: prefetch_inst_vgpr_base_vgpr_len_flat:
; GFX12-GISEL: ; %bb.0:
; GFX12-GISEL-NEXT: v_readfirstlane_b32 s0, v0
; GFX12-GISEL-NEXT: v_readfirstlane_b32 s1, v1
@@ -211,7 +490,7 @@ define amdgpu_ps void @prefetch_inst_vgpr_base_vgpr_len(ptr addrspace(4) %ptr, i
; 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-LABEL: prefetch_inst_vgpr_base_vgpr_len_flat:
; 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
@@ -220,7 +499,7 @@ define amdgpu_ps void @prefetch_inst_vgpr_base_vgpr_len(ptr addrspace(4) %ptr, i
; 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-LABEL: prefetch_inst_vgpr_base_vgpr_len_flat:
; 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
@@ -228,7 +507,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.p4(ptr addrspace(4) %ptr, i32 %len)
+ tail call void @llvm.amdgcn.s.prefetch.inst.p0(ptr %ptr, i32 %len)
ret void
}
>From 64354a9970180c211f31944dbc528d3430a86d57 Mon Sep 17 00:00:00 2001
From: Mirko Brkusanin <Mirko.Brkusanin at amd.com>
Date: Mon, 27 Apr 2026 15:20:06 +0200
Subject: [PATCH 4/4] test with funcion target
---
.../AMDGPU/llvm.amdgcn.s.prefetch.inst.ll | 28 +++++++++++++++++++
1 file changed, 28 insertions(+)
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 ac93af78e2890..1e5e3656be10a 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.prefetch.inst.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.prefetch.inst.ll
@@ -201,6 +201,7 @@ define amdgpu_ps void @prefetch_inst_vgpr_base_vgpr_len(ptr addrspace(4) %ptr, i
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_sgpr_len_global(ptr addrspace(1) inreg %ptr, i32 inreg %len) {
; GFX12-LABEL: prefetch_inst_sgpr_base_sgpr_len_global:
; GFX12: ; %bb.0:
@@ -511,5 +512,32 @@ define amdgpu_ps void @prefetch_inst_vgpr_base_vgpr_len_flat(ptr %ptr, i32 %len)
ret void
}
+define amdgpu_ps void @prefetch_inst_function_target() {
+; GFX12-LABEL: prefetch_inst_function_target:
+; GFX12: ; %bb.0:
+; GFX12-NEXT: s_getpc_b64 s[0:1]
+; GFX12-NEXT: s_sext_i32_i16 s1, s1
+; GFX12-NEXT: s_add_co_u32 s0, s0, target at gotpcrel32@lo+8
+; GFX12-NEXT: s_add_co_ci_u32 s1, s1, target at gotpcrel32@hi+16
+; GFX12-NEXT: s_load_b64 s[0:1], s[0:1], 0x0
+; GFX12-NEXT: s_wait_kmcnt 0x0
+; GFX12-NEXT: s_prefetch_inst s[0:1], 0x0, null, 31
+; GFX12-NEXT: s_endpgm
+;
+; GFX1250-LABEL: prefetch_inst_function_target:
+; 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_get_pc_i64 s[0:1]
+; GFX1250-NEXT: s_add_nc_u64 s[0:1], s[0:1], target at gotpcrel+4
+; GFX1250-NEXT: s_load_b64 s[0:1], s[0:1], 0x0 nv
+; GFX1250-NEXT: s_wait_kmcnt 0x0
+; 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 @target, i32 31)
+ ret void
+}
+
+declare amdgpu_ps void @target()
+
;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:
; GCN: {{.*}}
More information about the llvm-commits
mailing list