[llvm] [NFC][NVPTX] Fix links in NVPTX Usage document (PR #176482)
Rahul Joshi via llvm-commits
llvm-commits at lists.llvm.org
Fri Jan 16 13:57:52 PST 2026
https://github.com/jurahul created https://github.com/llvm/llvm-project/pull/176482
Fix links to not render the URL and instead use text with a link.
>From 12a244a5251374b436e7b870c28a0c7baee536a7 Mon Sep 17 00:00:00 2001
From: Rahul Joshi <rjoshi at nvidia.com>
Date: Fri, 16 Jan 2026 13:55:28 -0800
Subject: [PATCH] [NFC][NVPTX] Fix links in NVPTX Usage document
Fix links to not render the URL and instead use text with a link.
---
llvm/docs/NVPTXUsage.rst | 189 ++++++++++++++++++++-------------------
1 file changed, 98 insertions(+), 91 deletions(-)
diff --git a/llvm/docs/NVPTXUsage.rst b/llvm/docs/NVPTXUsage.rst
index 3256068078632..77eab72722dfd 100644
--- a/llvm/docs/NVPTXUsage.rst
+++ b/llvm/docs/NVPTXUsage.rst
@@ -431,8 +431,8 @@ actions are performed ``atomically``:
* The mbarrier object transitions to the next phase.
* The pending arrival count is reinitialized to the expected arrival count.
-For more information, refer PTX ISA
-`<https://docs.nvidia.com/cuda/parallel-thread-execution/#parallel-synchronization-and-communication-instructions-mbarrier>`_.
+For more information, refer `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/#parallel-synchronization-and-communication-instructions-mbarrier>`_.
'``llvm.nvvm.mbarrier.init``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -850,8 +850,8 @@ elected thread is captured in the i32 return value. The i1 return
value is set to ``True`` for the leader thread and ``False`` for all
the other threads. Election of a leader thread happens deterministically,
i.e. the same leader thread is elected for the same ``membermask``
-every time. For more information, refer PTX ISA
-`<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-elect-sync>`_.
+every time. For more information, refer `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-elect-sync>`_.
Membar/Fences
-------------
@@ -878,7 +878,8 @@ to operations performed on objects in `shared_cta` space. Likewise, when
`sync_restrict` is restricted to `shared_cluster`, then memory semantics must be
`acquire` and the effect of the fence operation only applies to operations
performed on objects in `shared_cluster` memory space. The scope for both
-operations is `cluster`. For more details, please refer the `PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-membar>`__
+operations is `cluster`. For more details, please refer the `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-membar>`__
'``llvm.nvvm.fence.mbarrier_init.release.cluster``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -897,7 +898,8 @@ Overview:
memory operations for which the fence instruction provides the memory ordering
guarantees. The `mbarrier_init` modifiers restricts the synchronizing effect to
the prior `mbarrier_init` operation executed by the same thread on mbarrier
-objects in `shared_cta` memory space. For more details, please refer the `PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-membar>`__
+objects in `shared_cta` memory space. For more details, please refer the `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-membar>`__
'``llvm.nvvm.fence.proxy.async_generic.acquire/release.sync_restrict``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -914,7 +916,8 @@ Overview:
"""""""""
`nvvm.fence.proxy.async_generic.{semantics}.sync_restrict` are used to establish
-ordering between a prior memory access performed via the `async proxy <https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#proxies>`__
+ordering between a prior memory access performed via the `async proxy
+<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#proxies>`__
and a subsequent memory access performed via the generic proxy.
``nvvm.fence.proxy.async_generic.release.sync_restrict`` can form a release
sequence that synchronizes with an acquire sequence that contains the
@@ -925,7 +928,8 @@ performed on objects in `shared_cta` space. Likewise, when `sync_restrict` is
restricted to `shared_cluster`, then memory semantics must be `acquire` and the
effect of the fence operation only applies to operations performed on objects in
`shared_cluster` memory space. The scope for both operations is `cluster`.
-For more details, please refer the `PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-membar>`__
+For more details, please refer the `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-membar>`__
'``llvm.nvvm.fence.proxy.<proxykind>``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -1513,8 +1517,8 @@ the amount of memory to be copied and it must be a multiple of 16.
presence of a multicast mask (``i16 %mc``) and generates the PTX instruction
with the ``.multicast::cluster`` modifier.
-For more information, refer PTX ISA
-`<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk>`_.
+For more information, refer `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk>`_.
'``llvm.nvvm.cp.async.bulk.global.to.shared.cta``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -1538,8 +1542,8 @@ to be copied and it must be a multiple of 16. The last argument (denoted by
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/index.html#data-movement-and-conversion-instructions-cp-async-bulk>`_.
+For more information, refer `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk>`_.
'``llvm.nvvm.cp.async.bulk.shared.cta.to.global``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -1568,8 +1572,8 @@ chunk of source data is copied to the destination.
When set, it 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/index.html#data-movement-and-conversion-instructions-cp-async-bulk>`_.
+For more information, refer `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk>`_.
'``llvm.nvvm.cp.async.bulk.shared.cta.to.cluster``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -1591,8 +1595,8 @@ memory. The destination has to be in the shared memory of a different CTA within
the cluster. The 32-bit operand ``%size`` specifies the amount of memory to be
copied and it must be a multiple of 16.
-For more information, refer PTX ISA
-`<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk>`_.
+For more information, refer `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk>`_.
'``llvm.nvvm.cp.async.bulk.prefetch.L2``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -1618,8 +1622,8 @@ prefetched in terms of bytes and it must be a multiple of 16.
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>`_.
+For more information, refer `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-cp-async-bulk-prefetch>`_.
'``llvm.nvvm.prefetch.*``'
^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -1666,8 +1670,8 @@ by the '``.level::eviction_priority``' modifier.
* A prefetch into the uniform cache requires a generic address, and no operation
occurs if the address maps to a const, local, or shared memory location.
-For more information, refer to the PTX ISA
-`<https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-prefetch-prefetchu>`_.
+For more information, refer to the `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-prefetch-prefetchu>`_.
'``llvm.nvvm.applypriority.*``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -1692,8 +1696,8 @@ operand size is an integer constant that specifies the amount of data, in bytes,
in the specified cache level on which the priority is to be applied. The only
supported value for the size operand is 128.
-For more information, refer to the PTX ISA
-`<https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-applypriority>`_.
+For more information, refer to the `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-applypriority>`_.
``llvm.nvvm.discard.*``'
^^^^^^^^^^^^^^^^^^^^^^^^
@@ -1732,7 +1736,10 @@ undefined.
%fb = freeze i64 %b ;; freezes undef to stable bit-pattern
;; %fa may compare different to %fb!
-For more information, refer to the `CUDA C++ discard documentation <https://nvidia.github.io/cccl/libcudacxx/extended_api/memory_access_properties/discard_memory.html>`__ and to the `PTX ISA discard documentation <https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-discard>`__ .
+For more information, refer to the `CUDA C++ discard documentation
+<https://nvidia.github.io/cccl/libcudacxx/extended_api/memory_access_properties/discard_memory.html>`__
+and to the `PTX ISA discard documentation
+<https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-discard>`__ .
'``llvm.nvvm.cp.async.bulk.tensor.g2s.tile.[1-5]d``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -1764,8 +1771,8 @@ coordinates specified by the ``i32 %d0 ... i32 %d4`` arguments. In
destination tensor. The first coordinate ``i32 %x0`` denotes the column index
followed by four coordinates indicating the four row-indices. So, this mode
takes a total of 5 coordinates as input arguments. For more information on
-``gather4`` mode, refer PTX ISA
-`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-tiled-scatter4-gather4-modes>`_.
+``gather4`` mode, refer `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-tiled-scatter4-gather4-modes>`_.
* The last three arguments to these intrinsics are flags indicating support for
multicast, cache_hint and cta_group::1/2 modifiers. These flag arguments must
@@ -1787,8 +1794,8 @@ takes a total of 5 coordinates as input arguments. For more information on
``cta_group::1`` and ``cta_group::2`` variants of the PTX instruction
respectively.
-For more information, refer PTX ISA
-`<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk-tensor>`_.
+For more information, refer `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk-tensor>`_.
'``llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.[3-5]d``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -1825,14 +1832,14 @@ are also specified (denoted by ``i16 im2col0...i16 %im2col2``). For the
of the tensor operation. For the ``im2col.w`` and ``im2col.w.128`` mode, the
number of offsets is always 2, denoted by ``i16 %wHalo`` and ``i16 %wOffset``
arguments. For more information on ``im2col.w`` and ``im2col.w.128`` modes,
-refer PTX ISA
-`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-im2col-w-w128-modes>`_.
+refer `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-im2col-w-w128-modes>`_.
The last three arguments to these intrinsics are flags, with the same
functionality as described in the ``tile`` mode intrinsics above.
-For more information, refer PTX ISA
-`<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk-tensor>`_.
+For more information, refer `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk-tensor>`_.
'``llvm.nvvm.cp.async.bulk.tensor.g2s.cta.tile.[1-5]d``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -1864,16 +1871,16 @@ coordinates specified by the ``i32 %d0 ... i32 %d4`` arguments. In
destination tensor. The first coordinate ``i32 %x0`` denotes the column index
followed by four coordinates indicating the four row-indices. So, this mode
takes a total of 5 coordinates as input arguments. For more information on
-``gather4`` mode, refer PTX ISA
-`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-tiled-scatter4-gather4-modes>`_.
+``gather4`` mode, refer `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-tiled-scatter4-gather4-modes>`_.
* The last argument to these intrinsics is a boolean flag indicating support for
cache_hint. This flag argument must be a compile-time constant. When set, it
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/index.html#data-movement-and-conversion-instructions-cp-async-bulk-tensor>`_.
+For more information, refer `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk-tensor>`_.
'``llvm.nvvm.cp.async.bulk.tensor.g2s.cta.im2col.[3-5]d``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -1909,16 +1916,16 @@ also specified (denoted by ``i16 im2col0...i16 %im2col2``). For the ``im2col``
mode, the number of offsets is two less than the number of dimensions of the
tensor operation. For the ``im2col.w`` and ``im2col.w.128`` mode, the number of
offsets is always 2, denoted by ``i16 %wHalo`` and ``i16 %wOffset`` arguments.
-For more information on ``im2col.w`` and ``im2col.w.128`` modes, refer PTX ISA
-`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-im2col-w-w128-modes>`_.
+For more information on ``im2col.w`` and ``im2col.w.128`` modes, refer `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-im2col-w-w128-modes>`_.
* The last argument to these intrinsics is a boolean flag indicating support
for cache_hint. This flag argument must be a compile-time constant. When set,
it 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/index.html#data-movement-and-conversion-instructions-cp-async-bulk-tensor>`_.
+For more information, refer `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk-tensor>`_.
'``llvm.nvvm.cp.async.bulk.tensor.s2g.tile.[1-5]d``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -1948,16 +1955,16 @@ of the tensor data ranges from 1d to 5d with the coordinates specified by the
tensor is divided into four rows in the 2D destination tensor. The first
coordinate ``i32 %x0`` denotes the column index followed by four coordinates
indicating the four row-indices. So, this mode takes a total of 5 coordinates as
-input arguments. For more information on ``scatter4`` mode, refer PTX ISA
-`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-tiled-scatter4-gather4-modes>`_.
+input arguments. For more information on ``scatter4`` mode, refer `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-tiled-scatter4-gather4-modes>`_.
* The last argument to these intrinsics is a boolean flag indicating support for
cache_hint. This flag argument must be a compile-time constant. When set, it
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/index.html#data-movement-and-conversion-instructions-cp-async-bulk-tensor>`_.
+For more information, refer `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk-tensor>`_.
'``llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.[3-5]d``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -1983,8 +1990,8 @@ mode. In this mode, the tensor has to be at least three-dimensional. Unlike the
argument to these intrinsics is a boolean flag, with the same functionality as
described in the ``s2g.tile`` mode intrinsics above.
-For more information, refer PTX ISA
-`<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk-tensor>`_.
+For more information, refer `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk-tensor>`_.
'``llvm.nvvm.cp.async.bulk.tensor.prefetch.tile.[1-5]d``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -2017,16 +2024,16 @@ In ``tile.gather4`` mode, four rows in the 2-dimnesional source tensor are
fetched to the L2 cache. The first coordinate ``i32 %x0`` denotes the column
index followed by four coordinates indicating the four row-indices. So, this
mode takes a total of 5 coordinates as input arguments. For more information
-on ``gather4`` mode, refer PTX ISA
-`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-tiled-scatter4-gather4-modes>`_.
+on ``gather4`` mode, refer `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-tiled-scatter4-gather4-modes>`_.
* The last argument to these intrinsics is a boolean flag indicating support
for cache_hint. This flag argument must be a compile-time constant. When set,
it 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-tensor>`_.
+For more information, refer `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-cp-async-bulk-prefetch-tensor>`_.
'``llvm.nvvm.cp.async.bulk.tensor.prefetch.im2col.[3-5]d``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -2062,15 +2069,15 @@ tensor coordinates, im2col offsets are also specified (denoted by
two less than the number of dimensions of the tensor operation. For the
``im2col.w`` and ``im2col.w.128`` modes, the number of offsets is always 2,
denoted by ``i16 %wHalo`` and ``i16 %wOffset`` arguments. For more information
-on ``im2col.w`` and ``im2col.w.128`` modes, refer PTX ISA
-`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-im2col-w-w128-modes>`_.
+on ``im2col.w`` and ``im2col.w.128`` modes, refer `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-im2col-w-w128-modes>`_.
The last argument to these intrinsics is a boolean flag, with the same
functionality as described in the ``tile`` mode intrinsics above.
-For more information, refer PTX ISA
-`<https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-cp-async-bulk-prefetch-tensor>`_.
+For more information, refer `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-cp-async-bulk-prefetch-tensor>`_.
'``llvm.nvvm.cp.async.bulk.tensor.reduce.[red_op].tile.[1-5]d``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -2111,8 +2118,8 @@ the ``tile.1d`` intrinsics.
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/index.html#data-movement-and-conversion-instructions-cp-reduce-async-bulk-tensor>`_.
+For more information, refer `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-reduce-async-bulk-tensor>`_.
'``llvm.nvvm.cp.async.bulk.tensor.reduce.[red_op].im2col.[3-5]d``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -2138,8 +2145,8 @@ The supported reduction operations supported are the same as the ones in the
tile mode. The last argument to these intrinsics is a boolean flag, with the
same functionality as described in the ``tile`` mode intrinsics above.
-For more information, refer PTX ISA
-`<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-reduce-async-bulk-tensor>`_.
+For more information, refer `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-reduce-async-bulk-tensor>`_.
Warp Group Intrinsics
---------------------
@@ -2172,8 +2179,8 @@ the following locations:
instructions of the same shape in which case an ordering guarantee is
provided by default.
-For more information, refer PTX ISA
-`<https://docs.nvidia.com/cuda/parallel-thread-execution/#asynchronous-warpgroup-level-matrix-instructions-wgmma-fence>`_.
+For more information, refer `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/#asynchronous-warpgroup-level-matrix-instructions-wgmma-fence>`_.
'``llvm.nvvm.wgmma.commit_group.sync.aligned``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -2199,8 +2206,8 @@ wgmma-group.
An executing thread can wait for the completion of all ``wgmma.mma_async``
operations in a wgmma-group by using ``wgmma.wait_group``.
-For more information, refer PTX ISA
-`<https://docs.nvidia.com/cuda/parallel-thread-execution/#asynchronous-warpgroup-level-matrix-instructions-wgmma-commit-group>`_.
+For more information, refer `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/#asynchronous-warpgroup-level-matrix-instructions-wgmma-commit-group>`_.
'``llvm.nvvm.wgmma.wait_group.sync.aligned``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -2228,8 +2235,8 @@ fragments of matrix A of a ``wgmma.mma_async`` instruction without first
performing a ``wgmma.wait_group`` instruction that waits on a wgmma-group
including that ``wgmma.mma_async`` instruction is undefined behavior.
-For more information, refer PTX ISA
-`<https://docs.nvidia.com/cuda/parallel-thread-execution/#asynchronous-warpgroup-level-matrix-instructions-wgmma-wait-group>`_.
+For more information, refer `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/#asynchronous-warpgroup-level-matrix-instructions-wgmma-wait-group>`_.
'``llvm.nvvm.griddepcontrol.*``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -2556,15 +2563,15 @@ The llvm.nvvm.tcgen05.* intrinsics model the TCGEN05 family of instructions
exposed by PTX. These intrinsics use 'Tensor Memory' (henceforth ``tmem``).
NVPTX represents this memory using ``addrspace(6)`` and is always 32-bits.
-For more information, refer to the PTX ISA
-`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-memory>`_.
+For more information, refer to the `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-memory>`_.
The tensor-memory pointers may only be used with the tcgen05 intrinsics.
There are specialized load/store instructions provided (tcgen05.ld/st) to
work with tensor-memory.
-See the PTX ISA for more information on tensor-memory load/store instructions
-`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-memory-and-register-load-store-instructions>`_.
+For more information on tensor-memory load/store instructions, refer to `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-memory-and-register-load-store-instructions>`_.
'``llvm.nvvm.tcgen05.alloc``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -2593,8 +2600,8 @@ explicitly uses shared memory address space for the ``%dst`` operand. The
``.cg1`` and ``.cg2`` variants generate ``cta_group::1`` and ``cta_group::2``
variants of the instruction respectively.
-For more information, refer to the PTX ISA
-`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-memory-allocation-and-management-instructions>`_.
+For more information, refer to the `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-memory-allocation-and-management-instructions>`_.
'``llvm.nvvm.tcgen05.dealloc``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -2618,8 +2625,8 @@ Tensor Memory allocation. The 32-bit operand ``%ncols`` specifies the number
of columns to be de-allocated. The ``.cg1`` and ``.cg2`` variants generate
``cta_group::1`` and ``cta_group::2`` variants of the instruction respectively.
-For more information, refer to the PTX ISA
-`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-memory-allocation-and-management-instructions>`_.
+For more information, refer to the `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-memory-allocation-and-management-instructions>`_.
'``llvm.nvvm.tcgen05.relinq.alloc.permit``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -2643,8 +2650,8 @@ after any of its constituent threads execute
``tcgen05.relinquish_alloc_permit``. The ``.cg1`` and ``.cg2`` variants generate
``cta_group::1`` and ``cta_group::2`` flavors of the instruction respectively.
-For more information, refer to the PTX ISA
-`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-memory-allocation-and-management-instructions>`_.
+For more information, refer to the `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-memory-allocation-and-management-instructions>`_.
'``llvm.nvvm.tcgen05.commit``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -2671,8 +2678,8 @@ multiple CTAs (specified by ``%mc``) in the cluster. The ``.cg1`` and ``.cg2``
variants generate ```cta_group::1`` and ``cta_group::2`` flavors of the
instruction respectively.
-For more information, refer to the PTX ISA
-`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen-async-sync-operations-commit>`_.
+For more information, refer to the `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen-async-sync-operations-commit>`_.
'``llvm.nvvm.tcgen05.wait``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -2695,8 +2702,8 @@ The '``@llvm.nvvm.tcgen05.wait.ld/st``' intrinsics correspond to the
``tcgen05.wait::st`` causes the executing thread to block until all prior
``tcgen05.st`` operations issued by the executing thread have completed.
-For more information, refer to the PTX ISA
-`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-instructions-tcgen05-wait>`_.
+For more information, refer to the `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-instructions-tcgen05-wait>`_.
'``llvm.nvvm.tcgen05.fence``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -2716,8 +2723,8 @@ The '``@llvm.nvvm.tcgen05.fence.*``' intrinsics correspond to the
``tcgen05.fence::{before/after}_thread_sync`` pair of PTX instructions. These
instructions act as code motion fences for asynchronous tcgen05 operations.
-For more information, refer to the PTX ISA
-`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensorcore-5th-generation-instructions-tcgen05-fence>`_.
+For more information, refer to the `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensorcore-5th-generation-instructions-tcgen05-fence>`_.
'``llvm.nvvm.tcgen05.shift``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -2740,8 +2747,8 @@ downwards across all the rows, except the last, by one row. The address operand
``%tmem_addr`` specifies the base address of the matrix in the Tensor Memory
whose rows must be down shifted.
-For more information, refer to the PTX ISA
-`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-instructions-tcgen05-shift>`_.
+For more information, refer to the `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-instructions-tcgen05-shift>`_.
'``llvm.nvvm.tcgen05.cp``'
^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -2802,11 +2809,11 @@ When the source format is ``.b4x16_p64``, a contiguous set of 16 elements of
4-bits each followed by eight bytes of padding (``_p64``) in shared memory is
decompressed into 16 elements of 8-bits (``.b8x16``) each in the Tensor Memory.
-For more information on the decompression schemes, refer to the PTX ISA
-`<https://docs.nvidia.com/cuda/parallel-thread-execution/#optional-decompression>`_.
+For more information on the decompression schemes, refer to the `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/#optional-decompression>`_.
-For more information on the tcgen05.cp instruction, refer to the PTX ISA
-`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-instructions-tcgen05-cp>`_.
+For more information on the tcgen05.cp instruction, refer to the `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-instructions-tcgen05-cp>`_.
'``llvm.nvvm.tcgen05.ld.*``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -2864,8 +2871,8 @@ The last argument `i1 %pack` is a compile-time constant which when set,
indicates that the adjacent columns are packed into a single 32-bit element
during the load.
-For more information, refer to the
-`PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-instructions-tcgen05-ld>`__.
+For more information, refer to the `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-instructions-tcgen05-ld>`__.
'``llvm.nvvm.tcgen05.st.*``'
@@ -2917,8 +2924,8 @@ The last argument `i1 %unpack` is a compile-time constant which when set,
indicates that a 32-bit element in the register to be unpacked into two 16-bit
elements and store them in adjacent columns.
-For more information, refer to the
-`PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-instructions-tcgen05-st>`__.
+For more information, refer to the `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-instructions-tcgen05-st>`__.
tcgen05.mma Intrinsics
----------------------
@@ -3591,8 +3598,8 @@ The '``llvm.nvvm.pm.event.mask``' intrinsic triggers one or more performance
monitor events. Each bit in the 16-bit immediate operand `%mask_val`` controls
an event.
-For more information on the pmevent instructions, refer to the PTX ISA
-`<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#miscellaneous-instructions-pmevent>`_.
+For more information on the pmevent instructions, refer to the `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#miscellaneous-instructions-pmevent>`_.
Other Intrinsics
----------------
More information about the llvm-commits
mailing list