[llvm] 8b2abcc - [NFC][NVPTX] Fix links in NVPTX Usage document (#176482)
via llvm-commits
llvm-commits at lists.llvm.org
Tue Jan 20 08:22:54 PST 2026
Author: Rahul Joshi
Date: 2026-01-20T08:22:49-08:00
New Revision: 8b2abcc766b2da74ea33e9c1f785b71b4fa6448f
URL: https://github.com/llvm/llvm-project/commit/8b2abcc766b2da74ea33e9c1f785b71b4fa6448f
DIFF: https://github.com/llvm/llvm-project/commit/8b2abcc766b2da74ea33e9c1f785b71b4fa6448f.diff
LOG: [NFC][NVPTX] Fix links in NVPTX Usage document (#176482)
Fix links to not render the URL and instead use text with a link.
Added:
Modified:
llvm/docs/NVPTXUsage.rst
Removed:
################################################################################
diff --git a/llvm/docs/NVPTXUsage.rst b/llvm/docs/NVPTXUsage.rst
index 0920d1726c407..0e7e21ad46b8d 100644
--- a/llvm/docs/NVPTXUsage.rst
+++ b/llvm/docs/NVPTXUsage.rst
@@ -21,7 +21,7 @@ LLVM IR.
This document assumes a basic familiarity with CUDA and the PTX
assembly language. Information about the CUDA Driver API and the PTX assembly
language can be found in the `CUDA documentation
- <http://docs.nvidia.com/cuda/index.html>`_.
+ <http://docs.nvidia.com/cuda/index.html>`__.
@@ -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>``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -1005,7 +1009,7 @@ memory accesses across the proxies is to be provided. The only supported value
for the ``size`` operand is ``128`` and must be an immediate. Generic Addressing
is used unconditionally, and the address specified by the operand addr must fall
within the ``.global`` state space. Otherwise, the behavior is undefined. For
-more information, see `PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#parallel-synchronization-and-communication-instructions-membar>`_.
+more information, see `PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#parallel-synchronization-and-communication-instructions-membar>`__.
Address Space Intrinsics
------------------------
@@ -1613,8 +1617,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``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -1638,8 +1642,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``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -1668,8 +1672,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``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -1691,8 +1695,8 @@ memory. The destination has to be in the shared memory of a
diff erent 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``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -1718,8 +1722,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.*``'
^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -1766,8 +1770,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.*``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -1792,8 +1796,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.*``'
^^^^^^^^^^^^^^^^^^^^^^^^
@@ -1832,7 +1836,10 @@ undefined.
%fb = freeze i64 %b ;; freezes undef to stable bit-pattern
;; %fa may compare
diff erent 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``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -1864,8 +1871,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
@@ -1887,8 +1894,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``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -1925,14 +1932,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``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -1964,16 +1971,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``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -2009,16 +2016,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``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -2048,16 +2055,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``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -2083,8 +2090,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``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -2117,16 +2124,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``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -2162,15 +2169,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``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -2211,8 +2218,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``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -2238,8 +2245,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
---------------------
@@ -2272,8 +2279,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``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -2299,8 +2306,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``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -2328,8 +2335,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.*``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -2656,15 +2663,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``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -2693,8 +2700,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``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -2718,8 +2725,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``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -2743,8 +2750,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``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -2771,8 +2778,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``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -2795,8 +2802,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``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -2816,8 +2823,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``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -2840,8 +2847,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``'
^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -2902,11 +2909,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.*``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -2986,8 +2993,8 @@ respectively
1 max
=========== =============
-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.*``'
@@ -3039,8 +3046,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
----------------------
@@ -3713,8 +3720,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
----------------
@@ -3735,7 +3742,7 @@ The library can be found under ``nvvm/libdevice/`` in the CUDA Toolkit and
there is a separate version for each compute architecture.
For a list of all math functions implemented in libdevice, see
-`libdevice Users Guide <http://docs.nvidia.com/cuda/libdevice-users-guide/index.html>`_.
+`libdevice Users Guide <http://docs.nvidia.com/cuda/libdevice-users-guide/index.html>`__.
To accommodate various math-related compiler flags that can affect code
generation of libdevice code, the library code depends on a special LLVM IR
More information about the llvm-commits
mailing list