[SelectionDAG] Add memory size for CSEMap ID calculation

[AlexM]

Summary

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!"

[AlexM]

Remove erroneously included file

[craig.topper]

Would memory VT work? That's what regular load/store use I think.

[AlexM]

Switch to memory VT

[AlexM]

In D153712#4447855, @craig.topper wrote:

Would memory VT work? That's what regular load/store use I think.

Yea, that would work as well, I've the diff over.

[craig.topper]

LGTM

[AlexM]

@craig.topper Thanks, can you please commit this on my behalf?

llvm/lib/CodeGen/SelectionDAG/SelectionDAG.cpp

static void AddNodeIDCustom(FoldingSetNodeID &ID, const SDNode *N) {
   } // end switch (N->getOpcode())
 
   // MemIntrinsic nodes could also have subclass data, address spaces, and flags
   // to check.
   if (auto *MN = dyn_cast<MemIntrinsicSDNode>(N)) {
     ID.AddInteger(MN->getRawSubclassData());
     ID.AddInteger(MN->getPointerInfo().getAddrSpace());
     ID.AddInteger(MN->getMemOperand()->getFlags());
+    ID.AddInteger(MN->getMemOperand()->getSize());
   }
 }
 
 /// AddNodeIDNode - Generic routine for adding a nodes info to the NodeID
 /// data.
 static void AddNodeIDNode(FoldingSetNodeID &ID, const SDNode *N) {
   AddNodeIDOpcode(ID, N->getOpcode());
   // Add the return value info.
SDValue SelectionDAG::getMemIntrinsicNode(unsigned Opcode, const SDLoc &dl,
   MemIntrinsicSDNode *N;
   if (VTList.VTs[VTList.NumVTs-1] != MVT::Glue) {
     FoldingSetNodeID ID;
     AddNodeIDNode(ID, Opcode, VTList, Ops);
     ID.AddInteger(getSyntheticNodeSubclassData<MemIntrinsicSDNode>(
         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);
       return SDValue(E, 0);
     }
 
     N = newSDNode<MemIntrinsicSDNode>(Opcode, dl.getIROrder(), dl.getDebugLoc(),
                                       VTList, MemVT, MMO);

llvm/test/CodeGen/NVPTX/dag-cse.ll

+; RUN: llc < %s -march=nvptx64 | FileCheck %s
+
+%st = type { i8, i8, i16 }
+
+@a = internal addrspace(1) global %st zeroinitializer, align 8
+@b = internal addrspace(1) global i32 0, align 8
+@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
+}

llvm/test/CodeGen/X86/dag-cse-target.ll

+; RUN: llc < %s -mtriple=x86_64-unknown -debug 2>&1 | FileCheck %s
+
+; Verify that CSE works for identical target memory nodes:
+;
+; CHECK:     Creating new node: t{{[0-9]+}}: v8f16,ch = X86ISD::VBROADCAST_LOAD<(load (s16) from `ptr undef` + 4)> t0, undef:i64
+; CHECK-NOT: Creating new node: t{{[0-9]+}}: v8f16,ch = X86ISD::VBROADCAST_LOAD<(load (s16) from `ptr undef` + 4)> t0, undef:i64
+;
+define void @t() #1 {
+entry:
+  %elem = load half, ptr undef, align 16
+  %load1 = insertelement <16 x half> poison, half %elem, i64 15
+  %load2 = load <16 x half>, ptr undef, align 2
+  %shuffle = shufflevector <16 x half> %load1, <16 x half> %load2, <16 x i32>
+    <i32 15, i32 16, i32 17, i32 18, i32 19, i32 20, i32 21, i32 22,
+     i32 23, i32 24, i32 25, i32 26, i32 27, i32 28, i32 29, i32 30>
+  %cmp = fcmp oeq <16 x half> %load2, %shuffle
+  %zext = zext <16 x i1> %cmp to <16 x i8>
+  store <16 x i8> %zext, ptr undef, align 16
+  ret void
+}
+
+attributes #1 = { nounwind uwtable "target-cpu"="skx" }