[PATCH] D153712: [SelectionDAG] Add memory size for CSEMap ID calculation

[Alex MacLean via Phabricator via llvm-commits]

AlexM created this revision.
Herald added subscribers: mattd, gchakrabarti, asavonic, pengfei, hiraditya.
Herald added a project: All.
AlexM updated this revision to Diff 534358.
AlexM added a comment.
AlexM added reviewers: tra, craig.topper, nikic.
Herald added a subscriber: StephenFan.
AlexM added a reviewer: goldstein.w.n.
AlexM published this revision for review.
Herald added subscribers: llvm-commits, jholewinski.
Herald added a project: LLVM.

Remove erroneously included file


In NVPTX `ReplaceVectorLoad()`, i1 and i8 types are promoted to i16,
followed by a truncate operation. Thus, v2i8 (or v2i1) and v2i16 will
have the same VTList, which causes a collision in CSEMap.

To differentiate the original VTList, let's add the size in generating
an ID. Otherwise the compiler crashes in refineAlignment:
`MMO->getSize() == getSize() && "Size mismatch!"`


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D153712

Files:
  llvm/lib/CodeGen/SelectionDAG/SelectionDAG.cpp
  llvm/test/CodeGen/NVPTX/dag-cse.ll


Index: llvm/test/CodeGen/NVPTX/dag-cse.ll
===================================================================
--- /dev/null
+++ llvm/test/CodeGen/NVPTX/dag-cse.ll
@@ -0,0 +1,24 @@
+; RUN: llc < %s -march=nvptx64 | FileCheck %s
+
+%st = type { i8, i8, i16 }
+
+ at a = internal addrspace(1) global %st zeroinitializer, align 8
+ at b = internal addrspace(1) global i32 0, align 8
+ at c = internal addrspace(1) global i32 0, align 8
+
+; Verify that loads with different memory types are not subject to CSE
+; once they are promoted to the same type.
+;
+; CHECK: ld.global.v2.u8  {%[[B1:rs[0-9]+]], %[[B2:rs[0-9]+]]}, [a];
+; CHECK: st.global.v2.u8  [b], {%[[B1]], %[[B2]]};
+;
+; CHECK: ld.global.v2.u16 {%[[C1:rs[0-9]+]], %[[C2:rs[0-9]+]]}, [a];
+; CHECK: st.global.v2.u16 [c], {%[[C1]], %[[C2]]};
+
+define void @test1() #0 {
+  %1 = load <2 x i8>, ptr addrspace(1) @a, align 8
+  store <2 x i8> %1, ptr addrspace(1) @b, align 8
+  %2 = load <2 x i16>, ptr addrspace(1) @a, align 8
+  store <2 x i16> %2, ptr addrspace(1) @c, align 8
+  ret void
+}
Index: llvm/lib/CodeGen/SelectionDAG/SelectionDAG.cpp
===================================================================
--- llvm/lib/CodeGen/SelectionDAG/SelectionDAG.cpp
+++ llvm/lib/CodeGen/SelectionDAG/SelectionDAG.cpp
@@ -927,6 +927,7 @@
     ID.AddInteger(MN->getRawSubclassData());
     ID.AddInteger(MN->getPointerInfo().getAddrSpace());
     ID.AddInteger(MN->getMemOperand()->getFlags());
+    ID.AddInteger(MN->getMemOperand()->getSize());
   }
 }
 
@@ -8066,6 +8067,7 @@
         Opcode, dl.getIROrder(), VTList, MemVT, MMO));
     ID.AddInteger(MMO->getPointerInfo().getAddrSpace());
     ID.AddInteger(MMO->getFlags());
+    ID.AddInteger(MMO->getSize());
     void *IP = nullptr;
     if (SDNode *E = FindNodeOrInsertPos(ID, dl, IP)) {
       cast<MemIntrinsicSDNode>(E)->refineAlignment(MMO);


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D153712.534358.patch
Type: text/x-patch
Size: 1853 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20230625/b3b6f232/attachment.bin>