[llvm] [LLVM][NVPTX] Add NVPTX codegen support for clusterlaunchcontrol instruction (PR #134568)

[Alex MacLean via llvm-commits]

================
@@ -6036,6 +6102,60 @@ static void ReplaceINTRINSIC_W_CHAIN(SDNode *N, SelectionDAG &DAG,
   case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x64:
   case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x128:
     return ReplaceTcgen05Ld(N, DAG, Results, /* Offset */ true);
+
+  case Intrinsic::nvvm_clusterlaunchcontrol_query_cancel_get_first_ctaid: {
+    // The intrinsic returns the CTAID of x, y and z dimension as a v4i32 value
+    EVT ResVT = N->getValueType(0);
+    if (!ResVT.isVector())
+      return; // already legalized.
+
+    const unsigned NumElts = ResVT.getVectorNumElements(); // v4i32
+
+    // 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(MVT::Other);
+
+    SDVTList ResVTs = DAG.getVTList(ListVTs);
+
+    SmallVector<SDValue, 8> Ops;
----------------
AlexMaclean wrote:

A `SmallVector` seems overkill given that there will always be exactly 4 operands. Can we instead use an initializer list directly in the `getMemIntrinsicNode` call?

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