diff --git a/llvm/include/llvm/Support/AMDGPUAddrSpace.h b/llvm/include/llvm/Support/AMDGPUAddrSpace.h index c9d9bdd2f2fa3..4a278d0acc23b 100644 --- a/llvm/include/llvm/Support/AMDGPUAddrSpace.h +++ b/llvm/include/llvm/Support/AMDGPUAddrSpace.h @@ -81,6 +81,20 @@ 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 d8f3bab45b2a6..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" @@ -6281,6 +6282,13 @@ 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: { + 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: case Intrinsic::nvvm_setmaxnreg_dec_sync_aligned_u32: { Value *V = Call.getArgOperand(0); 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/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 diff --git a/llvm/test/Verifier/AMDGPU/intrinsic-prefetch.ll b/llvm/test/Verifier/AMDGPU/intrinsic-prefetch.ll new file mode 100644 index 0000000000000..1a7e949c31e2a --- /dev/null +++ b/llvm/test/Verifier/AMDGPU/intrinsic-prefetch.ll @@ -0,0 +1,8 @@ +; RUN: not llvm-as %s -o /dev/null 2>&1 | FileCheck %s + +define amdgpu_ps void @prefetch_data_sgpr_base_imm_len_local(ptr addrspace(3) inreg %ptr) { +entry: + ; 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 +}