[PATCH] D115597: Prevent store value forwarding to distinct addrspace load

[Dmitry Borisenkov via Phabricator via llvm-commits]

dmitry created this revision.
dmitry added reviewers: hgreving, RKSimon, arsenm, spatel.
Herald added subscribers: asavonic, ecnelises, steven.zhang, hiraditya, jholewinski.
dmitry requested review of this revision.
Herald added subscribers: llvm-commits, wdng.
Herald added a project: LLVM.

DAGCombiner replaces `(load const_addr1)` directly chained with `(store (val, const_addr2))` with `val` if address space stripped const_addr1 == const_addr2. The patch fixes the issue by checking address spaces as well.
However, it might makes sense to not to chain together side effects that belong to different address spaces together in the first place and make `SelectionDAG::root`  address space aware.


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D115597

Files:
  llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp
  llvm/test/CodeGen/NVPTX/chain-different-as.ll


Index: llvm/test/CodeGen/NVPTX/chain-different-as.ll
===================================================================
--- /dev/null
+++ llvm/test/CodeGen/NVPTX/chain-different-as.ll
@@ -0,0 +1,10 @@
+; RUN: llc < %s -mtriple=nvptx64 | FileCheck %s
+define i64 @test() nounwind readnone {
+  %addr0 = inttoptr i64 1 to i64*
+; CHECK: mov.u64 %[[reg:rd[0-9]+]], 1
+; CHECK: ld.global.u64  %{{rd[0-9]+}}, [%[[reg]]]
+  %addr1 = inttoptr i64 1 to i64 addrspace(1)*
+  store i64 42, i64* %addr0
+  %res = load i64, i64 addrspace(1)* %addr1
+  ret i64 %res
+}
Index: llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp
===================================================================
--- llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp
+++ llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp
@@ -15980,7 +15980,7 @@
   SDValue Chain = LD->getOperand(0);
   StoreSDNode *ST = dyn_cast<StoreSDNode>(Chain.getNode());
   // TODO: Relax this restriction for unordered atomics (see D66309)
-  if (!ST || !ST->isSimple())
+  if (!ST || !ST->isSimple() || ST->getAddressSpace() != LD->getAddressSpace())
     return SDValue();
 
   EVT LDType = LD->getValueType(0);


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D115597.393742.patch
Type: text/x-patch
Size: 1145 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20211212/15a4a31f/attachment.bin>