[llvm] [LLVM][NVPTX] Add codegen support for tcgen05.{ld, st} instructions (PR #126740)

[Artem Belevich via llvm-commits]

================
@@ -5345,6 +5655,57 @@ static void ReplaceLoadVector(SDNode *N, SelectionDAG &DAG,
   Results.push_back(LoadChain);
 }
 
+// Lower vector return type of tcgen05.ld intrinsics
+static void ReplaceTcgen05Ld(SDNode *N, SelectionDAG &DAG,
+                             SmallVectorImpl<SDValue> &Results,
+                             bool hasOffset = false) {
+  SDLoc DL(N);
+  EVT ResVT = N->getValueType(0);
+  if (!ResVT.isVector())
+    return; // already legalized.
+
+  const unsigned NumElts = ResVT.getVectorNumElements();
+
+  // Create the return type of the instructions
+  SmallVector<EVT, 5> ListVTs;
+  for (unsigned i = 0; i < NumElts; ++i)
+    ListVTs.push_back(MVT::i32);
+
+  ListVTs.push_back(N->getValueType(1)); // Chain
+
+  SDVTList ResVTs = DAG.getVTList(ListVTs);
+
+  SmallVector<SDValue, 8> Ops;
+  // Add Chain and Intrinsic ID
+  Ops.push_back(N->getOperand(0)); // Chain
+  Ops.push_back(N->getOperand(1)); // Intrinsic ID
+  Ops.push_back(N->getOperand(2)); // taddr
----------------
Artem-B wrote:

`Ops` can be directly initialized with the first three values.

https://github.com/llvm/llvm-project/pull/126740