[llvm] [AMDGPU] Error on non-global pointer with s_prefetch_data (PR #107624)

[Stanislav Mekhanoshin via llvm-commits]

https://github.com/rampitec updated https://github.com/llvm/llvm-project/pull/107624

>From 5a58d1a3c5e39a748e268f086afe3811f2a543c2 Mon Sep 17 00:00:00 2001
From: Stanislav Mekhanoshin <Stanislav.Mekhanoshin at amd.com>
Date: Fri, 6 Sep 2024 10:53:01 -0700
Subject: [PATCH 1/2] [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 accc3084217f2b..27536dbcfd3a49 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<MemSDNode>(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 00000000000000..d135bc49db4937
--- /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 54c39d78adb583..b677f7863c14dc 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 <Stanislav.Mekhanoshin at amd.com>
Date: Mon, 9 Sep 2024 12:16:02 -0700
Subject: [PATCH 2/2] 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 d8f3bab45b2a65..2d8c3521b7a4f2 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 b2cd4377ad403c..25cb8341c51d53 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<MemSDNode>(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 d135bc49db4937..c50747c4070928 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
 }