From 5a58d1a3c5e39a748e268f086afe3811f2a543c2 Mon Sep 17 00:00:00 2001 From: Stanislav Mekhanoshin Date: Fri, 6 Sep 2024 10:53:01 -0700 Subject: [PATCH 1/4] [AMDGPU] Error on non-global pointer with s_prefetch_data --- llvm/lib/Target/AMDGPU/SIISelLowering.cpp | 4 ++-- .../CodeGen/AMDGPU/llvm.amdgcn.s.prefetch.data.err.ll | 8 ++++++++ llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.prefetch.data.ll | 9 --------- 3 files changed, 10 insertions(+), 11 deletions(-) create mode 100644 llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.prefetch.data.err.ll diff --git a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp index accc3084217f2..27536dbcfd3a4 100644 --- a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp +++ b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp @@ -9929,9 +9929,9 @@ SDValue SITargetLowering::LowerINTRINSIC_VOID(SDValue Op, return SDValue(NewMI, 0); } case Intrinsic::amdgcn_s_prefetch_data: { - // For non-global address space preserve the chain and remove the call. if (!AMDGPU::isFlatGlobalAddrSpace(cast(Op)->getAddressSpace())) - return Op.getOperand(0); + report_fatal_error("s_prefetch_data only supports global or constant" + " memory"); return Op; } default: { diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.prefetch.data.err.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.prefetch.data.err.ll new file mode 100644 index 0000000000000..d135bc49db493 --- /dev/null +++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.prefetch.data.err.ll @@ -0,0 +1,8 @@ +; RUN: not --crash llc -march=amdgcn -mcpu=gfx1200 < %s 2>&1 | FileCheck --check-prefixes=GCN-ERR %s + +; GCN-ERR: LLVM ERROR: s_prefetch_data only supports global or constant memory +define amdgpu_ps void @prefetch_data_sgpr_base_imm_len_local(ptr addrspace(3) inreg %ptr) { +entry: + tail call void @llvm.amdgcn.s.prefetch.data.p3(ptr addrspace(3) %ptr, i32 31) + ret void +} diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.prefetch.data.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.prefetch.data.ll index 54c39d78adb58..b677f7863c14d 100644 --- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.prefetch.data.ll +++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.prefetch.data.ll @@ -110,15 +110,6 @@ entry: ret void } -define amdgpu_ps void @prefetch_data_sgpr_base_imm_len_local(ptr addrspace(3) inreg %ptr) { -; GCN-LABEL: prefetch_data_sgpr_base_imm_len_local: -; GCN: ; %bb.0: ; %entry -; GCN-NEXT: s_endpgm -entry: - tail call void @llvm.amdgcn.s.prefetch.data.p3(ptr addrspace(3) %ptr, i32 31) - ret void -} - define amdgpu_ps void @prefetch_data_vgpr_base_imm_len(ptr addrspace(4) %ptr) { ; GCN-LABEL: prefetch_data_vgpr_base_imm_len: ; GCN: ; %bb.0: ; %entry From c50067c556b1274b3ba6a941ba393ec7f6fd3601 Mon Sep 17 00:00:00 2001 From: Stanislav Mekhanoshin Date: Mon, 9 Sep 2024 12:16:02 -0700 Subject: [PATCH 2/4] Moved to verifier --- llvm/lib/IR/Verifier.cpp | 6 ++++++ llvm/lib/Target/AMDGPU/SIISelLowering.cpp | 4 ++-- .../AMDGPU/intrinsic-prefetch.ll} | 4 ++-- 3 files changed, 10 insertions(+), 4 deletions(-) rename llvm/test/{CodeGen/AMDGPU/llvm.amdgcn.s.prefetch.data.err.ll => Verifier/AMDGPU/intrinsic-prefetch.ll} (51%) diff --git a/llvm/lib/IR/Verifier.cpp b/llvm/lib/IR/Verifier.cpp index d8f3bab45b2a6..2d8c3521b7a4f 100644 --- a/llvm/lib/IR/Verifier.cpp +++ b/llvm/lib/IR/Verifier.cpp @@ -6281,6 +6281,12 @@ void Verifier::visitIntrinsicCall(Intrinsic::ID ID, CallBase &Call) { "Value for inactive lanes must be a VGPR function argument", &Call); break; } + case Intrinsic::amdgcn_s_prefetch_data: { + unsigned AS = Call.getArgOperand(0)->getType()->getPointerAddressSpace(); + Check(AS == 0 || AS == 1 || AS == 4, + "s_prefetch_data only supports global or constant memory"); + break; + } case Intrinsic::nvvm_setmaxnreg_inc_sync_aligned_u32: case Intrinsic::nvvm_setmaxnreg_dec_sync_aligned_u32: { Value *V = Call.getArgOperand(0); diff --git a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp index b2cd4377ad403..25cb8341c51d5 100644 --- a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp +++ b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp @@ -9935,9 +9935,9 @@ SDValue SITargetLowering::LowerINTRINSIC_VOID(SDValue Op, return SDValue(NewMI, 0); } case Intrinsic::amdgcn_s_prefetch_data: { + // For non-global address space preserve the chain and remove the call. if (!AMDGPU::isFlatGlobalAddrSpace(cast(Op)->getAddressSpace())) - report_fatal_error("s_prefetch_data only supports global or constant" - " memory"); + return Op.getOperand(0); return Op; } case Intrinsic::amdgcn_s_buffer_prefetch_data: { diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.prefetch.data.err.ll b/llvm/test/Verifier/AMDGPU/intrinsic-prefetch.ll similarity index 51% rename from llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.prefetch.data.err.ll rename to llvm/test/Verifier/AMDGPU/intrinsic-prefetch.ll index d135bc49db493..c50747c407092 100644 --- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.prefetch.data.err.ll +++ b/llvm/test/Verifier/AMDGPU/intrinsic-prefetch.ll @@ -1,8 +1,8 @@ -; RUN: not --crash llc -march=amdgcn -mcpu=gfx1200 < %s 2>&1 | FileCheck --check-prefixes=GCN-ERR %s +; RUN: not llvm-as %s -o /dev/null 2>&1 | FileCheck %s -; GCN-ERR: LLVM ERROR: s_prefetch_data only supports global or constant memory define amdgpu_ps void @prefetch_data_sgpr_base_imm_len_local(ptr addrspace(3) inreg %ptr) { entry: + ; CHECK:s_prefetch_data only supports global or constant memory tail call void @llvm.amdgcn.s.prefetch.data.p3(ptr addrspace(3) %ptr, i32 31) ret void } From 13f2ceff269b090ffa1c4e4e648cfa2adedeb6aa Mon Sep 17 00:00:00 2001 From: Stanislav Mekhanoshin Date: Wed, 11 Sep 2024 10:32:54 -0700 Subject: [PATCH 3/4] Address review comments - Using llvm/Support/AMDGPUAddrSpace.h - Moved isFlatGlobalAddrSpace/isExtendedGlobalAddrSpace there - Using intrinsic name in the error message --- llvm/include/llvm/Support/AMDGPUAddrSpace.h | 16 ++++++++++++++++ llvm/lib/IR/Verifier.cpp | 8 +++++--- llvm/lib/Target/AMDGPU/AMDGPU.h | 14 -------------- llvm/test/Verifier/AMDGPU/intrinsic-prefetch.ll | 2 +- 4 files changed, 22 insertions(+), 18 deletions(-) diff --git a/llvm/include/llvm/Support/AMDGPUAddrSpace.h b/llvm/include/llvm/Support/AMDGPUAddrSpace.h index c9d9bdd2f2fa3..78a321a559e7a 100644 --- a/llvm/include/llvm/Support/AMDGPUAddrSpace.h +++ b/llvm/include/llvm/Support/AMDGPUAddrSpace.h @@ -81,6 +81,22 @@ enum : unsigned { UNKNOWN_ADDRESS_SPACE = ~0u, }; } // end namespace AMDGPUAS + +namespace AMDGPU { +inline bool isFlatGlobalAddrSpace(unsigned AS) { + return AS == AMDGPUAS::GLOBAL_ADDRESS || + AS == AMDGPUAS::FLAT_ADDRESS || + AS == AMDGPUAS::CONSTANT_ADDRESS || + AS > AMDGPUAS::MAX_AMDGPU_ADDRESS; +} + +inline bool isExtendedGlobalAddrSpace(unsigned AS) { + return AS == AMDGPUAS::GLOBAL_ADDRESS || AS == AMDGPUAS::CONSTANT_ADDRESS || + AS == AMDGPUAS::CONSTANT_ADDRESS_32BIT || + AS > AMDGPUAS::MAX_AMDGPU_ADDRESS; +} +} // end namespace AMDGPU + } // end namespace llvm #endif // LLVM_SUPPORT_AMDGPUADDRSPACE_H diff --git a/llvm/lib/IR/Verifier.cpp b/llvm/lib/IR/Verifier.cpp index 2d8c3521b7a4f..06a67346fbf95 100644 --- a/llvm/lib/IR/Verifier.cpp +++ b/llvm/lib/IR/Verifier.cpp @@ -114,6 +114,7 @@ #include "llvm/IR/Value.h" #include "llvm/InitializePasses.h" #include "llvm/Pass.h" +#include "llvm/Support/AMDGPUAddrSpace.h" #include "llvm/Support/AtomicOrdering.h" #include "llvm/Support/Casting.h" #include "llvm/Support/CommandLine.h" @@ -6282,9 +6283,10 @@ void Verifier::visitIntrinsicCall(Intrinsic::ID ID, CallBase &Call) { break; } case Intrinsic::amdgcn_s_prefetch_data: { - unsigned AS = Call.getArgOperand(0)->getType()->getPointerAddressSpace(); - Check(AS == 0 || AS == 1 || AS == 4, - "s_prefetch_data only supports global or constant memory"); + Check( + AMDGPU::isFlatGlobalAddrSpace( + Call.getArgOperand(0)->getType()->getPointerAddressSpace()), + "llvm.amdgcn.s.prefetch.data only supports global or constant memory"); break; } case Intrinsic::nvvm_setmaxnreg_inc_sync_aligned_u32: diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.h b/llvm/lib/Target/AMDGPU/AMDGPU.h index 8d6e022e1e4d4..b193b44cee074 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPU.h +++ b/llvm/lib/Target/AMDGPU/AMDGPU.h @@ -461,20 +461,6 @@ enum TargetIndex { TI_SCRATCH_RSRC_DWORD3 }; -// FIXME: Missing constant_32bit -inline bool isFlatGlobalAddrSpace(unsigned AS) { - return AS == AMDGPUAS::GLOBAL_ADDRESS || - AS == AMDGPUAS::FLAT_ADDRESS || - AS == AMDGPUAS::CONSTANT_ADDRESS || - AS > AMDGPUAS::MAX_AMDGPU_ADDRESS; -} - -inline bool isExtendedGlobalAddrSpace(unsigned AS) { - return AS == AMDGPUAS::GLOBAL_ADDRESS || AS == AMDGPUAS::CONSTANT_ADDRESS || - AS == AMDGPUAS::CONSTANT_ADDRESS_32BIT || - AS > AMDGPUAS::MAX_AMDGPU_ADDRESS; -} - static inline bool addrspacesMayAlias(unsigned AS1, unsigned AS2) { static_assert(AMDGPUAS::MAX_AMDGPU_ADDRESS <= 9, "Addr space out of range"); diff --git a/llvm/test/Verifier/AMDGPU/intrinsic-prefetch.ll b/llvm/test/Verifier/AMDGPU/intrinsic-prefetch.ll index c50747c407092..1a7e949c31e2a 100644 --- a/llvm/test/Verifier/AMDGPU/intrinsic-prefetch.ll +++ b/llvm/test/Verifier/AMDGPU/intrinsic-prefetch.ll @@ -2,7 +2,7 @@ define amdgpu_ps void @prefetch_data_sgpr_base_imm_len_local(ptr addrspace(3) inreg %ptr) { entry: - ; CHECK:s_prefetch_data only supports global or constant memory + ; CHECK: llvm.amdgcn.s.prefetch.data only supports global or constant memory tail call void @llvm.amdgcn.s.prefetch.data.p3(ptr addrspace(3) %ptr, i32 31) ret void } From a407189156bb9f4b8996639a12552c4281850e3a Mon Sep 17 00:00:00 2001 From: Stanislav Mekhanoshin Date: Wed, 11 Sep 2024 10:40:07 -0700 Subject: [PATCH 4/4] clang-format --- llvm/include/llvm/Support/AMDGPUAddrSpace.h | 6 ++---- 1 file changed, 2 insertions(+), 4 deletions(-) diff --git a/llvm/include/llvm/Support/AMDGPUAddrSpace.h b/llvm/include/llvm/Support/AMDGPUAddrSpace.h index 78a321a559e7a..4a278d0acc23b 100644 --- a/llvm/include/llvm/Support/AMDGPUAddrSpace.h +++ b/llvm/include/llvm/Support/AMDGPUAddrSpace.h @@ -84,10 +84,8 @@ enum : unsigned { namespace AMDGPU { inline bool isFlatGlobalAddrSpace(unsigned AS) { - return AS == AMDGPUAS::GLOBAL_ADDRESS || - AS == AMDGPUAS::FLAT_ADDRESS || - AS == AMDGPUAS::CONSTANT_ADDRESS || - AS > AMDGPUAS::MAX_AMDGPU_ADDRESS; + return AS == AMDGPUAS::GLOBAL_ADDRESS || AS == AMDGPUAS::FLAT_ADDRESS || + AS == AMDGPUAS::CONSTANT_ADDRESS || AS > AMDGPUAS::MAX_AMDGPU_ADDRESS; } inline bool isExtendedGlobalAddrSpace(unsigned AS) {