refine comments

Author: abhilash1910

llvm/docs/NVPTXUsage.rst

@@ -575,12 +575,8 @@ prefetched in terms of bytes and it must be a multiple of 16.
 
 * The last argument to these intrinsics is boolean flag indicating
   support for cache_hint. These flag argument must be compile-time
-  constant. The backend looks through this flag and lowers the
-  intrinsic appropriately.
-
-* The Nth argument (denoted by ``i1 %flag_ch``) when set, indicates
-  a valid cache_hint (``i64 %ch``) and generates the ``.L2::cache_hint``
-  variant of the PTX instruction.
+  constant. When set, indicates a valid cache_hint (``i64 %ch``) 
+  and generates the ``.L2::cache_hint`` variant of the PTX instruction.
 
 For more information, refer PTX ISA
 `<https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-cp-async-bulk-prefetch>`_.

llvm/include/llvm/IR/IntrinsicsNVVM.td

@@ -5036,7 +5036,7 @@ def int_nvvm_cp_async_bulk_shared_cta_to_global
 // Intrinsics for Bulk Copy Prefetch L2
 def int_nvvm_cp_async_bulk_prefetch_L2
   : DefaultAttrsIntrinsic<[],
-      [llvm_global_ptr_ty, // src_smem_ptr
+      [llvm_global_ptr_ty, // src_gmem_ptr
        llvm_i32_ty,        // copy_size
        llvm_i64_ty,        // cache_hint
        llvm_i1_ty],        // Flag for cache_hint

llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp

@@ -3115,11 +3115,8 @@ void NVPTXDAGToDAGISel::SelectCpAsyncBulkPrefetchL2(SDNode *N) {
   size_t NumArgs = IsCacheHint ? 3 : 2; // src, size, cache_hint
 
   SDLoc DL(N);
-  SmallVector<SDValue, 8> Ops(N->ops().slice(2, NumArgs));
+  SmallVector<SDValue, 4> Ops(N->ops().slice(2, NumArgs));
   Ops.push_back(N->getOperand(0)); // Chain operand
-  //if (IsCacheHint) {
-  //  Ops.push_back(N->getOperand(2));
-  //}
 
   unsigned Opcode;
   if (IsCacheHint)

llvm/lib/Target/NVPTX/NVPTXIntrinsics.td

@@ -552,12 +552,12 @@ defm CP_ASYNC_BULK_CTA_TO_CLUSTER_SHARED32 : CP_ASYNC_BULK_CTA_TO_CLUSTER<Int32R
 //------------------------------
 multiclass CP_ASYNC_BULK_PREFETCH_INTR {
   defvar prefetch = "cp.async.bulk.prefetch.L2.global";
-  def "": NVPTXInst<(outs),
+  def NAME: NVPTXInst<(outs),
             (ins Int64Regs:$src, Int32Regs:$size),
             !strconcat(prefetch," [$src], $size;"),
             []>,
             Requires<[hasPTX<80>, hasSM<90>]>;
-  def _CH: NVPTXInst<(outs),
+  def NAME # _CH: NVPTXInst<(outs),
                   (ins Int64Regs:$src, Int32Regs:$size, Int64Regs:$ch),
                   !strconcat(prefetch,".L2::cache_hint [$src], $size, $ch;"),
                   []>,