[llvm] [NVPTX][Docs] [NFC] Update docs on intrinsics (PR #133136)

[via llvm-commits]

llvmbot wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-backend-nvptx

Author: Durgadoss R (durga4github)

<details>
<summary>Changes</summary>

Recently, we have added a set of complex intrinsics on
the TMA, tcgen05, and Cvt family of instructions.

This patch captures the key learnings from our experience
so far and documents them as guidelines for future design.

---
Full diff: https://github.com/llvm/llvm-project/pull/133136.diff


1 Files Affected:

- (modified) llvm/docs/NVPTXUsage.rst (+62) 


``````````diff
diff --git a/llvm/docs/NVPTXUsage.rst b/llvm/docs/NVPTXUsage.rst
index 621879fc5648b..1efa72b649f0d 100644
--- a/llvm/docs/NVPTXUsage.rst
+++ b/llvm/docs/NVPTXUsage.rst
@@ -153,6 +153,68 @@ Example: 64-bit PTX for CUDA Driver API: ``nvptx64-nvidia-cuda``
 NVPTX Intrinsics
 ================
 
+Guidelines on NVPTX Intrinsic design
+------------------------------------
+
+The NVPTX intrinsics are used to model instructions in the PTX ISA.
+While simpler intrinsics can represent certain features effectively,
+more complex instructions like TMA and MMA are not as straightforward
+to model. A single variant of these complex instructions can expand
+into hundreds of intrinsics. Additionally, any expansion in the
+corresponding ISA can exponentially increase these numbers, making it
+difficult to manage them in the IR and backend passes. Therefore,
+a careful design of intrinsic interfaces can ease maintenance and
+contribute to a sustainable, long-term solution.
+
+Below are a set of guidelines that may help in choosing
+an appropriate design for the complex intrinsics:
+
+1. If there are only a few intrinsics, prefer a flat design
+   where the intrinsic name encodes all relevant details,
+   and includes only the arguments used by the actual instruction.
+2. As the number of intrinsics grows, it is desirable to consolidate
+   them. NVPTX uses a 'flags'-based design where each flag argument
+   represents one set of instruction modifiers. These flags are
+   compile-time integer constants.
+
+3. When an intrinsic uses flags, document it with details of the
+   flag usage in the ``NVPTXUsage.rst`` file.
+4. Use i1 for boolean flags and i8 for other flag types.
+5. Annotate all flag arguments with ImmArg<ArgIdx<>>.
+6. Place the flag arguments at the end of the (actual)argument list.
+
+7. Identify the key features of an intrinsic and distinguish between
+   first-order and supplementary information. Typically, encoding the
+   first-order information in the intrinsic name while using flags
+   for supplementary details improves readability.
+   For example:
+
+   i. For MMA intrinsics, 'dense' vs. 'sparse' is a fundamental feature,
+   whereas an optional scaling applied to matrices is relatively secondary.
+
+   ii. For TMAs, the mode of copy (e.g., 'Tile' or 'Im2col') is a first-order
+   information, while features like an optional cache hint tend to be
+   secondary.
+
+8. If there are invalid combinations within a set of modifiers, avoid
+   encoding them as flags, as much as possible. This helps reduce the
+   need for error handling of unsupported cases in the backend.
+   For example, some 'cvt' intrinsics support only a subset of the
+   possible rounding modes; so it is preferable not to encode the
+   rounding modes as flags.
+9. Similarly, when there are invalid combinations across a set of
+   modifiers, avoid encoding them as flags to prevent additional
+   complexity in error handling.
+
+10. Maintain a consistent design within an intrinsic family, including
+    argument ordering as well as the usage and ordering of flags.
+11. When designing an intrinsic corresponding to an instruction or its variant,
+    consider the entire instruction family. This may reveal common features
+    that can be modelled consistently across the family.
+
+In summary, strive to balance the aspects mentioned above, to achieve
+a scalable design with maximum readability.
+
 Reading PTX Special Registers
 -----------------------------
 

``````````

</details>


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