diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td --- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td +++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td @@ -514,7 +514,7 @@ 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">; @@ -528,6 +528,11 @@ 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; @@ -1436,8 +1441,8 @@ 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 = [{ @@ -1451,8 +1456,8 @@ 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;"); } @@ -1465,8 +1470,8 @@ 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;"); } @@ -1603,8 +1608,8 @@ |--------------|--------------|------------|--------------|---------------| ``` - 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;