[MLIR][NVVM] Make links for PTX ISA documents

[guraypp]

Summary

This works makes links more readable in NVVM dialect's tablegen file.

mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td

def NVVM_SyncWarpOp :
   NVVM_Op<"bar.warp.sync">,
   Arguments<(ins LLVM_Type:$mask)> {
   string llvmBuilder = [{
       createIntrinsicCall(builder, llvm::Intrinsic::nvvm_bar_warp_sync, {$mask});
   }];
   let assemblyFormat = "$mask attr-dict `:` type($mask)";
 }
 
-// https://docs.nvidia.com/cuda/parallel-thread-execution/#id62
+
 def LoadCacheModifierCA : I32EnumAttrCase<"CA", 0, "ca">;
 def LoadCacheModifierCG : I32EnumAttrCase<"CG", 1, "cg">;
 def LoadCacheModifierCS : I32EnumAttrCase<"CS", 2, "cs">;
 def LoadCacheModifierLU : I32EnumAttrCase<"LU", 3, "lu">;
 def LoadCacheModifierCV : I32EnumAttrCase<"CV", 4, "cv">;
 
 /// Enum attribute of the different kinds.
 def LoadCacheModifierKind : I32EnumAttr<"LoadCacheModifierKind", 
                                 "NVVM load cache modifier kind",
   [LoadCacheModifierCA, LoadCacheModifierCG, LoadCacheModifierCS, 
     LoadCacheModifierLU, LoadCacheModifierCV]> {
   let genSpecializedAttr = 0;
   let cppNamespace = "::mlir::NVVM";
+  let description = [{
+    Enum attribute of the different kinds of cache operators for load instructions.
+
+    [For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/#id62)    
+  }];
 }
 
 def LoadCacheModifierAttr : EnumAttr<NVVM_Dialect, LoadCacheModifierKind, "load_cache_modifier">;
 
 def NVVM_CpAsyncOp : NVVM_Op<"cp.async.shared.global", [DeclareOpInterfaceMethods<BasicPtxBuilderOpInterface>]>,
   Arguments<(ins LLVM_i8Ptr_shared:$dst,
                  LLVM_i8Ptr_global:$src,
                  I32Attr:$size,

 //===----------------------------------------------------------------------===//
 
 def NVVM_WgmmaFenceAlignedOp : NVVM_Op<"wgmma.fence.aligned", 
                     [DeclareOpInterfaceMethods<BasicPtxBuilderOpInterface>]> {
   let arguments = (ins);
   let description = [{
     Enforce an ordering of register accesses between warpgroup level matrix 
     multiplication and other operations. 
-    See for more information:
-    https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-instructions-wgmma-fence
+    
+    [For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-instructions-wgmma-fence)
   }];
   let assemblyFormat = "attr-dict";
   let extraClassDefinition = [{
     std::string $cppClass::getPtx() { return std::string("wgmma.fence.sync.aligned;"); }
   }];
 }
 
 def NVVM_WgmmaGroupSyncAlignedOp : NVVM_Op<"wgmma.commit.group.sync.aligned", 
                     [DeclareOpInterfaceMethods<BasicPtxBuilderOpInterface>]>,
   Arguments<(ins )> {
   let assemblyFormat = "attr-dict";
   let description = [{
     Commits all prior uncommitted warpgroup level matrix multiplication operations.
-    See for more information:
-    https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-instructions-wgmma-commit-group
+    
+    [For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-instructions-wgmma-commit-group)
   }];
   let extraClassDefinition = [{
     std::string $cppClass::getPtx() { return std::string("wgmma.commit_group.sync.aligned;"); }
   }];
 }
 
 def NVVM_WgmmaWaitGroupSyncOp : NVVM_Op<"wgmma.wait.group.sync.aligned", 
                     [DeclareOpInterfaceMethods<BasicPtxBuilderOpInterface>]>{
   let arguments = (ins I32Attr:$group);
   let assemblyFormat = "attr-dict $group";
   let description = [{
     Signal the completion of a preceding warpgroup operation.
-    See for more information:
-    https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-instructions-wgmma-wait-group
+    
+    [For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-instructions-wgmma-wait-group)
   }];
   let extraClassDefinition = [{
     std::string $cppClass::getPtx() { return std::string("wgmma.wait_group.sync.aligned %0;"); }
   }];
 }
 
 /// Enum attribute type for the negating of input operands
 def WGMMAScaleInNeg : I32EnumAttrCase<"neg", -1>;
let description = [{
     |   .m64n224k8 |  .m64n224k16 |            |              | .m64n224k32   |
     |   .m64n232k8 |  .m64n232k16 |            |              | .m64n232k32   |
     |   .m64n240k8 |  .m64n240k16 |            |              | .m64n240k32   |
     |   .m64n248k8 |  .m64n248k16 |            |              | .m64n248k32   |
     |   .m64n256k8 |  .m64n256k16 |            |              | .m64n256k32   |
     |--------------|--------------|------------|--------------|---------------|
     ```
 
-    See for more information:
-    https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-instructions
+    
+    [For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-instructions)
   }];
   
   let hasVerifier = 1;
 
   let extraClassDeclaration = [{
     void getAsmValues(RewriterBase &rewriter, 
         llvm::SmallVectorImpl<std::pair<mlir::Value, mlir::NVVM::PTXRegisterMod>> &asmValues);
   }];