brintos / llvm-project-archived public Read only

1=============================2User Guide for NVPTX Back-end3=============================4 5.. contents::6   :local:7   :depth: 38 9 10Introduction11============12 13To support GPU programming, the NVPTX back-end supports a subset of LLVM IR14along with a defined set of conventions used to represent GPU programming15concepts. This document provides an overview of the general usage of the back-16end, including a description of the conventions used and the set of accepted17LLVM IR.18 19.. note::20 21   This document assumes a basic familiarity with CUDA and the PTX22   assembly language. Information about the CUDA Driver API and the PTX assembly23   language can be found in the `CUDA documentation24   <http://docs.nvidia.com/cuda/index.html>`_.25 26 27 28Conventions29===========30 31Marking Functions as Kernels32----------------------------33 34In PTX, there are two types of functions: *device functions*, which are only35callable by device code, and *kernel functions*, which are callable by host36code. By default, the back-end will emit device functions. The ``ptx_kernel``37calling convention is used to declare a function as a kernel function.38 39The following example shows a kernel function calling a device function in LLVM40IR. The function ``@my_kernel`` is callable from host code, but ``@my_fmad`` is41not.42 43.. code-block:: llvm44 45    define float @my_fmad(float %x, float %y, float %z) {46      %mul = fmul float %x, %y47      %add = fadd float %mul, %z48      ret float %add49    }50 51    define ptx_kernel void @my_kernel(ptr %ptr) {52      %val = load float, ptr %ptr53      %ret = call float @my_fmad(float %val, float %val, float %val)54      store float %ret, ptr %ptr55      ret void56    }57 58When compiled, the PTX kernel functions are callable by host-side code.59 60 61Parameter Attributes62--------------------63 64``"nvvm.grid_constant"``65    This attribute may be attached to a ``byval`` parameter of a kernel function66    to indicate that the parameter should be lowered as a direct reference to67    the grid-constant memory of the parameter, as opposed to a copy of the68    parameter in local memory. Writing to a grid-constant parameter is69    undefined behavior. Unlike a normal ``byval`` parameter, the address of a70    grid-constant parameter is not unique to a given function invocation but71    instead is shared by all kernels in the grid.72 73.. _nvptx_fnattrs:74 75Function Attributes76-------------------77 78``"nvvm.maxclusterrank"="<n>"``79    This attribute specifies the maximum number of blocks per cluster. Must be 80    non-zero. Only supported for Hopper+.81 82``"nvvm.minctasm"="<n>"``83    This indicates a hint/directive to the compiler/driver, asking it to put at84    least these many CTAs on an SM.85 86``"nvvm.maxnreg"="<n>"``87    This attribute indicates the maximum number of registers to be used for the88    kernel function.89 90``"nvvm.maxntid"="<x>[,<y>[,<z>]]"``91    This attribute declares the maximum number of threads in the thread block92    (CTA). The maximum number of threads is the product of the maximum extent in93    each dimension. Exceeding the maximum number of threads results in a runtime94    error or kernel launch failure.95 96``"nvvm.reqntid"="<x>[,<y>[,<z>]]"``97    This attribute declares the exact number of threads in the thread block98    (CTA). The number of threads is the product of the value in each dimension.99    Specifying a different CTA dimension at launch will result in a runtime 100    error or kernel launch failure.101 102``"nvvm.cluster_dim"="<x>[,<y>[,<z>]]"``103    This attribute declares the number of thread blocks (CTAs) in the cluster.104    The total number of CTAs is the product of the number of CTAs in each 105    dimension. Specifying a different cluster dimension at launch will result in106    a runtime error or kernel launch failure. Only supported for Hopper+.107 108``"nvvm.blocksareclusters"``109    This attribute implies that the grid launch configuration for the corresponding110    kernel function is specifying the number of clusters instead of the number of thread111    blocks. This attribute is only allowed for kernel functions and requires112    ``nvvm.reqntid`` and ``nvvm.cluster_dim`` attributes.113 114.. _address_spaces:115 116Address Spaces117--------------118 119The NVPTX back-end uses the following address space mapping:120 121   ============= ======================122   Address Space Memory Space123   ============= ======================124   0             Generic125   1             Global126   2             Internal Use127   3             Shared128   4             Constant129   5             Local130   7             Shared Cluster131   ============= ======================132 133Every global variable and pointer type is assigned to one of these address134spaces, with 0 being the default address space. Intrinsics are provided which135can be used to convert pointers between the generic and non-generic address136spaces.137 138As an example, the following IR will define an array ``@g`` that resides in139global device memory.140 141.. code-block:: llvm142 143    @g = internal addrspace(1) global [4 x i32] [ i32 0, i32 1, i32 2, i32 3 ]144 145LLVM IR functions can read and write to this array, and host-side code can146copy data to it by name with the CUDA Driver API.147 148Note that since address space 0 is the generic space, it is illegal to have149global variables in address space 0.  Address space 0 is the default address150space in LLVM, so the ``addrspace(N)`` annotation is *required* for global151variables.152 153 154Triples155-------156 157The NVPTX target uses the module triple to select between 32/64-bit code158generation and the driver-compiler interface to use. The triple architecture159can be one of ``nvptx`` (32-bit PTX) or ``nvptx64`` (64-bit PTX). The160operating system should be one of ``cuda`` or ``nvcl``, which determines the161interface used by the generated code to communicate with the driver.  Most162users will want to use ``cuda`` as the operating system, which makes the163generated PTX compatible with the CUDA Driver API.164 165Example: 32-bit PTX for CUDA Driver API: ``nvptx-nvidia-cuda``166 167Example: 64-bit PTX for CUDA Driver API: ``nvptx64-nvidia-cuda``168 169.. _nvptx_arch_hierarchy:170 171NVPTX Architecture Hierarchy and Ordering172=========================================173 174GPU architectures: sm_2Y/sm_3Y/sm_5Y/sm_6Y/sm_7Y/sm_8Y/sm_9Y/sm_10Y/sm_12Y175('Y' represents version within the architecture)176The architectures have name of form ``sm_XYz`` where ``X`` represent the generation177number, ``Y`` represents the version within the architecture, and ``z`` represents178the optional feature suffix.179If ``X1Y1 <= X2Y2``, then GPU capabilities of ``sm_X1Y1`` are included in ``sm_X2Y2``.180For example, take ``sm_90`` (9 represents ``X``, 0 represents ``Y``, and no feature181suffix) and ``sm_103`` architectures (10 represents ``X``, 3 represents ``Y``, and no182feature suffix). Since 90 <= 103, ``sm_90`` is compatible with ``sm_103``.183 184The family-specific variants have ``f`` feature suffix and they follow185following order:186``sm_X{Y2}f > sm_X{Y1}f`` iff ``Y2 > Y1``187``sm_XY{f} > sm_{XY}{}``188 189For example, take ``sm_100f`` (10 represents ``X``, 0 represents ``Y``, and ``f``190represents ``z``) and ``sm_103f`` (10 represents ``X``, 3 represents ``Y``, and ``f``191represents ``z``) architecture variants. Since ``Y1 < Y2``, ``sm_100f`` is compatible with192``sm_103f``. Similarly based on the second rule, ``sm_90`` is compatible with ``sm_103f``.193 194Some counter examples, take ``sm_100f`` and ``sm_120f`` (12 represents ``X``, 0195represents ``Y``, and ``f`` represents ``z``) architecture variants. Since both196belongs to different family i.e. ``X1 != X2``, ``sm_100f`` is not compatible with197``sm_120f``.198 199The architecture-specific variants have ``a`` feature suffix and they follow200following order:201``sm_XY{a} > sm_XY{f} > sm_{XY}{}``202 203For example, take ``sm_103a`` (10 represents ``X``, 3 represents ``Y``, and ``a``204represents ``z``), ``sm_103f``, and ``sm_103`` architecture variants. The ``sm_103`` is205compatible with ``sm_103a`` and ``sm_103f``, and ``sm_103f`` is compatible with ``sm_103a``.206 207Encoding := Arch * 10 + 2 (for 'f') + 1 (for 'a')208Arch := X * 10 + Y209 210For example, ``sm_103f`` is encoded as 1032 (103 * 10 + 2) and ``sm_103a`` is211encoded as 1033 (103 * 10 + 2 + 1).212 213This encoding allows simple partial ordering of the architectures.214 215* Compare Family and Arch by dividing FullSMVersion by 100 and 10216  respectively before the comparison.217* Compare within the family by comparing FullSMVersion, given both belongs to218  the same family.219* Detect ``a`` variants by checking FullSMVersion & 1.220 221.. _nvptx_intrinsics:222 223NVPTX Intrinsics224================225 226Reading PTX Special Registers227-----------------------------228 229'``llvm.nvvm.read.ptx.sreg.*``'230^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^231 232Syntax:233"""""""234 235.. code-block:: llvm236 237    declare i32 @llvm.nvvm.read.ptx.sreg.tid.x()238    declare i32 @llvm.nvvm.read.ptx.sreg.tid.y()239    declare i32 @llvm.nvvm.read.ptx.sreg.tid.z()240    declare i32 @llvm.nvvm.read.ptx.sreg.ntid.x()241    declare i32 @llvm.nvvm.read.ptx.sreg.ntid.y()242    declare i32 @llvm.nvvm.read.ptx.sreg.ntid.z()243    declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()244    declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.y()245    declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.z()246    declare i32 @llvm.nvvm.read.ptx.sreg.nctaid.x()247    declare i32 @llvm.nvvm.read.ptx.sreg.nctaid.y()248    declare i32 @llvm.nvvm.read.ptx.sreg.nctaid.z()249    declare i32 @llvm.nvvm.read.ptx.sreg.warpsize()250 251Overview:252"""""""""253 254The '``@llvm.nvvm.read.ptx.sreg.*``' intrinsics provide access to the PTX255special registers, in particular the kernel launch bounds.  These registers256map in the following way to CUDA builtins:257 258   ============ =====================================259   CUDA Builtin PTX Special Register Intrinsic260   ============ =====================================261   ``threadId`` ``@llvm.nvvm.read.ptx.sreg.tid.*``262   ``blockIdx`` ``@llvm.nvvm.read.ptx.sreg.ctaid.*``263   ``blockDim`` ``@llvm.nvvm.read.ptx.sreg.ntid.*``264   ``gridDim``  ``@llvm.nvvm.read.ptx.sreg.nctaid.*``265   ============ =====================================266 267 268Barriers269--------270 271'``llvm.nvvm.barrier.cta.*``'272^^^^^^^^^^^^^^^^^^^^^^^^^^^^^273 274Syntax:275"""""""276 277.. code-block:: llvm278 279  declare void @llvm.nvvm.barrier.cta.sync.count(i32 %id, i32 %n)280  declare void @llvm.nvvm.barrier.cta.sync.all(i32 %id)281  declare void @llvm.nvvm.barrier.cta.arrive.count(i32 %id, i32 %n)282 283  declare void @llvm.nvvm.barrier.cta.sync.aligned.count(i32 %id, i32 %n)284  declare void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 %id)285  declare void @llvm.nvvm.barrier.cta.arrive.aligned.count(i32 %id, i32 %n)286 287Overview:288"""""""""289 290The '``@llvm.nvvm.barrier.cta.*``' family of intrinsics perform barrier291synchronization and communication within a CTA. They can be used by the threads292within the CTA for synchronization and communication.293 294Semantics:295""""""""""296 297Operand %id specifies a logical barrier resource and must fall within the range2980 through 15. When present, operand %n specifies the number of threads299participating in the barrier. When specifying a thread count, the value must be300a multiple of the warp size. With the '``@llvm.nvvm.barrier.cta.sync.*``'301variants, the '``.all``' suffix indicates that all threads in the CTA should302participate in the barrier while the '``.count``' suffix indicates that only303the threads specified by the %n operand should participate in the barrier.304 305All forms of the '``@llvm.nvvm.barrier.cta.*``' intrinsic cause the executing306thread to wait for all non-exited threads from its warp and then marks the307warp's arrival at the barrier. In addition to signaling its arrival at the 308barrier, the '``@llvm.nvvm.barrier.cta.sync.*``' intrinsics cause the executing309thread to wait for non-exited threads of all other warps participating in the310barrier to arrive. On the other hand, the '``@llvm.nvvm.barrier.cta.arrive.*``'311intrinsic does not cause the executing thread to wait for threads of other312participating warps.313 314When a barrier completes, the waiting threads are restarted without delay,315and the barrier is reinitialized so that it can be immediately reused.316 317The '``@llvm.nvvm.barrier.cta.*``' intrinsic has an optional '``.aligned``'318modifier to indicate textual alignment of the barrier. When specified, it319indicates that all threads in the CTA will execute the same320'``@llvm.nvvm.barrier.cta.*``' instruction. In conditionally executed code, an321aligned '``@llvm.nvvm.barrier.cta.*``' instruction should only be used if it is322known that all threads in the CTA evaluate the condition identically, otherwise323behavior is undefined.324 325MBarrier family of Intrinsics326-----------------------------327 328Overview:329^^^^^^^^^330 331An ``mbarrier`` is a barrier created in shared memory that supports:332 333* Synchronizing any subset of threads within a CTA.334* One-way synchronization of threads across CTAs of a cluster.335  Threads can perform only ``arrive`` operations but not ``*_wait`` on an336  mbarrier located in shared::cluster space.337* Waiting for completion of asynchronous memory operations initiated by a338  thread and making them visible to other threads.339 340Unlike ``bar{.cta}/barrier{.cta}`` instructions which can access a limited341number of barriers per CTA, ``mbarrier`` objects are user-defined and are342only limited by the total shared memory size available.343 344An mbarrier object is an opaque object in shared memory with an345alignment of 8-bytes. It keeps track of:346 347* Current phase of the mbarrier object348* Count of pending arrivals for the current phase of the mbarrier object349* Count of expected arrivals for the next phase of the mbarrier object350* Count of pending asynchronous memory operations (or transactions)351  tracked by the current phase of the mbarrier object. This is also352  referred to as ``tx-count``. The unit of ``tx-count`` is specified353  by the asynchronous memory operation (for example,354  ``llvm.nvvm.cp.async.bulk.tensor.g2s.*``).355 356The ``phase`` of an mbarrier object is the number of times the mbarrier357object has been used to synchronize threads/track async operations.358In each phase, threads perform:359 360* arrive/expect-tx/complete-tx operations to progress the current phase.361* test_wait/try_wait operations to check for completion of the current phase.362 363An mbarrier object completes the current phase when:364 365* The count of the pending arrivals has reached zero AND366* The tx-count has reached zero.367 368When an mbarrier object completes the current phase, below369actions are performed ``atomically``:370 371* The mbarrier object transitions to the next phase.372* The pending arrival count is reinitialized to the expected arrival count.373 374For more information, refer PTX ISA375`<https://docs.nvidia.com/cuda/parallel-thread-execution/#parallel-synchronization-and-communication-instructions-mbarrier>`_.376 377'``llvm.nvvm.mbarrier.init``'378^^^^^^^^^^^^^^^^^^^^^^^^^^^^^379 380Syntax:381"""""""382 383.. code-block:: llvm384 385  declare void @llvm.nvvm.mbarrier.init(ptr %addr, i32 %count)386  declare void @llvm.nvvm.mbarrier.init.shared(ptr addrspace(3) %addr, i32 %count)387 388Overview:389"""""""""390 391The '``@llvm.nvvm.mbarrier.init.*``' intrinsics are used to initialize392an mbarrier object located at ``addr`` with the value ``count``.393``count`` is a 32-bit unsigned integer value and must be within394the range [1...2^20-1]. During initialization:395 396* The tx-count and the current phase of the mbarrier object are set to 0.397* The expected and pending arrival counts are set to ``count``.398 399Semantics:400""""""""""401 402The ``.shared`` variant explicitly uses shared memory address space for403the ``addr`` operand. If the ``addr`` does not fall within the404shared::cta space, then the behavior of this intrinsic is undefined.405Performing ``mbarrier.init`` on a valid mbarrier object is undefined;406use ``mbarrier.inval`` before reusing the memory for another mbarrier407or any other purpose.408 409'``llvm.nvvm.mbarrier.inval``'410^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^411 412Syntax:413"""""""414 415.. code-block:: llvm416 417  declare void @llvm.nvvm.mbarrier.inval(ptr %addr)418  declare void @llvm.nvvm.mbarrier.inval.shared(ptr addrspace(3) %addr)419 420Overview:421"""""""""422 423The '``@llvm.nvvm.mbarrier.inval.*``' intrinsics invalidate the mbarrier424object at the address specified by ``addr``.425 426Semantics:427""""""""""428 429The ``.shared`` variant explicitly uses shared memory address space for430the ``addr`` operand. If the ``addr`` does not fall within the431shared::cta space, then the behavior of this intrinsic is undefined.432It is expected that ``addr`` was previously initialized using433``mbarrier.init``; otherwise, the behavior is undefined.434 435'``llvm.nvvm.mbarrier.expect.tx``'436^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^437 438Syntax:439"""""""440 441.. code-block:: llvm442 443  declare void @llvm.nvvm.mbarrier.expect.tx.scope.cta.space.cta(ptr addrspace(3) %addr, i32 %tx_count)444  declare void @llvm.nvvm.mbarrier.expect.tx.scope.cluster.space.cta(ptr addrspace(3) %addr, i32 %tx_count)445  declare void @llvm.nvvm.mbarrier.expect.tx.scope.cta.space.cluster(ptr addrspace(7) %addr, i32 %tx_count)446  declare void @llvm.nvvm.mbarrier.expect.tx.scope.cluster.space.cluster(ptr addrspace(7) %addr, i32 %tx_count)447 448Overview:449"""""""""450 451The '``@llvm.nvvm.mbarrier.expect.tx.*``' intrinsics increase the transaction452count of the mbarrier object at ``%addr`` by ``%tx_count``. The ``%tx_count``453is a 32-bit unsigned integer value.454 455Semantics:456""""""""""457 458The ``.space.{cta/cluster}`` indicates the address space where the mbarrier459object resides.460 461The ``.scope.{cta/cluster}`` denotes the set of threads that can directly462observe the synchronizing effect of the mbarrier operation. When scope is463"cta", all threads executing in the same CTA (as the current thread) can464directly observe the effect of the ``expect.tx`` operation. Similarly,465when scope is "cluster", all threads executing in the same Cluster466(as the current thread) can directly observe the effect of the operation.467 468If the ``addr`` does not fall within shared::cta or shared::cluster space,469then the behavior of this intrinsic is undefined. This intrinsic has470``relaxed`` semantics and hence does not provide any memory ordering471or visibility guarantees.472 473'``llvm.nvvm.mbarrier.complete.tx``'474^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^475 476Syntax:477"""""""478 479.. code-block:: llvm480 481  declare void @llvm.nvvm.mbarrier.complete.tx.scope.cta.space.cta(ptr addrspace(3) %addr, i32 %tx_count)482  declare void @llvm.nvvm.mbarrier.complete.tx.scope.cluster.space.cta(ptr addrspace(3) %addr, i32 %tx_count)483  declare void @llvm.nvvm.mbarrier.complete.tx.scope.cta.space.cluster(ptr addrspace(7) %addr, i32 %tx_count)484  declare void @llvm.nvvm.mbarrier.complete.tx.scope.cluster.space.cluster(ptr addrspace(7) %addr, i32 %tx_count)485 486Overview:487"""""""""488 489The '``@llvm.nvvm.mbarrier.complete.tx.*``' intrinsics decrease the transaction490count of the mbarrier object at ``%addr`` by ``%tx_count``. The ``%tx_count``491is a 32-bit unsigned integer value. As a result of this decrement,492the mbarrier can potentially complete its current phase and transition493to the next phase.494 495Semantics:496""""""""""497 498The semantics of these intrinsics are identical to those of the499``llvm.nvvm.mbarrier.expect.tx.*`` intrinsics described above.500 501'``llvm.nvvm.mbarrier.arrive``'502^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^503 504Syntax:505"""""""506 507.. code-block:: llvm508 509  declare i64  @llvm.nvvm.mbarrier.arrive.scope.cta.space.cta(ptr addrspace(3) %addr, i32 %count)510  declare i64  @llvm.nvvm.mbarrier.arrive.scope.cluster.space.cta(ptr addrspace(3) %addr, i32 %count)511  declare void @llvm.nvvm.mbarrier.arrive.scope.cta.space.cluster(ptr addrspace(7) %addr, i32 %count)512  declare void @llvm.nvvm.mbarrier.arrive.scope.cluster.space.cluster(ptr addrspace(7) %addr, i32 %count)513 514  declare i64  @llvm.nvvm.mbarrier.arrive.relaxed.scope.cta.space.cta(ptr addrspace(3) %addr, i32 %count)515  declare i64  @llvm.nvvm.mbarrier.arrive.relaxed.scope.cluster.space.cta(ptr addrspace(3) %addr, i32 %count)516  declare void @llvm.nvvm.mbarrier.arrive.relaxed.scope.cta.space.cluster(ptr addrspace(7) %addr, i32 %count)517  declare void @llvm.nvvm.mbarrier.arrive.relaxed.scope.cluster.space.cluster(ptr addrspace(7) %addr, i32 %count)518 519Overview:520"""""""""521 522The ``@llvm.nvvm.mbarrier.arrive.*`` intrinsics signal the arrival of the523executing thread or completion of an asynchronous instruction associated with524an arrive operation on the mbarrier object at ``%addr``. This operation525decrements the pending arrival count by ``%count``, a 32-bit unsigned integer,526potentially completing the current phase and triggering a transition to the527next phase.528 529Semantics:530""""""""""531 532The ``.space.{cta/cluster}`` indicates the address space where the mbarrier533object resides. When the mbarrier is in shared::cta space, the intrinsics534return an opaque 64-bit value capturing the phase of the mbarrier object535_prior_ to this arrive operation. This value can be used with a try_wait536or test_wait operation to check for the completion of the mbarrier.537 538The ``.scope.{cta/cluster}`` denotes the set of threads that can directly539observe the synchronizing effect of the mbarrier operation. When scope is540"cta", all threads executing in the same CTA (as the current thread) can541directly observe the effect of the ``arrive`` operation. Similarly,542when scope is "cluster", all threads executing in the same Cluster543(as the current thread) can directly observe the effect of the operation.544 545If the ``addr`` does not fall within shared::cta or shared::cluster space,546then the behavior of this intrinsic is undefined.547 548These intrinsics have ``release`` semantics by default. The release semantics549ensure ordering of operations that occur in program order _before_ this arrive550instruction, making their effects visible to subsequent operations in other551threads of the CTA (or cluster, depending on scope). Threads performing552corresponding acquire operations (such as mbarrier.test.wait) synchronize553with this release. The ``relaxed`` variants of these intrinsics do not554provide any memory ordering or visibility guarantees.555 556'``llvm.nvvm.mbarrier.arrive.expect.tx``'557^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^558 559Syntax:560"""""""561 562.. code-block:: llvm563 564  declare i64  @llvm.nvvm.mbarrier.arrive.expect.tx.scope.cta.space.cta(ptr addrspace(3) %addr, i32 %tx_count)565  declare i64  @llvm.nvvm.mbarrier.arrive.expect.tx.scope.cluster.space.cta(ptr addrspace(3) %addr, i32 %tx_count)566  declare void @llvm.nvvm.mbarrier.arrive.expect.tx.scope.cta.space.cluster(ptr addrspace(7) %addr, i32 %tx_count)567  declare void @llvm.nvvm.mbarrier.arrive.expect.tx.scope.cluster.space.cluster(ptr addrspace(7) %addr, i32 %tx_count)568 569  declare i64  @llvm.nvvm.mbarrier.arrive.expect.tx.relaxed.scope.cta.space.cta(ptr addrspace(3) %addr, i32 %tx_count)570  declare i64  @llvm.nvvm.mbarrier.arrive.expect.tx.relaxed.scope.cluster.space.cta(ptr addrspace(3) %addr, i32 %tx_count)571  declare void @llvm.nvvm.mbarrier.arrive.expect.tx.relaxed.scope.cta.space.cluster(ptr addrspace(7) %addr, i32 %tx_count)572  declare void @llvm.nvvm.mbarrier.arrive.expect.tx.relaxed.scope.cluster.space.cluster(ptr addrspace(7) %addr, i32 %tx_count)573 574Overview:575"""""""""576 577The ``@llvm.nvvm.mbarrier.arrive.expect.tx.*`` intrinsics are similar to578the ``@llvm.nvvm.mbarrier.arrive`` intrinsics except that they also579perform an ``expect-tx`` operation _prior_ to the ``arrive`` operation.580The ``%tx_count`` specifies the transaction count for the ``expect-tx``581operation and the count for the ``arrive`` operation is assumed to be 1.582 583Semantics:584""""""""""585 586The semantics of these intrinsics are identical to those of the587``llvm.nvvm.mbarrier.arrive.*`` intrinsics described above.588 589'``llvm.nvvm.mbarrier.arrive.drop``'590^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^591 592Syntax:593"""""""594 595.. code-block:: llvm596 597  declare i64  @llvm.nvvm.mbarrier.arrive.drop.scope.cta.space.cta(ptr addrspace(3) %addr, i32 %count)598  declare i64  @llvm.nvvm.mbarrier.arrive.drop.scope.cluster.space.cta(ptr addrspace(3) %addr, i32 %count)599  declare void @llvm.nvvm.mbarrier.arrive.drop.scope.cta.space.cluster(ptr addrspace(7) %addr, i32 %count)600  declare void @llvm.nvvm.mbarrier.arrive.drop.scope.cluster.space.cluster(ptr addrspace(7) %addr, i32 %count)601 602  declare i64  @llvm.nvvm.mbarrier.arrive.drop.relaxed.scope.cta.space.cta(ptr addrspace(3) %addr, i32 %count)603  declare i64  @llvm.nvvm.mbarrier.arrive.drop.relaxed.scope.cluster.space.cta(ptr addrspace(3) %addr, i32 %count)604  declare void @llvm.nvvm.mbarrier.arrive.drop.relaxed.scope.cta.space.cluster(ptr addrspace(7) %addr, i32 %count)605  declare void @llvm.nvvm.mbarrier.arrive.drop.relaxed.scope.cluster.space.cluster(ptr addrspace(7) %addr, i32 %count)606 607Overview:608"""""""""609 610The ``@llvm.nvvm.mbarrier.arrive.drop.*`` intrinsics decrement the611expected arrival count of the mbarrier object at ``%addr`` by612``%count`` and then perform an ``arrive`` operation with ``%count``.613The ``%count`` is a 32-bit integer.614 615Semantics:616""""""""""617 618The semantics of these intrinsics are identical to those of the619``llvm.nvvm.mbarrier.arrive.*`` intrinsics described above.620 621'``llvm.nvvm.mbarrier.arrive.drop.expect.tx``'622^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^623 624Syntax:625"""""""626 627.. code-block:: llvm628 629  declare i64  @llvm.nvvm.mbarrier.arrive.drop.expect.tx.scope.cta.space.cta(ptr addrspace(3) %addr, i32 %tx_count)630  declare i64  @llvm.nvvm.mbarrier.arrive.drop.expect.tx.scope.cluster.space.cta(ptr addrspace(3) %addr, i32 %tx_count)631  declare void @llvm.nvvm.mbarrier.arrive.drop.expect.tx.scope.cta.space.cluster(ptr addrspace(7) %addr, i32 %tx_count)632  declare void @llvm.nvvm.mbarrier.arrive.drop.expect.tx.scope.cluster.space.cluster(ptr addrspace(7) %addr, i32 %tx_count)633 634  declare i64  @llvm.nvvm.mbarrier.arrive.drop.expect.tx.relaxed.scope.cta.space.cta(ptr addrspace(3) %addr, i32 %tx_count)635  declare i64  @llvm.nvvm.mbarrier.arrive.drop.expect.tx.relaxed.scope.cluster.space.cta(ptr addrspace(3) %addr, i32 %tx_count)636  declare void @llvm.nvvm.mbarrier.arrive.drop.expect.tx.relaxed.scope.cta.space.cluster(ptr addrspace(7) %addr, i32 %tx_count)637  declare void @llvm.nvvm.mbarrier.arrive.drop.expect.tx.relaxed.scope.cluster.space.cluster(ptr addrspace(7) %addr, i32 %tx_count)638 639Overview:640"""""""""641 642The ``@llvm.nvvm.mbarrier.arrive.drop.expect.tx.*`` intrinsics perform643the below operations on the mbarrier located at ``%addr``.644 645* Perform an ``expect-tx`` operation i.e. increase the transaction count646  of the mbarrier by ``%tx_count``, a 32-bit unsigned integer value.647* Decrement the expected arrival count of the mbarrier by 1.648* Perform an ``arrive`` operation on the mbarrier with a value of 1.649 650Semantics:651""""""""""652 653The semantics of these intrinsics are identical to those of the654``llvm.nvvm.mbarrier.arrive.*`` intrinsics described above.655 656'``llvm.nvvm.mbarrier.test.wait``'657^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^658 659Syntax:660"""""""661 662.. code-block:: llvm663 664  declare i1 @llvm.nvvm.mbarrier.test.wait.scope.cta.space.cta(ptr addrspace(3) %addr, i64 %state)665  declare i1 @llvm.nvvm.mbarrier.test.wait.scope.cluster.space.cta(ptr addrspace(3) %addr, i64 %state)666  declare i1 @llvm.nvvm.mbarrier.test.wait.parity.scope.cta.space.cta(ptr addrspace(3) %addr, i32 %phase)667  declare i1 @llvm.nvvm.mbarrier.test.wait.parity.scope.cluster.space.cta(ptr addrspace(3) %addr, i32 %phase)668 669  declare i1 @llvm.nvvm.mbarrier.test.wait.relaxed.scope.cta.space.cta(ptr addrspace(3) %addr, i64 %state)670  declare i1 @llvm.nvvm.mbarrier.test.wait.relaxed.scope.cluster.space.cta(ptr addrspace(3) %addr, i64 %state)671  declare i1 @llvm.nvvm.mbarrier.test.wait.parity.relaxed.scope.cta.space.cta(ptr addrspace(3) %addr, i32 %phase)672  declare i1 @llvm.nvvm.mbarrier.test.wait.parity.relaxed.scope.cluster.space.cta(ptr addrspace(3) %addr, i32 %phase)673 674Overview:675"""""""""676 677The ``@llvm.nvvm.mbarrier.test.wait.*`` intrinsics test for the completion678of the current or the immediately preceding phase of an mbarrier object at679``%addr``. The test for completion can be done with either the ``state`` or680the ``phase-parity`` of the mbarrier object.681 682* When done through the ``i64 %state`` operand, the state must be683  returned by an ``llvm.nvvm.mbarrier.arrive.*`` on the _same_684  mbarrier object.685* The ``.parity`` variant of these intrinsics test for completion686  of the phase indicated by the operand ``i32 %phase``, which is687  the integer parity of either the current phase or the immediately688  preceding phase of the mbarrier object. An even phase has integer689  parity 0 and an odd phase has integer parity of 1. So the valid690  values for phase-parity are 0 and 1.691 692Semantics:693""""""""""694 695The ``.scope.{cta/cluster}`` denotes the set of threads that the696test_wait operation can directly synchronize with.697 698If the ``addr`` does not fall within shared::cta space, then the699the behavior of this intrinsic is undefined.700 701These intrinsics have ``acquire`` semantics by default. This acquire702pattern establishes memory ordering for operations occurring in program703order after this ``test_wait`` instruction by making operations from704other threads in the CTA (or cluster, depending on scope) visible to705subsequent operations in the current thread. When this wait completes,706it synchronizes with the corresponding release pattern from the707``mbarrier.arrive`` operation. The ``relaxed`` variants of these intrinsics708do not provide any memory ordering or visibility guarantees.709 710This ``test.wait`` intrinsic is non-blocking and immediately returns711the completion status without suspending the executing thread.712 713The boolean return value indicates:714 715* True: The immediately preceding phase has completed716* False: The current phase is still incomplete717 718When this wait returns true, the following ordering guarantees hold:719 720* All memory accesses (except async operations) requested prior to721  ``mbarrier.arrive`` having release semantics by participating722  threads of a CTA (or cluster, depending on scope) are visible to723  the executing thread.724* All ``cp.async`` operations requested prior to ``cp.async.mbarrier.arrive``725  by participating threads of a CTA are visible to the executing thread.726* All ``cp.async.bulk`` operations using the same mbarrier object requested727  prior to ``mbarrier.arrive`` having release semantics by participating CTA728  threads are visible to the executing thread.729* Memory accesses requested after this wait are not visible to memory730  accesses performed prior to ``mbarrier.arrive`` by other participating731  threads.732* No ordering guarantee exists for memory accesses by the same thread733  between an ``mbarrier.arrive`` and this wait.734 735'``llvm.nvvm.mbarrier.try.wait``'736^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^737 738Syntax:739"""""""740 741.. code-block:: llvm742 743  declare i1 @llvm.nvvm.mbarrier.try.wait{.relaxed}.scope.cta.space.cta(ptr addrspace(3) %addr, i64 %state)744  declare i1 @llvm.nvvm.mbarrier.try.wait{.relaxed}.scope.cluster.space.cta(ptr addrspace(3) %addr, i64 %state)745 746  declare i1 @llvm.nvvm.mbarrier.try.wait.parity{.relaxed}.scope.cta.space.cta(ptr addrspace(3) %addr, i32 %phase)747  declare i1 @llvm.nvvm.mbarrier.try.wait.parity{.relaxed}.scope.cluster.space.cta(ptr addrspace(3) %addr, i32 %phase)748 749  declare i1 @llvm.nvvm.mbarrier.try.wait.tl{.relaxed}.scope.cta.space.cta(ptr addrspace(3) %addr, i64 %state, i32 %timelimit)750  declare i1 @llvm.nvvm.mbarrier.try.wait.tl{.relaxed}.scope.cluster.space.cta(ptr addrspace(3) %addr, i64 %state, i32 %timelimit)751 752  declare i1 @llvm.nvvm.mbarrier.try.wait.parity.tl{.relaxed}.scope.cta.space.cta(ptr addrspace(3) %addr, i32 %phase, i32 %timelimit)753  declare i1 @llvm.nvvm.mbarrier.try.wait.parity.tl{.relaxed}.scope.cluster.space.cta(ptr addrspace(3) %addr, i32 %phase, i32 %timelimit)754 755Overview:756"""""""""757 758The ``@llvm.nvvm.mbarrier.try.wait.*`` intrinsics test for the completion of759the current or immediately preceding phase of an mbarrier object at ``%addr``.760Unlike the ``test.wait`` intrinsics, which perform a non-blocking test, these761intrinsics may block the executing thread until the specified phase completes762or a system-dependent time limit expires. Suspended threads resume execution763when the phase completes or the time limit elapses. This time limit is764configurable through the ``.tl`` variants of these intrinsics, where the765``%timelimit`` operand (an unsigned integer) specifies the limit in766nanoseconds. Other semantics are identical to those of the ``test.wait``767intrinsics described above.768 769Electing a thread770-----------------771 772'``llvm.nvvm.elect.sync``'773^^^^^^^^^^^^^^^^^^^^^^^^^^774 775Syntax:776"""""""777 778.. code-block:: llvm779 780  declare {i32, i1} @llvm.nvvm.elect.sync(i32 %membermask)781 782Overview:783"""""""""784 785The '``@llvm.nvvm.elect.sync``' intrinsic generates the ``elect.sync``786PTX instruction, which elects one predicated active leader thread from787a set of threads specified by ``membermask``. The behavior is undefined788if the executing thread is not in ``membermask``. The laneid of the789elected thread is captured in the i32 return value. The i1 return790value is set to ``True`` for the leader thread and ``False`` for all791the other threads. Election of a leader thread happens deterministically,792i.e. the same leader thread is elected for the same ``membermask``793every time. For more information, refer PTX ISA794`<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-elect-sync>`_.795 796Membar/Fences797-------------798 799'``llvm.nvvm.fence.acquire/release.sync_restrict.*``'800^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^801 802Syntax:803"""""""804 805.. code-block:: llvm806 807  declare void @llvm.nvvm.fence.acquire.sync_restrict.space.cluster.scope.cluster()808  declare void @llvm.nvvm.fence.release.sync_restrict.space.cta.scope.cluster()809 810Overview:811"""""""""812 813The `nvvm.fence.{semantics}.sync_restrict.*` restrict the class of memory814operations for which the fence instruction provides the memory ordering guarantees.815When `.sync_restrict` is restricted to `shared_cta`, then memory semantics must816be `release` and the effect of the fence operation only applies to operations817performed on objects in `shared_cta` space. Likewise, when `sync_restrict` is818restricted to `shared_cluster`, then memory semantics must be `acquire` and the819effect of the fence operation only applies to operations performed on objects in820`shared_cluster` memory space. The scope for both operations is `cluster`. For more details,821please refer the `PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-membar>`__822 823'``llvm.nvvm.fence.mbarrier_init.release.cluster``'824^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^825 826Syntax:827"""""""828 829.. code-block:: llvm830 831  declare void @llvm.nvvm.fence.mbarrier_init.release.cluster()832 833Overview:834"""""""""835 836`nvvm.fence.mbarrier_init.release.cluster` intrinsic restrict the class of837memory operations for which the fence instruction provides the memory ordering838guarantees. The `mbarrier_init` modifiers restricts the synchronizing effect to839the prior `mbarrier_init` operation executed by the same thread on mbarrier objects840in `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>`__841 842'``llvm.nvvm.fence.proxy.async_generic.acquire/release.sync_restrict``'843^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^844 845Syntax:846"""""""847 848.. code-block:: llvm849 850  declare void @llvm.nvvm.fence.proxy.async.generic.acquire.sync_restrict.space.cluster.scope.cluster()851  declare void @llvm.nvvm.fence.proxy.async.generic.release.sync_restrict.space.cta.scope.cluster()852 853Overview:854"""""""""855 856`nvvm.fence.proxy.async_generic.{semantics}.sync_restrict` are used to establish857ordering between a prior memory access performed via the `async proxy<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#proxies>__`858and a subsequent memory access performed via the generic proxy.859``nvvm.fence.proxy.async_generic.release.sync_restrict`` can form a release860sequence that synchronizes with an acquire sequence that contains the861``nvvm.fence.proxy.async_generic.acquire.sync_restrict`` proxy fence. When862`.sync_restrict` is restricted to `shared_cta`, then memory semantics must863be `release` and the effect of the fence operation only applies to operations864performed on objects in `shared_cta` space. Likewise, when `sync_restrict` is865restricted to `shared_cluster`, then memory semantics must be `acquire` and the866effect of the fence operation only applies to operations performed on objects in867`shared_cluster` memory space. The scope for both operations is `cluster`.868For more details, please refer the `PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-membar>`__869 870'``llvm.nvvm.fence.proxy.<proxykind>``'871^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^872 873Syntax:874"""""""875 876.. code-block:: llvm877 878  declare void @llvm.nvvm.fence.proxy.alias()879  declare void @llvm.nvvm.fence.proxy.async()880  declare void @llvm.nvvm.fence.proxy.async.global()881  declare void @llvm.nvvm.fence.proxy.async.shared_cluster()882  declare void @llvm.nvvm.fence.proxy.async.shared_cta()883 884Overview:885"""""""""886 887`nvvm.fence.proxy.{proxykind}` intrinsics represent a fence with bi-directional888proxy ordering that is established between the memory accesses done between the889`generic proxy<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#proxies>__`890and the proxy specified by `proxykind`. A `bi-directional proxy` ordering between891two proxykinds establishes two `uni-directional` proxy orderings: one from the892first proxykind to the second proxykind and the other from the second proxykind893to the first proxykind.894 895`alias` proxykind refers to memory accesses performed using virtually aliased896addresses to the same memory location897 898`async` proxykind specifies that the memory ordering is established between the899`async proxy` and the `generic proxy`. The memory ordering is limited only to900operations performed on objects in the state space specified (`generic`, `global`,901`shared_cluster`, `shared_cta`). If no state space is specified, then the memory902ordering applies on all state spaces. For more details, please refer the903`PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-membar>`__904 905'``llvm.nvvm.fence.proxy.tensormap_generic.*``'906^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^907 908Syntax:909"""""""910 911.. code-block:: llvm912 913  declare void @llvm.nvvm.fence.proxy.tensormap_generic.release.cta()914  declare void @llvm.nvvm.fence.proxy.tensormap_generic.release.cluster()915  declare void @llvm.nvvm.fence.proxy.tensormap_generic.release.gpu()916  declare void @llvm.nvvm.fence.proxy.tensormap_generic.release.sys()917 918  declare void @llvm.nvvm.fence.proxy.tensormap_generic.acquire.cta(ptr %addr, i32 %size)919  declare void @llvm.nvvm.fence.proxy.tensormap_generic.acquire.cluster(ptr %addr, i32 %size)920  declare void @llvm.nvvm.fence.proxy.tensormap_generic.acquire.gpu(ptr %addr, i32 %size)921  declare void @llvm.nvvm.fence.proxy.tensormap_generic.acquire.sys(ptr %addr, i32 %size)922 923Overview:924"""""""""925 926The ``@llvm.nvvm.fence.proxy.tensormap_generic.*`` is a uni-directional fence used to establish ordering between a prior memory access performed via the generic `proxy<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#proxies>_` and a subsequent memory access performed via the tensormap proxy. ``nvvm.fence.proxy.tensormap_generic.release`` can form a release sequence that synchronizes with an acquire sequence that contains the ``nvvm.fence.proxy.tensormap_generic.acquire`` proxy fence. The following table describes the mapping between LLVM Intrinsic and the PTX instruction:927 928  ====================================================== =========================================================929  NVVM Intrinsic                                         PTX Instruction930  ====================================================== =========================================================931  ``@llvm.nvvm.fence.proxy.tensormap_generic.release.*`` ``fence.proxy.tensormap::generic.release.*``932  ``@llvm.nvvm.fence.proxy.tensormap_generic.acquire.*`` ``fence.proxy.tensormap::generic.acquire.* [addr], size``933  ====================================================== =========================================================934 935The address operand ``addr`` and the operand ``size`` together specify the memory range ``[addr, addr+size)`` on which the ordering guarantees on the 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>`_.936 937Address Space Intrinsics938------------------------939 940'``llvm.nvvm.isspacep.*``' Intrinsics941^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^942 943Syntax:944"""""""945 946.. code-block:: llvm947 948    declare i1 @llvm.nvvm.isspacep.const(ptr %p)949    declare i1 @llvm.nvvm.isspacep.global(ptr %p)950    declare i1 @llvm.nvvm.isspacep.local(ptr %p)951    declare i1 @llvm.nvvm.isspacep.shared(ptr %p)952    declare i1 @llvm.nvvm.isspacep.shared.cluster(ptr %p)953 954Overview:955"""""""""956 957The '``llvm.nvvm.isspacep.*``' intrinsics determine whether the provided generic958pointer references memory which falls within a particular address space.959 960Semantics:961""""""""""962 963If the given pointer in the generic address space refers to memory which falls964within the state space of the intrinsic (and therefore could be safely address965space casted to this space), 1 is returned, otherwise 0 is returned.966 967'``llvm.nvvm.mapa.*``' Intrinsics968^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^969 970Syntax:971"""""""972 973.. code-block:: llvm974 975    declare ptr @llvm.nvvm.mapa(ptr %p, i32 %rank)976    declare ptr addrspace(7) @llvm.nvvm.mapa.shared.cluster(ptr addrspace(3) %p, i32 %rank)977 978Overview:979"""""""""980 981The '``llvm.nvvm.mapa.*``' intrinsics map a shared memory pointer ``p`` of another CTA with ``%rank`` to the current CTA.982The ``llvm.nvvm.mapa`` form expects a generic pointer to shared memory and returns a generic pointer to shared cluster memory.983The ``llvm.nvvm.mapa.shared.cluster`` form expects a pointer to shared memory and returns a pointer to shared cluster memory.984They corresponds directly to the ``mapa`` and ``mapa.shared.cluster`` PTX instructions.985 986Semantics:987""""""""""988 989If the given pointer in the generic address space refers to memory which falls990within the state space of the intrinsic (and therefore could be safely address991space casted to this space), 1 is returned, otherwise 0 is returned.992 993Arithmetic Intrinsics994---------------------995 996'``llvm.nvvm.fabs.*``' Intrinsic997^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^998 999Syntax:1000"""""""1001 1002.. code-block:: llvm1003 1004    declare float @llvm.nvvm.fabs.f32(float %a)1005    declare double @llvm.nvvm.fabs.f64(double %a)1006    declare half @llvm.nvvm.fabs.f16(half %a)1007    declare <2 x half> @llvm.nvvm.fabs.v2f16(<2 x half> %a)1008    declare bfloat @llvm.nvvm.fabs.bf16(bfloat %a)1009    declare <2 x bfloat> @llvm.nvvm.fabs.v2bf16(<2 x bfloat> %a)1010 1011Overview:1012"""""""""1013 1014The '``llvm.nvvm.fabs.*``' intrinsics return the absolute value of the operand.1015 1016Semantics:1017""""""""""1018 1019Unlike, '``llvm.fabs.*``', these intrinsics do not perfectly preserve NaN1020values. Instead, a NaN input yields an unspecified NaN output.1021 1022 1023'``llvm.nvvm.fabs.ftz.*``' Intrinsic1024^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^1025 1026Syntax:1027"""""""1028 1029.. code-block:: llvm1030 1031    declare float @llvm.nvvm.fabs.ftz.f32(float %a)1032    declare half @llvm.nvvm.fabs.ftz.f16(half %a)1033    declare <2 x half> @llvm.nvvm.fabs.ftz.v2f16(<2 x half> %a)1034 1035Overview:1036"""""""""1037 1038The '``llvm.nvvm.fabs.ftz.*``' intrinsics return the absolute value of the1039operand, flushing subnormals to sign preserving zero.1040 1041Semantics:1042""""""""""1043 1044Before the absolute value is taken, the input is flushed to sign preserving1045zero if it is a subnormal. In addition, unlike '``llvm.fabs.*``', a NaN input1046yields an unspecified NaN output.1047 1048 1049'``llvm.nvvm.idp2a.[us].[us]``' Intrinsics1050^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^1051 1052Syntax:1053"""""""1054 1055.. code-block:: llvm1056 1057    declare i32 @llvm.nvvm.idp2a.s.s(i32 %a, i32 %b, i1 immarg %is.hi, i32 %c)1058    declare i32 @llvm.nvvm.idp2a.s.u(i32 %a, i32 %b, i1 immarg %is.hi, i32 %c)1059    declare i32 @llvm.nvvm.idp2a.u.s(i32 %a, i32 %b, i1 immarg %is.hi, i32 %c)1060    declare i32 @llvm.nvvm.idp2a.u.u(i32 %a, i32 %b, i1 immarg %is.hi, i32 %c)1061 1062 1063Overview:1064"""""""""1065 1066The '``llvm.nvvm.idp2a.[us].[us]``' intrinsics performs a 2-element vector dot1067product followed by addition. They corresponds directly to the ``dp2a`` PTX 1068instruction.1069 1070Semantics:1071""""""""""1072 1073The 32-bit value in ``%a`` is broken into 2 16-bit values which are extended to107432 bits. For the '``llvm.nvvm.idp2a.u.[us]``' variants zero-extension is used,1075while for the '``llvm.nvvm.idp2a.s.[us]``' sign-extension is used. Two bytes are1076selected from ``%b``, if ``%is.hi`` is true, the most significant bytes are1077selected, otherwise the least significant bytes are selected. These bytes are1078then extended to 32-bits. For the '``llvm.nvvm.idp2a.[us].u``' variants1079zero-extension is used, while for the '``llvm.nvvm.idp2a.[us].s``'1080sign-extension is used. The dot product of these 2-element vectors is added to1081``%c`` to produce the return.1082 1083 1084'``llvm.nvvm.idp4a.[us].[us]``' Intrinsics1085^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^1086 1087Syntax:1088"""""""1089 1090.. code-block:: llvm1091 1092    declare i32 @llvm.nvvm.idp4a.s.s(i32 %a, i32 %b, i32 %c)1093    declare i32 @llvm.nvvm.idp4a.s.u(i32 %a, i32 %b, i32 %c)1094    declare i32 @llvm.nvvm.idp4a.u.s(i32 %a, i32 %b, i32 %c)1095    declare i32 @llvm.nvvm.idp4a.u.u(i32 %a, i32 %b, i32 %c)1096 1097Overview:1098"""""""""1099 1100The '``llvm.nvvm.idp4a.[us].[us]``' intrinsics perform a 4-element vector dot1101product followed by addition. They corresponds directly to the ``dp4a`` PTX1102instruction.1103 1104Semantics:1105""""""""""1106 1107Each of the 4 bytes in both ``%a`` and ``%b`` are extended to 32-bit integers1108forming 2 ``<4 x i32>``. For ``%a``, zero-extension is used in the1109'``llvm.nvvm.idp4a.u.[us]``' variants, while sign-extension is used with1110'``llvm.nvvm.idp4a.s.[us]``' variants. Similarly, for ``%b``, zero-extension is1111used in the '``llvm.nvvm.idp4a.[us].u``' variants, while sign-extension is used1112with '``llvm.nvvm.idp4a.[us].s``' variants. The dot product of these 4-element1113vectors is added to ``%c`` to produce the return.1114 1115Bit Manipulation Intrinsics1116---------------------------1117 1118'``llvm.nvvm.fshl.clamp.*``' Intrinsic1119^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^1120 1121Syntax:1122"""""""1123 1124.. code-block:: llvm1125 1126    declare i32 @llvm.nvvm.fshl.clamp.i32(i32 %hi, i32 %lo, i32 %n)1127 1128Overview:1129"""""""""1130 1131The '``llvm.nvvm.fshl.clamp``' family of intrinsics performs a clamped funnel1132shift left. These intrinsics are very similar to '``llvm.fshl``', except the1133shift amount is clamped at the integer width (instead of modulo it). Currently,1134only ``i32`` is supported.1135 1136Semantics:1137""""""""""1138 1139The '``llvm.nvvm.fshl.clamp``' family of intrinsic functions performs a clamped1140funnel shift left: the first two values are concatenated as { %hi : %lo } (%hi1141is the most significant bits of the wide value), the combined value is shifted1142left, and the most significant bits are extracted to produce a result that is1143the same size as the original arguments. The shift amount is the minimum of the1144value of %n and the bit width of the integer type.1145 1146'``llvm.nvvm.fshr.clamp.*``' Intrinsic1147^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^1148 1149Syntax:1150"""""""1151 1152.. code-block:: llvm1153 1154    declare i32 @llvm.nvvm.fshr.clamp.i32(i32 %hi, i32 %lo, i32 %n)1155 1156Overview:1157"""""""""1158 1159The '``llvm.nvvm.fshr.clamp``' family of intrinsics perform a clamped funnel1160shift right. These intrinsics are very similar to '``llvm.fshr``', except the1161shift amount is clamped at the integer width (instead of modulo it). Currently,1162only ``i32`` is supported.1163 1164Semantics:1165""""""""""1166 1167The '``llvm.nvvm.fshr.clamp``' family of intrinsic functions performs a clamped1168funnel shift right: the first two values are concatenated as { %hi : %lo } (%hi1169is the most significant bits of the wide value), the combined value is shifted1170right, and the least significant bits are extracted to produce a result that is1171the same size as the original arguments. The shift amount is the minimum of the1172value of %n and the bit width of the integer type.1173 1174'``llvm.nvvm.flo.u.*``' Intrinsic1175^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^1176 1177Syntax:1178"""""""1179 1180.. code-block:: llvm1181 1182    declare i32 @llvm.nvvm.flo.u.i32(i32 %a, i1 %shiftamt)1183    declare i32 @llvm.nvvm.flo.u.i64(i64 %a, i1 %shiftamt)1184 1185Overview:1186"""""""""1187 1188The '``llvm.nvvm.flo.u``' family of intrinsics identifies the bit position of the1189leading one, returning either it's offset from the most or least significant bit.1190 1191Semantics:1192""""""""""1193 1194The '``llvm.nvvm.flo.u``' family of intrinsics returns the bit position of the1195most significant 1. If %shiftamt is true, The result is the shift amount needed1196to left-shift the found bit into the most-significant bit position, otherwise1197the result is the shift amount needed to right-shift the found bit into the1198least-significant bit position. 0xffffffff is returned if no 1 bit is found.1199 1200'``llvm.nvvm.flo.s.*``' Intrinsic1201^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^1202 1203Syntax:1204"""""""1205 1206.. code-block:: llvm1207 1208    declare i32 @llvm.nvvm.flo.s.i32(i32 %a, i1 %shiftamt)1209    declare i32 @llvm.nvvm.flo.s.i64(i64 %a, i1 %shiftamt)1210 1211Overview:1212"""""""""1213 1214The '``llvm.nvvm.flo.s``' family of intrinsics identifies the bit position of the1215leading non-sign bit, returning either it's offset from the most or least1216significant bit.1217 1218Semantics:1219""""""""""1220 1221The '``llvm.nvvm.flo.s``' family of intrinsics returns the bit position of the1222most significant 0 for negative inputs and the most significant 1 for 1223non-negative inputs. If %shiftamt is true, The result is the shift amount needed1224to left-shift the found bit into the most-significant bit position, otherwise1225the result is the shift amount needed to right-shift the found bit into the1226least-significant bit position. 0xffffffff is returned if no 1 bit is found.1227 1228'``llvm.nvvm.{zext,sext}.{wrap,clamp}``' Intrinsics1229^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^1230 1231Syntax:1232"""""""1233 1234.. code-block:: llvm1235 1236    declare i32 @llvm.nvvm.zext.wrap(i32 %a, i32 %b)1237    declare i32 @llvm.nvvm.zext.clamp(i32 %a, i32 %b)1238    declare i32 @llvm.nvvm.sext.wrap(i32 %a, i32 %b)1239    declare i32 @llvm.nvvm.sext.clamp(i32 %a, i32 %b)1240 1241Overview:1242"""""""""1243 1244The '``llvm.nvvm.{zext,sext}.{wrap,clamp}``' family of intrinsics extracts the1245low bits of the input value, and zero- or sign-extends them back to the original1246width.1247 1248Semantics:1249""""""""""1250 1251The '``llvm.nvvm.{zext,sext}.{wrap,clamp}``' family of intrinsics returns1252extension of N lowest bits of operand %a. For the '``wrap``' variants, N is the1253value of operand %b modulo 32. For the '``clamp``' variants, N is the value of1254operand %b clamped to the range [0, 32]. The N lowest bits are then1255zero-extended the case of the '``zext``' variants, or sign-extended the case of1256the '``sext``' variants. If N is 0, the result is 0.1257 1258'``llvm.nvvm.bmsk.{wrap,clamp}``' Intrinsic1259^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^1260 1261Syntax:1262"""""""1263 1264.. code-block:: llvm1265 1266    declare i32 @llvm.nvvm.bmsk.wrap(i32 %a, i32 %b)1267    declare i32 @llvm.nvvm.bmsk.clamp(i32 %a, i32 %b)1268 1269Overview:1270"""""""""1271 1272The '``llvm.nvvm.bmsk.{wrap,clamp}``' family of intrinsics creates a bit mask1273given a starting bit position and a bit width.1274 1275Semantics:1276""""""""""1277 1278The '``llvm.nvvm.bmsk.{wrap,clamp}``' family of intrinsics returns a value with1279all bits set to 0 except for %b bits starting at bit position %a. For the1280'``wrap``' variants, the values of %a and %b modulo 32 are used. For the1281'``clamp``' variants, the values of %a and %b are clamped to the range [0, 32],1282which in practice is equivalent to using them as is.1283 1284'``llvm.nvvm.prmt``' Intrinsic1285^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^1286 1287Syntax:1288"""""""1289 1290.. code-block:: llvm1291 1292    declare i32 @llvm.nvvm.prmt(i32 %lo, i32 %hi, i32 %selector)1293 1294Overview:1295"""""""""1296 1297The '``llvm.nvvm.prmt``' constructs a permutation of the bytes of the first two1298operands, selecting based on the third operand.1299 1300Semantics:1301""""""""""1302 1303The bytes in the first two source operands are numbered from 0 to 7:1304{%hi, %lo} = {{b7, b6, b5, b4}, {b3, b2, b1, b0}}. For each byte in the target1305register, a 4-bit selection value is defined.1306 1307The 3 lsbs of the selection value specify which of the 8 source bytes should be1308moved into the target position. The msb defines if the byte value should be1309copied, or if the sign (msb of the byte) should be replicated over all 8 bits1310of the target position (sign extend of the byte value); msb=0 means copy the1311literal value; msb=1 means replicate the sign.1312 1313These 4-bit selection values are pulled from the lower 16-bits of the %selector1314operand, with the least significant selection value corresponding to the least1315significant byte of the destination.1316 1317 1318'``llvm.nvvm.prmt.*``' Intrinsics1319^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^1320 1321Syntax:1322"""""""1323 1324.. code-block:: llvm1325 1326    declare i32 @llvm.nvvm.prmt.f4e(i32 %lo, i32 %hi, i32 %selector)1327    declare i32 @llvm.nvvm.prmt.b4e(i32 %lo, i32 %hi, i32 %selector)1328 1329    declare i32 @llvm.nvvm.prmt.rc8(i32 %lo, i32 %selector)1330    declare i32 @llvm.nvvm.prmt.ecl(i32 %lo, i32 %selector)1331    declare i32 @llvm.nvvm.prmt.ecr(i32 %lo, i32 %selector)1332    declare i32 @llvm.nvvm.prmt.rc16(i32 %lo, i32 %selector)1333 1334Overview:1335"""""""""1336 1337The '``llvm.nvvm.prmt.*``' family of intrinsics constructs a permutation of the1338bytes of the first one or two operands, selecting based on the 2 least1339significant bits of the final operand.1340 1341Semantics:1342""""""""""1343 1344As with the generic '``llvm.nvvm.prmt``' intrinsic, the bytes in the first one1345or two source operands are numbered. The first source operand (%lo) is numbered1346{b3, b2, b1, b0}, in the case of the '``f4e``' and '``b4e``' variants, the1347second source operand (%hi) is numbered {b7, b6, b5, b4}.1348 1349Depending on the 2 least significant bits of the %selector operand, the result1350of the permutation is defined as follows:1351 1352+------------+----------------+--------------+1353|    Mode    | %selector[1:0] |    Output    |1354+------------+----------------+--------------+1355| '``f4e``'  | 0              | {3, 2, 1, 0} |1356|            +----------------+--------------+1357|            | 1              | {4, 3, 2, 1} |1358|            +----------------+--------------+1359|            | 2              | {5, 4, 3, 2} |1360|            +----------------+--------------+1361|            | 3              | {6, 5, 4, 3} |1362+------------+----------------+--------------+1363| '``b4e``'  | 0              | {5, 6, 7, 0} |1364|            +----------------+--------------+1365|            | 1              | {6, 7, 0, 1} |1366|            +----------------+--------------+1367|            | 2              | {7, 0, 1, 2} |1368|            +----------------+--------------+1369|            | 3              | {0, 1, 2, 3} |1370+------------+----------------+--------------+1371| '``rc8``'  | 0              | {0, 0, 0, 0} |1372|            +----------------+--------------+1373|            | 1              | {1, 1, 1, 1} |1374|            +----------------+--------------+1375|            | 2              | {2, 2, 2, 2} |1376|            +----------------+--------------+1377|            | 3              | {3, 3, 3, 3} |1378+------------+----------------+--------------+1379| '``ecl``'  | 0              | {3, 2, 1, 0} |1380|            +----------------+--------------+1381|            | 1              | {3, 2, 1, 1} |1382|            +----------------+--------------+1383|            | 2              | {3, 2, 2, 2} |1384|            +----------------+--------------+1385|            | 3              | {3, 3, 3, 3} |1386+------------+----------------+--------------+1387| '``ecr``'  | 0              | {0, 0, 0, 0} |1388|            +----------------+--------------+1389|            | 1              | {1, 1, 1, 0} |1390|            +----------------+--------------+1391|            | 2              | {2, 2, 1, 0} |1392|            +----------------+--------------+1393|            | 3              | {3, 2, 1, 0} |1394+------------+----------------+--------------+1395| '``rc16``' | 0              | {1, 0, 1, 0} |1396|            +----------------+--------------+1397|            | 1              | {3, 2, 3, 2} |1398|            +----------------+--------------+1399|            | 2              | {1, 0, 1, 0} |1400|            +----------------+--------------+1401|            | 3              | {3, 2, 3, 2} |1402+------------+----------------+--------------+1403 1404TMA family of Intrinsics1405------------------------1406 1407'``llvm.nvvm.cp.async.bulk.global.to.shared.cluster``'1408^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^1409 1410Syntax:1411"""""""1412 1413.. code-block:: llvm1414 1415  declare void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster(ptr addrspace(7) %dst, ptr addrspace(3) %mbar, ptr addrspace(1) %src, i32 %size, i16 %mc, i64 %ch, i1 %flag_mc, i1 %flag_ch)1416 1417Overview:1418"""""""""1419 1420The '``@llvm.nvvm.cp.async.bulk.global.to.shared.cluster``' intrinsic1421corresponds to the ``cp.async.bulk.shared::cluster.global.*`` family1422of PTX instructions. These instructions initiate an asynchronous1423copy of bulk data from global memory to shared::cluster memory.1424The 32-bit operand ``%size`` specifies the amount of memory to be1425copied and it must be a multiple of 16.1426 1427* The last two arguments to these intrinsics are boolean flags1428  indicating support for cache_hint and/or multicast modifiers.1429  These flag arguments must be compile-time constants. The backend1430  looks through these flags and lowers the intrinsics appropriately.1431 1432* The Nth argument (denoted by ``i1 %flag_ch``) when set, indicates1433  a valid cache_hint (``i64 %ch``) and generates the ``.L2::cache_hint``1434  variant of the PTX instruction.1435 1436* The [N-1]th argument (denoted by ``i1 %flag_mc``) when set, indicates1437  the presence of a multicast mask (``i16 %mc``) and generates the PTX1438  instruction with the ``.multicast::cluster`` modifier.1439 1440For more information, refer PTX ISA1441`<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk>`_.1442 1443'``llvm.nvvm.cp.async.bulk.global.to.shared.cta``'1444^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^1445 1446Syntax:1447"""""""1448 1449.. code-block:: llvm1450 1451  declare void @llvm.nvvm.cp.async.bulk.global.to.shared.cta(ptr addrspace(3) %dst, ptr addrspace(3) %mbar, ptr addrspace(1) %src, i32 %size, i64 %ch, i1 %flag_ch)1452 1453Overview:1454"""""""""1455 1456The '``@llvm.nvvm.cp.async.bulk.global.to.shared.cta``' intrinsic1457corresponds to the ``cp.async.bulk.shared::cta.global.*`` family1458of PTX instructions. These instructions initiate an asynchronous1459copy of bulk data from global memory to shared::cta memory.1460The 32-bit operand ``%size`` specifies the amount of memory to be1461copied and it must be a multiple of 16. The last argument1462(denoted by ``i1 %flag_ch``) is a compile-time constant. When set,1463it indicates a valid cache_hint (``i64 %ch``) and generates the1464``.L2::cache_hint`` variant of the PTX instruction.1465 1466For more information, refer PTX ISA1467`<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk>`_.1468 1469'``llvm.nvvm.cp.async.bulk.shared.cta.to.global``'1470^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^1471 1472Syntax:1473"""""""1474 1475.. code-block:: llvm1476 1477  declare void @llvm.nvvm.cp.async.bulk.shared.cta.to.global(ptr addrspace(1) %dst, ptr addrspace(3) %src, i32 %size, i64 %ch, i1 %flag_ch)1478  declare void @llvm.nvvm.cp.async.bulk.shared.cta.to.global.bytemask(..., i32 %size, i64 %ch, i1 %flag_ch, i16 %mask)1479 1480Overview:1481"""""""""1482 1483The '``@llvm.nvvm.cp.async.bulk.shared.cta.to.global``' intrinsic1484corresponds to the ``cp.async.bulk.global.shared::cta.*`` set of PTX1485instructions. These instructions initiate an asynchronous copy from1486shared::cta to global memory. The 32-bit operand ``%size`` specifies1487the amount of memory to be copied (in bytes) and it must be a multiple1488of 16. For the ``.bytemask`` variant, the 16-bit wide mask operand1489specifies whether the i-th byte of each 16-byte wide chunk of source1490data is copied to the destination.1491 1492* The ``i1 %flag_ch`` argument to these intrinsics is a boolean1493  flag indicating support for cache_hint. This flag argument must1494  be a compile-time constant. When set, it indicates a valid1495  cache_hint (``i64 %ch``) and generates the ``.L2::cache_hint``1496  variant of the PTX instruction.1497 1498For more information, refer PTX ISA1499`<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk>`_.1500 1501'``llvm.nvvm.cp.async.bulk.shared.cta.to.cluster``'1502^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^1503 1504Syntax:1505"""""""1506 1507.. code-block:: llvm1508 1509  declare void @llvm.nvvm.cp.async.bulk.shared.cta.to.cluster(ptr addrspace(7) %dst, ptr addrspace(3) %mbar, ptr addrspace(3) %src, i32 %size)1510 1511Overview:1512"""""""""1513 1514The '``@llvm.nvvm.cp.async.bulk.shared.cta.to.cluster``' intrinsic1515corresponds to the ``cp.async.bulk.shared::cluster.shared::cta.*``1516PTX instruction. This instruction initiates an asynchronous copy from1517shared::cta to shared::cluster memory. The destination has to be in1518the shared memory of a different CTA within the cluster. The 32-bit1519operand ``%size`` specifies the amount of memory to be copied and1520it must be a multiple of 16.1521 1522For more information, refer PTX ISA1523`<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk>`_.1524 1525'``llvm.nvvm.cp.async.bulk.prefetch.L2``'1526^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^1527 1528Syntax:1529"""""""1530 1531.. code-block:: llvm1532 1533  declare void @llvm.nvvm.cp.async.bulk.prefetch.L2(ptr addrspace(1) %src, i32 %size, i64 %ch, i1 %flag_ch)1534 1535Overview:1536"""""""""1537 1538The '``@llvm.nvvm.cp.async.bulk.prefetch.L2``' intrinsic1539corresponds to the ``cp.async.bulk.prefetch.L2.*`` family1540of PTX instructions. These instructions initiate an asynchronous1541prefetch of bulk data from global memory to the L2 cache.1542The 32-bit operand ``%size`` specifies the amount of memory to be1543prefetched in terms of bytes and it must be a multiple of 16.1544 1545* The last argument to these intrinsics is boolean flag indicating1546  support for cache_hint. These flag argument must be compile-time1547  constant. When set, it indicates a valid cache_hint (``i64 %ch``)1548  and generates the ``.L2::cache_hint`` variant of the PTX instruction.1549 1550For more information, refer PTX ISA1551`<https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-cp-async-bulk-prefetch>`_.1552 1553'``llvm.nvvm.prefetch.*``'1554^^^^^^^^^^^^^^^^^^^^^^^^^^1555 1556Syntax:1557"""""""1558 1559.. code-block:: llvm1560 1561  declare void  @llvm.nvvm.prefetch.global.L1(ptr addrspace(1) %global_ptr)1562  declare void  @llvm.nvvm.prefetch.global.L2(ptr addrspace(1) %global_ptr)1563  declare void  @llvm.nvvm.prefetch.local.L1(ptr addrspace(5) %local_ptr)1564  declare void  @llvm.nvvm.prefetch.local.L2(ptr addrspace(5) %local_ptr)1565  1566  declare void  @llvm.nvvm.prefetch.L1(ptr %ptr)1567  declare void  @llvm.nvvm.prefetch.L2(ptr %ptr)1568  1569  declare void  @llvm.nvvm.prefetch.tensormap.p0(ptr %ptr)1570  declare void  @llvm.nvvm.prefetch.tensormap.p4(ptr addrspace(4) %const_ptr)1571  declare void  @llvm.nvvm.prefetch.tensormap.p101(ptr addrspace(101) %param_ptr)  1572  1573  declare void  @llvm.nvvm.prefetch.global.L2.evict.normal(ptr addrspace(1) %global_ptr)1574  declare void  @llvm.nvvm.prefetch.global.L2.evict.last(ptr addrspace(1) %global_ptr)1575 1576  declare void  @llvm.nvvm.prefetchu.L1(ptr %ptr)1577 1578Overview:1579"""""""""1580 1581The '``@llvm.nvvm.prefetch.*``' and '``@llvm.nvvm.prefetchu.*``' intrinsic1582correspond to the '``prefetch.*``;' and '``prefetchu.*``' family of PTX instructions. 1583The '``prefetch.*``' instructions bring the cache line containing the1584specified address in global or local memory address space into the 1585specified cache level (L1 or L2). If the '``.tensormap``' qualifier is specified then the 1586prefetch instruction brings the cache line containing the specified address in the 1587'``.const``' or '``.param memory``' state space for subsequent use by the '``cp.async.bulk.tensor``' 1588instruction. The '`prefetchu.*``' instruction brings the cache line 1589containing the specified generic address into the specified uniform cache level.1590If no address space is specified, it is assumed to be generic address. The intrinsic 1591uses and eviction priority which can be accessed by the '``.level::eviction_priority``' modifier.1592 1593* A prefetch to a shared memory location performs no operation.1594* A prefetch into the uniform cache requires a generic address, 1595  and no operation occurs if the address maps to a const, local, or shared memory location.1596 1597For more information, refer to the PTX ISA1598`<https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-prefetch-prefetchu>`_.1599 1600'``llvm.nvvm.applypriority.*``'1601^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^1602 1603Syntax:1604"""""""1605 1606.. code-block:: llvm1607 1608  declare void  @llvm.nvvm.applypriority.global.L2.evict.normal(ptr addrspace(1) %global_ptr, i64 %size)1609  declare void  @llvm.nvvm.applypriority.L2.evict.normal(ptr %ptr, i64 %size)1610 1611Overview:1612"""""""""1613 1614The '``@llvm.nvvm.applypriority.*``'  applies the cache eviction priority specified by the1615.level::eviction_priority qualifier to the address range [a..a+size) in the specified cache 1616level. If no state space is specified then Generic Addressing is used. If the specified address 1617does not fall within the address window of .global state space then the behavior is undefined.1618The operand size is an integer constant that specifies the amount of data, in bytes, in the specified cache1619level on which the priority is to be applied. The only supported value for the size operand is 128.1620 1621For more information, refer to the PTX ISA1622`<https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-applypriority>`_.1623 1624``llvm.nvvm.discard.*``'1625^^^^^^^^^^^^^^^^^^^^^^^^1626 1627Syntax:1628"""""""1629 1630.. code-block:: llvm1631 1632  declare void  @llvm.nvvm.discard.global.L2(ptr addrspace(1) %global_ptr, i64 immarg)1633  declare void  @llvm.nvvm.discard.L2(ptr %ptr, i64 immarg)1634 1635Overview:1636"""""""""1637 1638The *effects* of the ``@llvm.nvvm.discard.L2*`` intrinsics are those of a non-atomic 1639non-volatile ``llvm.memset`` that writes ``undef`` to the destination 1640address range ``[%ptr, %ptr + immarg)``. The ``%ptr`` must be aligned by 128 bytes.1641Subsequent reads from the address range may read ``undef`` until the memory is overwritten 1642with a different value.1643These operations *hint* the implementation that data in the L2 cache can be destructively 1644discarded without writing it back to memory. 1645The operand ``immarg`` is an integer constant that specifies the length in bytes of the 1646address range ``[%ptr, %ptr + immarg)`` to write ``undef`` into. 1647The only supported value for the ``immarg`` operand is ``128``. 1648If generic addressing is used and the specified address does not fall within the 1649address window of global memory (``addrspace(1)``) the behavior is undefined.1650 1651.. code-block:: llvm1652 1653   call void @llvm.nvvm.discard.L2(ptr %p, i64 128)  ;; writes `undef` to [p, p+128)1654   %a = load i64, ptr %p. ;; loads 8 bytes containing undef1655   %b = load i64, ptr %p  ;; loads 8 bytes containing undef1656   ;; comparing %a and %b compares `undef` values!1657   %fa = freeze i64 %a  ;; freezes undef to stable bit-pattern1658   %fb = freeze i64 %b  ;; freezes undef to stable bit-pattern1659   ;; %fa may compare different to %fb!1660   1661For 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>`__ .1662 1663'``llvm.nvvm.cp.async.bulk.tensor.g2s.tile.[1-5]d``'1664^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^1665 1666Syntax:1667"""""""1668 1669.. code-block:: llvm1670 1671  declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.1d(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr %tensor_map, i32 %d0, i16 %mc, i64 %ch, i1 %flag_mc, i1 %flag_ch, i32 %flag_cta_group)1672  declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.2d(..., i32 %d0, i32 %d1, ...)1673  declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.3d(..., i32 %d0, i32 %d1, i32 %d2, ...)1674  declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.4d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, ...)1675  declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.5d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, ...)1676 1677  declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.gather4.2d(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr %tensor_map, i32 %x0, i32 %y0, i32 %y1, i32 %y2, i32 %y3, i16 %mc, i64 %ch, i1 %flag_mc, i1 %flag_ch, i32 %flag_cta_group)1678 1679Overview:1680"""""""""1681 1682The '``@llvm.nvvm.cp.async.bulk.tensor.g2s.tile.[1-5]d``' intrinsics1683correspond to the ``cp.async.bulk.tensor.[1-5]d.*`` set of PTX instructions.1684These instructions initiate an asynchronous copy of tensor data from1685global memory to shared::cluster memory (indicated by the ``g2s`` prefix)1686in ``tile`` mode. In tile mode, the multi-dimensional layout of the1687source tensor is preserved at the destination. The dimension of the1688tensor data ranges from 1d to 5d with the coordinates specified1689by the ``i32 %d0 ... i32 %d4`` arguments. In ``tile.gather4`` mode,1690four rows in a 2D tensor are combined to form a single 2D destination1691tensor. The first coordinate ``i32 %x0`` denotes the column index1692followed by four coordinates indicating the four row-indices.1693So, this mode takes a total of 5 coordinates as input arguments.1694For more information on ``gather4`` mode, refer PTX ISA1695`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-tiled-scatter4-gather4-modes>`_.1696 1697* The last three arguments to these intrinsics are flags1698  indicating support for multicast, cache_hint and cta_group::1/21699  modifiers. These flag arguments must be compile-time constants.1700  The backend looks through these flags and lowers the intrinsics1701  appropriately.1702 1703* The argument denoted by ``i1 %flag_ch`` when set, indicates1704  a valid cache_hint (``i64 %ch``) and generates the ``.L2::cache_hint``1705  variant of the PTX instruction.1706 1707* The argument denoted by ``i1 %flag_mc`` when set, indicates1708  the presence of a multicast mask (``i16 %mc``) and generates1709  the PTX instruction with the ``.multicast::cluster`` modifier.1710 1711* The argument denoted by ``i32 %flag_cta_group`` takes values within1712  the range [0, 3) i.e. {0,1,2}. When the value of ``%flag_cta_group``1713  is not within the range, it may raise an error from the Verifier.1714  The default value is '0' with no cta_group modifier in the1715  instruction. The values of '1' and '2' lower to ``cta_group::1``1716  and ``cta_group::2`` variants of the PTX instruction respectively.1717 1718For more information, refer PTX ISA1719`<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk-tensor>`_.1720 1721'``llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.[3-5]d``'1722^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^1723 1724Syntax:1725"""""""1726 1727.. code-block:: llvm1728 1729  declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.3d(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr %tensor_map, i32 %d0, i32 %d1, i32 %d2, i16 %im2col0, i16 %mc, i64 %ch, i1 %flag_mc, i1 %flag_ch, i32 %flag_cta_group)1730  declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.4d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i16 %im2col0, i16 %im2col1, ...)1731  declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.5d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, i16 %im2col0, i16 %im2col1, i16 %im2col2, ...)1732 1733  declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.w.3d(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr %tensor_map, i32 %d0, i32 %d1, i32 %d2, i16 %wHalo, i16 %wOffset, i16 %mc, i64 %ch, i1 %flag_mc, i1 %flag_ch, i32 %flag_cta_group)1734  declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.w.4d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, ...)1735  declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.w.5d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, ...)1736 1737  declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.w.128.3d(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr %tensor_map, i32 %d0, i32 %d1, i32 %d2, i16 %wHalo, i16 %wOffset, i16 %mc, i64 %ch, i1 %flag_mc, i1 %flag_ch, i32 %flag_cta_group)1738  declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.w.128.4d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, ...)1739  declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.w.128.5d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, ...)1740 1741Overview:1742"""""""""1743 1744The '``@llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.[3-5]d``' intrinsics1745correspond to the ``cp.async.bulk.tensor.[1-5]d.*`` set of PTX instructions.1746These instructions initiate an asynchronous copy of tensor data from1747global memory to shared::cluster memory (indicated by the ``g2s`` prefix)1748in ``im2col`` mode. In im2col mode, some dimensions of the source tensor1749are unrolled into a single dimensional column at the destination. In this1750mode, the tensor has to be at least three-dimensional. Along with the tensor1751coordinates, im2col offsets are also specified (denoted by1752``i16 im2col0...i16 %im2col2``). For the ``im2col`` mode, the number of offsets1753is two less than the number of dimensions of the tensor operation. For the1754``im2col.w`` and ``im2col.w.128`` mode, the number of offsets is always 2,1755denoted by ``i16 %wHalo`` and ``i16 %wOffset`` arguments. For more information1756on ``im2col.w`` and ``im2col.w.128`` modes, refer PTX ISA1757`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-im2col-w-w128-modes>`_.1758 1759The last three arguments to these intrinsics are flags, with the same functionality1760as described in the ``tile`` mode intrinsics above.1761 1762For more information, refer PTX ISA1763`<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk-tensor>`_.1764 1765'``llvm.nvvm.cp.async.bulk.tensor.g2s.cta.tile.[1-5]d``'1766^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^1767 1768Syntax:1769"""""""1770 1771.. code-block:: llvm1772 1773  declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.cta.tile.1d(ptr addrspace(3) %dst, ptr addrspace(3) %bar, ptr %tensor_map, i32 %d0, i64 %ch, i1 %flag_ch)1774  declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.cta.tile.2d(..., i32 %d0, i32 %d1, ...)1775  declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.cta.tile.3d(..., i32 %d0, i32 %d1, i32 %d2, ...)1776  declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.cta.tile.4d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, ...)1777  declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.cta.tile.5d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, ...)1778 1779  declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.cta.tile.gather4.2d(ptr addrspace(3) %dst, ptr addrspace(3) %bar, ptr %tensor_map, i32 %x0, i32 %y0, i32 %y1, i32 %y2, i32 %y3, i64 %ch, i1 %flag_ch)1780 1781Overview:1782"""""""""1783 1784The '``@llvm.nvvm.cp.async.bulk.tensor.g2s.cta.tile.[1-5]d``' intrinsics1785correspond to the ``cp.async.bulk.tensor.[1-5]d.shared::cta.global.*``1786set of PTX instructions. These instructions initiate an asynchronous1787copy of tensor data from global memory to shared::cta memory in1788``tile`` mode. In tile mode, the multi-dimensional layout of the1789source tensor is preserved at the destination. The dimension of the1790tensor data ranges from 1d to 5d with the coordinates specified1791by the ``i32 %d0 ... i32 %d4`` arguments. In ``tile.gather4`` mode,1792four rows in a 2D tensor are combined to form a single 2D destination1793tensor. The first coordinate ``i32 %x0`` denotes the column index1794followed by four coordinates indicating the four row-indices.1795So, this mode takes a total of 5 coordinates as input arguments.1796For more information on ``gather4`` mode, refer PTX ISA1797`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-tiled-scatter4-gather4-modes>`_.1798 1799* The last argument to these intrinsics is a boolean flag1800  indicating support for cache_hint. This flag argument must1801  be a compile-time constant. When set, it indicates a valid1802  cache_hint (``i64 %ch``) and generates the ``.L2::cache_hint``1803  variant of the PTX instruction.1804 1805For more information, refer PTX ISA1806`<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk-tensor>`_.1807 1808'``llvm.nvvm.cp.async.bulk.tensor.g2s.cta.im2col.[3-5]d``'1809^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^1810 1811Syntax:1812"""""""1813 1814.. code-block:: llvm1815 1816  declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.cta.im2col.3d(ptr addrspace(3) %dst, ptr addrspace(3) %bar, ptr %tensor_map, i32 %d0, i32 %d1, i32 %d2, i16 %im2col0, i64 %ch, i1 %flag_ch)1817  declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.cta.im2col.4d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i16 %im2col0, i16 %im2col1, ...)1818  declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.cta.im2col.5d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, i16 %im2col0, i16 %im2col1, i16 %im2col2, ...)1819 1820  declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.cta.im2col.w.3d(ptr addrspace(3) %dst, ptr addrspace(3) %bar, ptr %tensor_map, i32 %d0, i32 %d1, i32 %d2, i16 %wHalo, i16 %wOffset, i64 %ch, i1 %flag_ch)1821  declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.cta.im2col.w.4d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, ...)1822  declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.cta.im2col.w.5d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, ...)1823 1824  declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.cta.im2col.w.128.3d(ptr addrspace(3) %dst, ptr addrspace(3) %bar, ptr %tensor_map, i32 %d0, i32 %d1, i32 %d2, i16 %wHalo, i16 %wOffset, i64 %ch, i1 %flag_ch)1825  declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.cta.im2col.w.128.4d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, ...)1826  declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.cta.im2col.w.128.5d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, ...)1827 1828Overview:1829"""""""""1830 1831The '``@llvm.nvvm.cp.async.bulk.tensor.g2s.cta.im2col.[3-5]d``' intrinsics1832correspond to the ``cp.async.bulk.tensor.[1-5]d.shared::cta.global.*``1833set of PTX instructions. These instructions initiate an asynchronous copy1834of tensor data from global memory to shared::cta memory in ``im2col`` mode.1835In im2col mode, some dimensions of the source tensor are unrolled into a1836single dimensional column at the destination. In this mode, the tensor has1837to be at least three-dimensional. Along with the tensor coordinates, im2col1838offsets are also specified (denoted by ``i16 im2col0...i16 %im2col2``).1839For the ``im2col`` mode, the number of offsets is two less than the number1840of dimensions of the tensor operation. For the ``im2col.w`` and ``im2col.w.128``1841mode, the number of offsets is always 2, denoted by ``i16 %wHalo`` and1842``i16 %wOffset`` arguments. For more information on ``im2col.w`` and1843``im2col.w.128`` modes, refer PTX ISA1844`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-im2col-w-w128-modes>`_.1845 1846* The last argument to these intrinsics is a boolean flag1847  indicating support for cache_hint. This flag argument must1848  be a compile-time constant. When set, it indicates a valid1849  cache_hint (``i64 %ch``) and generates the ``.L2::cache_hint``1850  variant of the PTX instruction.1851 1852For more information, refer PTX ISA1853`<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk-tensor>`_.1854 1855'``llvm.nvvm.cp.async.bulk.tensor.s2g.tile.[1-5]d``'1856^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^1857 1858Syntax:1859"""""""1860 1861.. code-block:: llvm1862 1863  declare void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.1d(ptr addrspace(3) %src, ptr %tensor_map, i32 %d0, i64 %ch, i1 %flag_ch)1864  declare void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.2d(..., i32 %d0, i32 %d1, ...)1865  declare void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.3d(..., i32 %d0, i32 %d1, i32 %d2, ...)1866  declare void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.4d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, ...)1867  declare void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.5d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, ...)1868 1869  declare void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.scatter4.2d(ptr addrspace(3) %src, ptr %tensor_map, i32 %x0, i32 %y0, i32 %y1, i32 %y2, i32 %y3, i64 %ch, i1 %flag_ch)1870 1871Overview:1872"""""""""1873 1874The '``@llvm.nvvm.cp.async.bulk.tensor.s2g.tile.[1-5]d``' intrinsics1875correspond to the ``cp.async.bulk.tensor.[1-5]d.*`` set of PTX instructions.1876These instructions initiate an asynchronous copy of tensor data from1877shared::cta to global memory (indicated by the ``s2g`` prefix)1878in ``tile`` mode. The dimension of the tensor data ranges from 1d to 5d1879with the coordinates specified by the ``i32 %d0 ... i32 %d4`` arguments.1880In ``tile.scatter4`` mode, a single 2D source tensor is divided into1881four rows in the 2D destination tensor. The first coordinate ``i32 %x0``1882denotes the column index followed by four coordinates indicating the1883four row-indices. So, this mode takes a total of 5 coordinates as input arguments.1884For more information on ``scatter4`` mode, refer PTX ISA1885`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-tiled-scatter4-gather4-modes>`_.1886 1887* The last argument to these intrinsics is a boolean flag1888  indicating support for cache_hint. This flag argument must1889  be a compile-time constant. When set, it indicates a valid1890  cache_hint (``i64 %ch``) and generates the ``.L2::cache_hint``1891  variant of the PTX instruction.1892 1893For more information, refer PTX ISA1894`<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk-tensor>`_.1895 1896'``llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.[3-5]d``'1897^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^1898 1899Syntax:1900"""""""1901 1902.. code-block:: llvm1903 1904  declare void @llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.3d(ptr addrspace(3) %src, ptr %tensor_map, i32 %d0, i32 %d1, i32 %d2, i64 %ch, i1 %flag_ch)1905  declare void @llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.4d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, ...)1906  declare void @llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.5d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, ...)1907 1908Overview:1909"""""""""1910 1911The '``@llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.[1-5]d``' intrinsics1912correspond to the ``cp.async.bulk.tensor.[1-5]d.*`` set of PTX instructions.1913These instructions initiate an asynchronous copy of tensor data from1914shared::cta to global memory (indicated by the ``s2g`` prefix)1915in ``im2col`` mode. In this mode, the tensor has to be at least1916three-dimensional. Unlike the ``g2s`` variants, there are no1917im2col_offsets for these intrinsics. The last argument to these1918intrinsics is a boolean flag, with the same functionality as1919described in the ``s2g.tile`` mode intrinsics above.1920 1921For more information, refer PTX ISA1922`<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk-tensor>`_.1923 1924'``llvm.nvvm.cp.async.bulk.tensor.prefetch.tile.[1-5]d``'1925^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^1926 1927Syntax:1928"""""""1929 1930.. code-block:: llvm1931 1932  declare void @llvm.nvvm.cp.async.bulk.tensor.prefetch.tile.1d(ptr %tensor_map, i32 %d0, i64 %ch, i1 %flag_ch)1933  declare void @llvm.nvvm.cp.async.bulk.tensor.prefetch.tile.2d(..., i32 %d0, i32 %d1, ...)1934  declare void @llvm.nvvm.cp.async.bulk.tensor.prefetch.tile.3d(..., i32 %d0, i32 %d1, i32 %d2, ...)1935  declare void @llvm.nvvm.cp.async.bulk.tensor.prefetch.tile.4d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, ...)1936  declare void @llvm.nvvm.cp.async.bulk.tensor.prefetch.tile.5d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, ...)1937 1938  declare void @llvm.nvvm.cp.async.bulk.tensor.prefetch.tile.gather4.2d(ptr %tensor_map, i32 %x0, i32 %y0, i32 %y1, i32 %y2, i32 %y3, i64 %ch, i1 %flag_ch)1939 1940Overview:1941"""""""""1942 1943The '``@llvm.nvvm.cp.async.bulk.tensor.prefetch.tile.[1-5]d``' intrinsics1944correspond to the ``cp.async.bulk.prefetch.tensor.[1-5]d.L2.global*`` set1945of PTX instructions. These instructions initiate an asynchronous prefetch1946of tensor data from global memory to the L2 cache. In tile mode, the1947multi-dimensional layout of the source tensor is preserved at the destination.1948The dimension of the tensor data ranges from 1d to 5d with the coordinates1949specified by the ``i32 %d0 ... i32 %d4`` arguments.1950 1951In ``tile.gather4`` mode, four rows in the 2-dimnesional source tensor are1952fetched to the L2 cache. The first coordinate ``i32 %x0`` denotes the column index1953followed by four coordinates indicating the four row-indices. So, this mode takes1954a total of 5 coordinates as input arguments.1955For more information on ``gather4`` mode, refer PTX ISA1956`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-tiled-scatter4-gather4-modes>`_.1957 1958* The last argument to these intrinsics is a boolean flag1959  indicating support for cache_hint. This flag argument must1960  be a compile-time constant. When set, it indicates a valid1961  cache_hint (``i64 %ch``) and generates the ``.L2::cache_hint``1962  variant of the PTX instruction.1963 1964For more information, refer PTX ISA1965`<https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-cp-async-bulk-prefetch-tensor>`_.1966 1967'``llvm.nvvm.cp.async.bulk.tensor.prefetch.im2col.[3-5]d``'1968^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^1969 1970Syntax:1971"""""""1972 1973.. code-block:: llvm1974 1975  declare void @llvm.nvvm.cp.async.bulk.tensor.prefetch.im2col.3d(ptr %tensor_map, i32 %d0, i32 %d1, i32 %d2, i16 %im2col0, i64 %ch, i1 %flag_ch)1976  declare void @llvm.nvvm.cp.async.bulk.tensor.prefetch.im2col.4d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i16 %im2col0, i16 %im2col1, ...)1977  declare void @llvm.nvvm.cp.async.bulk.tensor.prefetch.im2col.5d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, i16 %im2col0, i16 %im2col1, i16 %im2col2, ...)1978 1979  declare void @llvm.nvvm.cp.async.bulk.tensor.prefetch.im2col.w.3d(ptr %tensor_map, i32 %d0, i32 %d1, i32 %d2, i16 %wHalo, i16 %wOffset, i64 %ch, i1 %flag_ch)1980  declare void @llvm.nvvm.cp.async.bulk.tensor.prefetch.im2col.w.4d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, ...)1981  declare void @llvm.nvvm.cp.async.bulk.tensor.prefetch.im2col.w.5d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, ...)1982 1983  declare void @llvm.nvvm.cp.async.bulk.tensor.prefetch.im2col.w.128.3d(ptr %tensor_map, i32 %d0, i32 %d1, i32 %d2, i16 %wHalo, i16 %wOffset, i64 %ch, i1 %flag_ch)1984  declare void @llvm.nvvm.cp.async.bulk.tensor.prefetch.im2col.w.128.4d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, ...)1985  declare void @llvm.nvvm.cp.async.bulk.tensor.prefetch.im2col.w.128.5d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, ...)1986 1987Overview:1988"""""""""1989 1990The '``@llvm.nvvm.cp.async.bulk.tensor.prefetch.im2col.[3-5]d``' intrinsics1991correspond to the ``cp.async.bulk.prefetch.tensor.[1-5]d.L2.global*`` set1992of PTX instructions. These instructions initiate an asynchronous prefetch1993of tensor data from global memory to the L2 cache. In im2col mode, some1994dimensions of the source tensor are unrolled into a single dimensional1995column at the destination. In this mode, the tensor has to be at least1996three-dimensional. Along with the tensor coordinates, im2col offsets are1997also specified (denoted by ``i16 im2col0...i16 %im2col2``). For ``im2col``1998mode, the number of offsets is two less than the number of dimensions of1999the tensor operation. For the ``im2col.w`` and ``im2col.w.128`` modes,2000the number of offsets is always 2, denoted by ``i16 %wHalo`` and2001``i16 %wOffset`` arguments. For more information on ``im2col.w`` and2002``im2col.w.128`` modes, refer PTX ISA2003`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-im2col-w-w128-modes>`_.2004 2005 2006The last argument to these intrinsics is a boolean flag, with2007the same functionality as described in the ``tile`` mode intrinsics above.2008 2009For more information, refer PTX ISA2010`<https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-cp-async-bulk-prefetch-tensor>`_.2011 2012'``llvm.nvvm.cp.async.bulk.tensor.reduce.[red_op].tile.[1-5]d``'2013^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^2014 2015Syntax:2016"""""""2017 2018.. code-block:: llvm2019 2020  declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.add.tile.1d(ptr addrspace(3) %src, ptr %tensor_map, i32 %d0, i64 %ch, i1 %flag_ch)2021  declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.min.tile.1d(ptr addrspace(3) %src, ptr %tensor_map, i32 %d0, i64 %ch, i1 %flag_ch)2022  declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.max.tile.1d(ptr addrspace(3) %src, ptr %tensor_map, i32 %d0, i64 %ch, i1 %flag_ch)2023  declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.inc.tile.1d(ptr addrspace(3) %src, ptr %tensor_map, i32 %d0, i64 %ch, i1 %flag_ch)2024  declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.dec.tile.1d(ptr addrspace(3) %src, ptr %tensor_map, i32 %d0, i64 %ch, i1 %flag_ch)2025  declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.and.tile.1d(ptr addrspace(3) %src, ptr %tensor_map, i32 %d0, i64 %ch, i1 %flag_ch)2026  declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.or.tile.1d(ptr addrspace(3) %src, ptr %tensor_map, i32 %d0, i64 %ch, i1 %flag_ch)2027  declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.xor.tile.1d(ptr addrspace(3) %src, ptr %tensor_map, i32 %d0, i64 %ch, i1 %flag_ch)2028 2029  declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.<red_op>.tile.2d(..., i32 %d0, i32 %d1, ...)2030  declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.<red_op>.tile.3d(..., i32 %d0, i32 %d1, i32 %d2, ...)2031  declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.<red_op>.tile.4d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, ...)2032  declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.<red_op>.tile.5d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, ...)2033 2034Overview:2035"""""""""2036 2037The '``@llvm.nvvm.cp.async.bulk.tensor.reduce.<red_op>.tile.[1-5]d``' intrinsics2038correspond to the ``cp.reduce.async.bulk.tensor.[1-5]d.*`` set of PTX instructions.2039These instructions initiate an asynchronous reduction operation of tensor data2040in global memory with the tensor data in shared{::cta} memory, using ``tile`` mode.2041The dimension of the tensor data ranges from 1d to 5d with the coordinates2042specified by the ``i32 %d0 ... i32 %d4`` arguments. The supported reduction2043operations are {add, min, max, inc, dec, and, or, xor} as described in the2044``tile.1d`` intrinsics.2045 2046* The last argument to these intrinsics is a boolean flag2047  indicating support for cache_hint. This flag argument must2048  be a compile-time constant. When set, it indicates a valid2049  cache_hint (``i64 %ch``) and generates the ``.L2::cache_hint``2050  variant of the PTX instruction.2051 2052For more information, refer PTX ISA2053`<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-reduce-async-bulk-tensor>`_.2054 2055'``llvm.nvvm.cp.async.bulk.tensor.reduce.[red_op].im2col.[3-5]d``'2056^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^2057 2058Syntax:2059"""""""2060 2061.. code-block:: llvm2062 2063  declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.<red_op>.im2col.3d(ptr addrspace(3) %src, ptr %tensor_map, i32 %d0, i32 %d1, i32 %d2, i64 %ch, i1 %flag_ch)2064  declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.<red_op>.im2col.4d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, ...)2065  declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.<red_op>.im2col.5d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, ...)2066 2067Overview:2068"""""""""2069 2070The '``@llvm.nvvm.cp.async.bulk.tensor.reduce.<red_op>.im2col.[3-5]d``' intrinsics2071correspond to the ``cp.reduce.async.bulk.tensor.[3-5]d.*`` set of PTX instructions.2072These instructions initiate an asynchronous reduction operation of tensor data2073in global memory with the tensor data in shared{::cta} memory, using ``im2col`` mode.2074In this mode, the tensor has to be at least three-dimensional. The supported reduction2075operations supported are the same as the ones in the tile mode. The last argument to2076these intrinsics is a boolean flag, with the same functionality as described in the2077``tile`` mode intrinsics above.2078 2079For more information, refer PTX ISA2080`<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-reduce-async-bulk-tensor>`_.2081 2082Warp Group Intrinsics2083---------------------2084 2085'``llvm.nvvm.wgmma.fence.sync.aligned``'2086^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^2087 2088Syntax:2089"""""""2090 2091.. code-block:: llvm2092 2093  declare void @llvm.nvvm.wgmma.fence.sync.aligned()2094 2095Overview:2096"""""""""2097 2098The '``@llvm.nvvm.wgmma.fence.sync.aligned``' intrinsic generates the2099``wgmma.fence.sync.aligned`` PTX instruction, which establishes an ordering2100between prior accesses to any warpgroup registers and subsequent accesses to2101the same registers by a ``wgmma.mma_async`` instruction.2102 2103The ``wgmma.fence`` instruction must be issued by all warps of the warpgroup in2104the following locations:2105 2106* Before the first ``wgmma.mma_async`` operation in a warpgroup.2107* Between a register access by a thread in the warpgroup and any2108  ``wgmma.mma_async`` instruction that accesses the same registers, except when2109  these are accumulator register accesses across multiple ``wgmma.mma_async``2110  instructions of the same shape in which case an ordering guarantee is2111  provided by default.2112 2113For more information, refer PTX ISA2114`<https://docs.nvidia.com/cuda/parallel-thread-execution/#asynchronous-warpgroup-level-matrix-instructions-wgmma-fence>`_.2115 2116'``llvm.nvvm.wgmma.commit_group.sync.aligned``'2117^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^2118 2119Syntax:2120"""""""2121 2122.. code-block:: llvm2123 2124  declare void @llvm.nvvm.wgmma.commit_group.sync.aligned()2125 2126Overview:2127"""""""""2128 2129The '``@llvm.nvvm.wgmma.commit_group.sync.aligned``' intrinsic generates the2130``wgmma.commit_group.sync.aligned`` PTX instruction, which creates a new2131wgmma-group per warpgroup and batches all prior ``wgmma.mma_async``2132instructions initiated by the executing warp but not committed to any2133wgmma-group into the new wgmma-group. If there are no uncommitted ``wgmma2134mma_async`` instructions then, ``wgmma.commit_group`` results in an empty2135wgmma-group.2136 2137An executing thread can wait for the completion of all ``wgmma.mma_async``2138operations in a wgmma-group by using ``wgmma.wait_group``.2139 2140For more information, refer PTX ISA2141`<https://docs.nvidia.com/cuda/parallel-thread-execution/#asynchronous-warpgroup-level-matrix-instructions-wgmma-commit-group>`_.2142 2143'``llvm.nvvm.wgmma.wait_group.sync.aligned``'2144^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^2145 2146Syntax:2147"""""""2148 2149.. code-block:: llvm2150 2151  declare void @llvm.nvvm.wgmma.wait_group.sync.aligned(i64 immarg N)2152 2153Overview:2154"""""""""2155 2156The '``@llvm.nvvm.wgmma.wait_group.sync.aligned``' intrinsic generates the2157``wgmma.commit_group.sync.aligned N`` PTX instruction, which will cause the2158executing thread to wait until only ``N`` or fewer of the most recent2159wgmma-groups are pending and all the prior wgmma-groups committed by the2160executing threads are complete. For example, when ``N`` is 0, the executing2161thread waits on all the prior wgmma-groups to complete. Operand ``N`` is an2162integer constant.2163 2164Accessing the accumulator register or the input register containing the2165fragments of matrix A of a ``wgmma.mma_async`` instruction without first2166performing a ``wgmma.wait_group`` instruction that waits on a wgmma-group2167including that ``wgmma.mma_async`` instruction is undefined behavior.2168 2169For more information, refer PTX ISA2170`<https://docs.nvidia.com/cuda/parallel-thread-execution/#asynchronous-warpgroup-level-matrix-instructions-wgmma-wait-group>`_.2171 2172'``llvm.nvvm.griddepcontrol.*``'2173^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^2174 2175Syntax:2176"""""""2177 2178.. code-block:: llvm2179 2180  declare void @llvm.nvvm.griddepcontrol.launch_dependents()2181  declare void @llvm.nvvm.griddepcontrol.wait()2182 2183Overview:2184"""""""""2185 2186The ``griddepcontrol`` intrinsics allows the dependent grids and prerequisite grids as defined by the runtime, to control execution in the following way:2187 2188``griddepcontrol.launch_dependents`` intrinsic signals that the dependents can be scheduled, before the current grid completes. The intrinsic can be invoked by multiple threads in the current CTA and repeated invocations of the intrinsic will have no additional side effects past that of the first invocation.2189 2190``griddepcontrol.wait`` intrinsic causes the executing thread to wait until all prerequisite grids in flight have completed and all the memory operations from the prerequisite grids are performed and made visible to the current grid.2191 2192For more information, refer 2193`PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#parallel-synchronization-and-communication-instructions-griddepcontrol>`__.2194 2195TCGEN05 family of Intrinsics2196----------------------------2197 2198The llvm.nvvm.tcgen05.* intrinsics model the TCGEN05 family of instructions2199exposed by PTX. These intrinsics use 'Tensor Memory' (henceforth ``tmem``).2200NVPTX represents this memory using ``addrspace(6)`` and is always 32-bits.2201 2202For more information, refer to the PTX ISA2203`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-memory>`_.2204 2205The tensor-memory pointers may only be used with the tcgen05 intrinsics.2206There are specialized load/store instructions provided (tcgen05.ld/st) to2207work with tensor-memory.2208 2209See the PTX ISA for more information on tensor-memory load/store instructions2210`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-memory-and-register-load-store-instructions>`_.2211 2212'``llvm.nvvm.tcgen05.alloc``'2213^^^^^^^^^^^^^^^^^^^^^^^^^^^^^2214 2215Syntax:2216"""""""2217 2218.. code-block:: llvm2219 2220  declare void @llvm.nvvm.tcgen05.alloc.cg1(ptr %dst, i32 %ncols)2221  declare void @llvm.nvvm.tcgen05.alloc.cg2(ptr %dst, i32 %ncols)2222  declare void @llvm.nvvm.tcgen05.alloc.shared.cg1(ptr addrspace(3) %dst, i32 %ncols)2223  declare void @llvm.nvvm.tcgen05.alloc.shared.cg2(ptr addrspace(3) %dst, i32 %ncols)2224 2225Overview:2226"""""""""2227 2228The '``@llvm.nvvm.tcgen05.alloc.*``' intrinsics correspond to the2229``tcgen05.alloc.cta_group*.sync.aligned.b32`` family of PTX instructions.2230The ``tcgen05.alloc`` is a potentially blocking instruction which dynamically2231allocates the specified number of columns in the Tensor Memory and writes2232the address of the allocated Tensor Memory into shared memory at the2233location specified by ``%dst``. The 32-bit operand ``%ncols`` specifies2234the number of columns to be allocated and it must be a power-of-two.2235The ``.shared`` variant explicitly uses shared memory address space for2236the ``%dst`` operand. The ``.cg1`` and ``.cg2`` variants generate2237``cta_group::1`` and ``cta_group::2`` variants of the instruction respectively.2238 2239For more information, refer to the PTX ISA2240`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-memory-allocation-and-management-instructions>`_.2241 2242'``llvm.nvvm.tcgen05.dealloc``'2243^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^2244 2245Syntax:2246"""""""2247 2248.. code-block:: llvm2249 2250  declare void @llvm.nvvm.tcgen05.dealloc.cg1(ptr addrspace(6) %tmem_addr, i32 %ncols)2251  declare void @llvm.nvvm.tcgen05.dealloc.cg2(ptr addrspace(6) %tmem_addr, i32 %ncols)2252 2253Overview:2254"""""""""2255 2256The '``@llvm.nvvm.tcgen05.dealloc.*``' intrinsics correspond to the2257``tcgen05.dealloc.*`` set of PTX instructions. The ``tcgen05.dealloc``2258instructions deallocates the Tensor Memory specified by the Tensor Memory2259address ``%tmem_addr``. The operand ``%tmem_addr`` must point to a previous2260Tensor Memory allocation. The 32-bit operand ``%ncols`` specifies the number2261of columns to be de-allocated. The ``.cg1`` and ``.cg2`` variants generate2262``cta_group::1`` and ``cta_group::2`` variants of the instruction respectively.2263 2264For more information, refer to the PTX ISA2265`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-memory-allocation-and-management-instructions>`_.2266 2267'``llvm.nvvm.tcgen05.relinq.alloc.permit``'2268^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^2269 2270Syntax:2271"""""""2272 2273.. code-block:: llvm2274 2275  declare void @llvm.nvvm.tcgen05.relinq.alloc.permit.cg1()2276  declare void @llvm.nvvm.tcgen05.relinq.alloc.permit.cg2()2277 2278Overview:2279"""""""""2280 2281The '``@llvm.nvvm.tcgen05.relinq.alloc.permit.*``' intrinsics correspond2282to the ``tcgen05.relinquish_alloc_permit.*`` set of PTX instructions.2283This instruction specifies that the CTA of the executing thread is2284relinquishing the right to allocate Tensor Memory. So, it is illegal2285for a CTA to perform ``tcgen05.alloc`` after any of its constituent2286threads execute ``tcgen05.relinquish_alloc_permit``. The ``.cg1``2287and ``.cg2`` variants generate ``cta_group::1`` and ``cta_group::2``2288flavors of the instruction respectively.2289 2290For more information, refer to the PTX ISA2291`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-memory-allocation-and-management-instructions>`_.2292 2293'``llvm.nvvm.tcgen05.commit``'2294^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^2295 2296Syntax:2297"""""""2298 2299.. code-block:: llvm2300 2301  declare void @llvm.nvvm.tcgen05.commit.{cg1,cg2}(ptr %mbar)2302  declare void @llvm.nvvm.tcgen05.commit.shared.{cg1,cg2}(ptr addrspace(3) %mbar)2303  declare void @llvm.nvvm.tcgen05.commit.mc.{cg1,cg2}(ptr %mbar, i16 %mc)2304  declare void @llvm.nvvm.tcgen05.commit.mc.shared.{cg1,cg2}(ptr addrspace(3) %mbar, i16 %mc)2305 2306Overview:2307"""""""""2308 2309The '``@llvm.nvvm.tcgen05.commit.*``' intrinsics correspond to the2310``tcgen05.commit.{cg1/cg2}.mbarrier::arrive::one.*`` set of PTX instructions.2311The ``tcgen05.commit`` is an asynchronous instruction which makes the mbarrier2312object (``%mbar``) track the completion of all prior asynchronous tcgen05 operations.2313The ``.mc`` variants allow signaling on the mbarrier objects of multiple CTAs2314(specified by ``%mc``) in the cluster. The ``.cg1`` and ``.cg2`` variants generate2315``cta_group::1`` and ``cta_group::2`` flavors of the instruction respectively.2316 2317For more information, refer to the PTX ISA2318`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen-async-sync-operations-commit>`_.2319 2320'``llvm.nvvm.tcgen05.wait``'2321^^^^^^^^^^^^^^^^^^^^^^^^^^^^2322 2323Syntax:2324"""""""2325 2326.. code-block:: llvm2327 2328  declare void @llvm.nvvm.tcgen05.wait.ld()2329  declare void @llvm.nvvm.tcgen05.wait.st()2330 2331Overview:2332"""""""""2333 2334The '``@llvm.nvvm.tcgen05.wait.ld/st``' intrinsics correspond to2335the ``tcgen05.wait::{ld/st}.sync.aligned`` pair of PTX instructions.2336The ``tcgen05.wait::ld`` causes the executing thread to block until2337all prior ``tcgen05.ld`` operations issued by the executing thread2338have completed. The ``tcgen05.wait::st`` causes the executing thread2339to block until all prior ``tcgen05.st`` operations issued by the2340executing thread have completed.2341 2342For more information, refer to the PTX ISA2343`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-instructions-tcgen05-wait>`_.2344 2345'``llvm.nvvm.tcgen05.fence``'2346^^^^^^^^^^^^^^^^^^^^^^^^^^^^^2347 2348Syntax:2349"""""""2350 2351.. code-block:: llvm2352 2353  declare void @llvm.nvvm.tcgen05.fence.before.thread.sync()2354  declare void @llvm.nvvm.tcgen05.fence.after.thread.sync()2355 2356Overview:2357"""""""""2358 2359The '``@llvm.nvvm.tcgen05.fence.*``' intrinsics correspond to2360the ``tcgen05.fence::{before/after}_thread_sync`` pair of PTX instructions.2361These instructions act as code motion fences for asynchronous tcgen052362operations.2363 2364For more information, refer to the PTX ISA2365`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensorcore-5th-generation-instructions-tcgen05-fence>`_.2366 2367'``llvm.nvvm.tcgen05.shift``'2368^^^^^^^^^^^^^^^^^^^^^^^^^^^^^2369 2370Syntax:2371"""""""2372 2373.. code-block:: llvm2374 2375  declare void @llvm.nvvm.tcgen05.shift.down.cg1(ptr addrspace(6) %tmem_addr)2376  declare void @llvm.nvvm.tcgen05.shift.down.cg2(ptr addrspace(6) %tmem_addr)2377 2378Overview:2379"""""""""2380 2381The '``@llvm.nvvm.tcgen05.shift.{cg1/cg2}``' intrinsics correspond to2382the ``tcgen05.shift.{cg1/cg2}`` PTX instructions. The ``tcgen05.shift``2383is an asynchronous instruction which initiates the shifting of 32-byte2384elements downwards across all the rows, except the last, by one row.2385The address operand ``%tmem_addr`` specifies the base address of the2386matrix in the Tensor Memory whose rows must be down shifted.2387 2388For more information, refer to the PTX ISA2389`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-instructions-tcgen05-shift>`_.2390 2391'``llvm.nvvm.tcgen05.cp``'2392^^^^^^^^^^^^^^^^^^^^^^^^^^2393 2394Syntax:2395"""""""2396 2397.. code-block:: llvm2398 2399  declare void @llvm.nvvm.tcgen05.cp.4x256b.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i64 %sdesc)2400  declare void @llvm.nvvm.tcgen05.cp.128x256b.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i64 %sdesc)2401  declare void @llvm.nvvm.tcgen05.cp.128x128b.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i64 %sdesc)2402  declare void @llvm.nvvm.tcgen05.cp.32x128b_warpx4.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i64 %sdesc)2403  declare void @llvm.nvvm.tcgen05.cp.64x128b_warpx2_02_13.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i64 %sdesc)2404  declare void @llvm.nvvm.tcgen05.cp.64x128b_warpx2_01_23.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i64 %sdesc)2405 2406  declare void @llvm.nvvm.tcgen05.cp.4x256b.b6x16_p32.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i64 %sdesc)2407  declare void @llvm.nvvm.tcgen05.cp.128x256b.b6x16_p32.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i64 %sdesc)2408  declare void @llvm.nvvm.tcgen05.cp.128x128b.b6x16_p32.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i64 %sdesc)2409  declare void @llvm.nvvm.tcgen05.cp.32x128b_warpx4.b6x16_p32.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i64 %sdesc)2410  declare void @llvm.nvvm.tcgen05.cp.64x128b_warpx2_02_13.b6x16_p32.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i64 %sdesc)2411  declare void @llvm.nvvm.tcgen05.cp.64x128b_warpx2_01_23.b6x16_p32.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i64 %sdesc)2412 2413  declare void @llvm.nvvm.tcgen05.cp.4x256b.b4x16_p64.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i64 %sdesc)2414  declare void @llvm.nvvm.tcgen05.cp.128x256b.b4x16_p64.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i64 %sdesc)2415  declare void @llvm.nvvm.tcgen05.cp.128x128b.b4x16_p64.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i64 %sdesc)2416  declare void @llvm.nvvm.tcgen05.cp.32x128b_warpx4.b4x16_p64.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i64 %sdesc)2417  declare void @llvm.nvvm.tcgen05.cp.64x128b_warpx2_02_13.b4x16_p64.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i64 %sdesc)2418  declare void @llvm.nvvm.tcgen05.cp.64x128b_warpx2_01_23.b4x16_p64.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i64 %sdesc)2419 2420Overview:2421"""""""""2422 2423The '``@llvm.nvvm.tcgen05.cp.{shape}.{src_fmt}.{cg1/cg2}``' intrinsics2424correspond to the ``tcgen05.cp.*`` family of PTX instructions.2425The ``tcgen05.cp`` instruction initiates an asynchronous copy operation from2426shared memory to the location specified by ``%tmem_addr`` in Tensor Memory.2427The 64-bit register operand ``%sdesc`` is the matrix descriptor representing2428the source matrix in shared memory that needs to be copied.2429 2430The valid shapes for the copy operation are:2431{128x256b, 4x256b, 128x128b, 64x128b_warpx2_02_13, 64x128b_warpx2_01_23, 32x128b_warpx4}.2432 2433Shapes ``64x128b`` and ``32x128b`` require dedicated multicast qualifiers,2434which are appended to the corresponding intrinsic names.2435 2436Optionally, the data can be decompressed from the source format in the shared memory2437to the destination format in Tensor Memory during the copy operation. Currently,2438only ``.b8x16`` is supported as destination format. The valid source formats are2439``.b6x16_p32`` and ``.b4x16_p64``.2440 2441When the source format is ``.b6x16_p32``, a contiguous set of 16 elements of 6-bits2442each followed by four bytes of padding (``_p32``) in shared memory is decompressed2443into 16 elements of 8-bits (``.b8x16``) each in the Tensor Memory.2444 2445When the source format is ``.b4x16_p64``, a contiguous set of 16 elements of 4-bits2446each followed by eight bytes of padding (``_p64``) in shared memory is decompressed2447into 16 elements of 8-bits (``.b8x16``) each in the Tensor Memory.2448 2449For more information on the decompression schemes, refer to the PTX ISA2450`<https://docs.nvidia.com/cuda/parallel-thread-execution/#optional-decompression>`_.2451 2452For more information on the tcgen05.cp instruction, refer to the PTX ISA2453`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-instructions-tcgen05-cp>`_.2454 2455'``llvm.nvvm.tcgen05.ld.*``'2456^^^^^^^^^^^^^^^^^^^^^^^^^^^^2457 2458Syntax:2459"""""""2460 2461.. code-block:: llvm2462 2463  declare <n x i32> @llvm.nvvm.tcgen05.ld.<shape>.<num>(ptr addrspace(6) %tmem_addr, i1 %pack)2464 2465  declare <n x i32> @llvm.nvvm.tcgen05.ld.16x32bx2.<num>(ptr addrspace(6) %tmem_addr, i64 %offset, i1 %pack)2466 2467Overview:2468"""""""""2469 2470This group of intrinsics asynchronously load data from the Tensor Memory at the location specified2471by the 32-bit address operand `tmem_addr` into the destination registers, collectively across all threads2472of the warps.2473 2474All the threads in the warp must specify the same value of `tmem_addr`, which must be the base address2475of the collective load operation. Otherwise, the behavior is undefined.2476 2477The `shape` qualifier and the `num` qualifier together determines the total dimension of the data ('n') which2478is loaded from the Tensor Memory. The `shape` qualifier indicates the base dimension of data. The `num` qualifier2479indicates the repeat factor on the base dimension resulting in the total dimension of the data that is accessed.2480 2481Allowed values for the 'num' are `x1, x2, x4, x8, x16, x32, x64, x128`.2482 2483Allowed values for the 'shape' in the first intrinsic are `16x64b, 16x128b, 16x256b, 32x32b`.2484 2485Allowed value for the 'shape' in the second intrinsic is `16x32bx2`.2486 2487The result of the intrinsic is a vector consisting of one or more 32-bit registers derived from `shape` and2488`num` as shown below.2489 2490=========== =========================  ==========  ==========2491 num/shape     16x32bx2/16x64b/32x32b    16x128b    16x256b2492=========== =========================  ==========  ==========2493 x1                 1                      2           42494 x2                 2                      4           82495 x4                 4                      8           162496 x8                 8                      16          322497 x16                16                     32          642498 x32                32                     64          1282499 x64                64                     128         NA2500 x128               128                    NA          NA2501=========== =========================  ==========  ==========2502 2503The 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 load2504 2505For more information, refer to the2506`PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-instructions-tcgen05-ld>`__.2507 2508 2509'``llvm.nvvm.tcgen05.st.*``'2510^^^^^^^^^^^^^^^^^^^^^^^^^^^^2511 2512Syntax:2513"""""""2514 2515.. code-block:: llvm2516 2517  declare void @llvm.nvvm.tcgen05.st.<shape>.<num>(ptr addrspace(6) %tmem_addr, <n x i32> %args, i1 %unpack)2518 2519  declare void @llvm.nvvm.tcgen05.st.16x32bx2.<num>(ptr addrspace(6) %tmem_addr, <n x i32> %args, i64 %offset, i1 %unpack)2520 2521Overview:2522"""""""""2523 2524This group of intrinsics asynchronously store data from the source vector into the Tensor Memory at the location2525specified by the 32-bit address operand 'tmem_addr` collectively across all threads of the warps.2526 2527All the threads in the warp must specify the same value of `tmem_addr`, which must be the base address of the2528collective load operation. Otherwise, the behavior is undefined.2529 2530The `shape` qualifier and the `num` qualifier together determines the total dimension of the data ('n') which2531is loaded from the Tensor Memory. The `shape` qualifier indicates the base dimension of data. The `num` qualifier2532indicates the repeat factor on the base dimension resulting in the total dimension of the data that is accessed.2533 2534Allowed values for the 'num' are `x1, x2, x4, x8, x16, x32, x64, x128`.2535 2536Allowed values for the 'shape' in the first intrinsic are `16x64b, 16x128b, 16x256b, 32x32b`.2537 2538Allowed value for the 'shape' in the second intrinsic is `16x32bx2`.2539 2540`args` argument is a vector consisting of one or more 32-bit registers derived from `shape` and2541`num` as listed in the table listed in the `tcgen05.ld` section.2542 2543Each shape support an `unpack` mode to allow a 32-bit element in the register to be unpacked into two 16-bit elements and store them in adjacent columns. `unpack` mode can be enabled by setting the `%unpack` operand to 1 and can be disabled by setting it to 0.2544 2545The 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.2546 2547For more information, refer to the2548`PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-instructions-tcgen05-st>`__.2549 2550tcgen05.mma Intrinsics2551----------------------2552 2553Overview2554^^^^^^^^2555 2556`tcgen05.mma` operation of shape `M x N x K` perform matrix multiplication and2557accumulation of the form: `D =  A * B + D` where:2558 2559  - the `A` matrix has shape `M x K`, in either `Tensor Memory` or `Shared Memory`2560  - the `B` matrix has shape `K x N`, in `Shared Memory` of the current CTA and, optionally in peer CTA2561  - the `D` matrix is of the shape `M x N`, in `Tensor Memory`2562 2563Optionally an input predicate can be used to disable the input (`%enable_inp_d`)2564from the accumulator matrix and the following operation can be performed as `D = A * B`2565 2566The matrix multiplication and accumulation operations are categorized into various2567kinds based on input types and the throughput of the multiplication operation.2568The following table shows the different kinds of MMA operations that are supported:2569 2570+------------+--------------------------------------------+2571| .kind      | Supported Input Types                      |2572+============+============================================+2573| f16        | F16 and BF16                               |2574+------------+--------------------------------------------+2575| tf32       | TF32                                       |2576+------------+--------------------------------------------+2577| f8f6f4     | All combinations of F8, F6, and F4         |2578+------------+--------------------------------------------+2579| i8         | Signed and Unsigned 8-bit Integers         |2580+------------+--------------------------------------------+2581| mxf8f6f4   | MX-floating point formats                  |2582+------------+--------------------------------------------+2583| mxf4       | MX-floating point formats (FP4)            |2584+------------+--------------------------------------------+2585| mxf4nvf4   | MXF4 + custom NVIDIA 4-bit floating point  |2586|            | (with common scaling factor)               |2587+------------+--------------------------------------------+2588 2589`tcgen05.mma.sp` supports sparse variant of `A` with shape `M x K` stored in packed2590form as `M X (K / 2)` in memory. The `%spmetadata` specifies the mapping of the2591`K / 2` non-zero elements to the `K` elements before performing the MMA operation.2592 2593`tcgen05.mma.block_scale` perform matrix multiplication with block scaling2594`D = (A * scale_A)  * (B * scale_B) + D` where scaling of input matrices from2595memory to form the matrix `A` and matrix `B` before performing the MMA operation.2596Scale factors for `A` and `B` matrices need to be duplicated to all 32 lane partitions2597of tensor memory. The shape of `%scale_a` and `%scale_b` matrices depend on the2598`.scale_vectorsize` described in `here <https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-mma-scale-valid-comb>`__2599 2600The sparsity metadata (`%spmetadata`) as well as the block-scale inputs for `A / B`2601matrices (`%scale_a` and `%scale_b`) reside in Tensor Memory.2602 2603To facilitate opportunistic re-use of `A / B` matrix data across a sequence of MMA2604operations, the `A/B` matrices are loaded into a collector buffer2605(`%collector_usage_a_op_flag`, `%collector_usage_b_buffer_flag`, and `%collector_usage_b_op_flag`).2606The flag value of the collector_usage flag in the intrinsic specifies the nature of the re-use2607 2608There are three kinds of matrix descriptors used by the tcgen05 family of instructions:2609 2610+----------------------------+-----------------------------------------------------------------------------------------------------------+-------------+2611| Descriptor                 | Description                                                                                               | Size (bits) |2612+============================+===========================================================================================================+=============+2613| Shared Memory Descriptor   | Describes properties of multiplicand matrix                                                               |             |2614|                            | in shared memory, including its location                                                                  |             |2615|                            | within the CTA's shared memory.                                                                           |     64      |2616|                            | `PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-shared-memory-descriptor>`__    |             |2617+----------------------------+-----------------------------------------------+-------------+---------------------------------------------+-------------+2618| Instruction Descriptor     | Describes shapes, types, and details of                                                                   |             |2619|                            | all matrices and the MMA operation.                                                                       |     32      |2620|                            | `PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-zero-column-mask-descriptor>`__ |             |2621+----------------------------+-----------------------------------------------+-------------+---------------------------------------------+-------------+2622| Zero-Column Mask Descriptor| Generates a mask specifying which columns of                                                              |             |2623|                            | B matrix are zeroed in the MMA operation,                                                                 |             |2624|                            | regardless of values in shared memory.                                                                    |     64      |2625|                            | Total mask size = N bits                                                                                  |             |2626|                            | `PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-instruction-descriptor>`__      |             |2627+----------------------------+-----------------------------------------------+-------------+---------------------------------------------+-------------+2628 2629`tcgen05.mma` can be used for general matrix multiplication or for convolution operations.2630In case of convolutions, the `activations` can be stored in either matrix `A` or matrix `B`2631while the `weights` will be stored in the other matrix2632 2633`tcgen05.mma` has an optional collector qualifier to specify when an `A` or `B` matrix2634is new to the sequence and should be loaded, unchanged within the sequence and,2635should be reused, or the last use in the sequence and should be discarded.2636The collector qualifier is used to give the TensorCore permission to reuse a2637previously loaded `A` or `B` matrix; however reuse is opportunistic in that the2638TensorCore may reload a matrix even when it has permission to reuse that matrix.2639Thus, the source memory of an A or B matrix must not be modified while the MMA2640instruction using those matrices has not completed - regardless of collector2641qualifier permissions.2642 2643The `cta_group::1` specifies that the operation is performed on the Tensor Memory2644of the executing thread’s CTA only. The `cta_group::2` specifies that the MMA2645operation is performed on the Tensor Memory of the executing thread’s CTA and its peer CTA.2646 2647The vector operand `%disable_output_lane` specifies the lane(s) in the Tensor Memory2648that should be not be updated with the resultant matrix D. Elements of the vector operand2649disable-output-lane forms a mask where each bit corresponds to a lane of the Tensor Memory,2650with least significant bit of the first element of the vector (leftmost in syntax)2651corresponding to the lane 0 of the Tensor Memory. If a bit in the mask is 1, then2652the corresponding lane in the Tensor Memory for the resultant matrix D will not be2653updated2654 2655Intrinsic Design:2656^^^^^^^^^^^^^^^^^2657 2658Given the broad feature set of `tcgen05.mma` instruction modeling these2659through intrinsics is highly complex, and the following table outlines the large2660number of intrinsics required to fully support the `tcgen05.mma` instruction set.2661 2662+------------------------------------+---------------------------------------------------------------------------------------------------+----------------+2663| variant                            | Configuration                                                                                     | Total Variants |2664+====================================+===================================================================================================+================+2665| tcgen05.mma.shared                 | 2 (space) x 2 (sp) x 4 (kind) x 2 (cta_group) x 4 (collector_usage)                               | 128            |2666+------------------------------------+---------------------------------------------------------------------------------------------------+----------------+2667| tcgen05.mma.tensor.ashift          | 2 (sp) x 4 (kind) x 2 (cta_group) x 2 (collector_usage)                                           | 32             |2668+------------------------------------+---------------------------------------------------------------------------------------------------+----------------+2669| tcgen05.mma.scale_d                | 2 (space) x 2 (sp) x 2 (kind) x 2 (cta_group) x 4 (collector_usage)                               | 128            |2670+------------------------------------+---------------------------------------------------------------------------------------------------+----------------+2671| tcgen05.mma.scale_d.tensor.ashift  | 2 (sp) x 2 (kind) x 2 (cta_group) x 2 (collector_usage)                                           | 16             |2672+------------------------------------+---------------------------------------------------------------------------------------------------+----------------+2673| tcgen05.mma.disable_output_lane    | 2 (space) x 2 (sp) x 4 (kind) x 2 (cta_group) x 4 (collector_usage)                               | 128            |2674+------------------------------------+---------------------------------------------------------------------------------------------------+----------------+2675| tcgen05.mma.disable_output_lane... | 2 (sp) x 4 (kind) x 2 (cta_group) x 2 (collector_usage)                                           | 32             |2676+------------------------------------+---------------------------------------------------------------------------------------------------+----------------+2677| tcgen05.mma.block_scale            | 2 (space) x 1 (mxf4nvf4) x 2 (cta_group) x 2 (scale_vec_size) x 4 (collector_usage)               | 32             |2678+------------------------------------+---------------------------------------------------------------------------------------------------+----------------+2679| tcgen05.mma.block_scale            | 2 (space) x 1 (mxf4) x 2 (cta_group) x 2 (scale_vec_size) x 4 (collector_usage)                   | 32             |2680+------------------------------------+---------------------------------------------------------------------------------------------------+----------------+2681| tcgen05.mma.block_scale            | 2 (space) x 1 (mxf8f6f4) x 2 (cta_group) x 2 (scale_vec_size) x 4 (collector_usage)               | 32             |2682+------------------------------------+---------------------------------------------------------------------------------------------------+----------------+2683| tcgen05.mma.ws                     | 2 (space) x 2 (sp) x 4 (kind) x 2 (zero_col_mask) x 4 (collector_usage_op) x 4 (collector_buffer) | 256            |2684+------------------------------------+---------------------------------------------------------------------------------------------------+----------------+2685| Total                              |                                                                                                   | 816            |2686+------------------------------------+---------------------------------------------------------------------------------------------------+----------------+2687 2688 2689To reduce the number of possible intrinsic variations, we've modeled the `tcgen05.mma`2690instructions using flag operands. We've added range checks to these flags to prevent2691invalid values. We also expanded some flags back into intrinsic modifiers to avoid2692supporting invalid combinations of features.2693 2694 2695'``llvm.nvvm.tcgen05.mma.*``'2696^^^^^^^^^^^^^^^^^^^^^^^^^^^^^2697 2698Syntax:2699"""""""2700 2701.. code-block:: llvm2702 2703  declare void @llvm.nvvm.tcgen05.mma.shared(ptr addrspace(6) %d, i64 %adesc, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, i32 %kind_flag, i32 %cta_group_flag, i32 %collector_usage_a_op_flag)2704  declare void @llvm.nvvm.tcgen05.mma.tensor<.ashift>(ptr addrspace(6) %d, ptr addrspace(6) %atensor, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, i32 %kind_flag, i32 %cta_group_flag, i32 %collector_usage_a_op_flag)2705 2706  ; .sp variants2707  declare void @llvm.nvvm.tcgen05.mma.sp.shared(ptr addrspace(6) %d, i64 %adesc, i64 %bdesc, i32 %idesc, ptr addrspace(6) %spmetadata, i1 %enable_inp_d, i32 %kind_flag, i32 %cta_group_flag, i32 %collector_usage_op_flag)2708  declare void @llvm.nvvm.tcgen05.mma.sp.tensor<.ashift>(ptr addrspace(6) %d, ptr addrspace(6) %atensor, i64 %bdesc, i32 %idesc, ptr addrspace(6) %spmetadata, i1 %enable_inp_d, i32 %kind_flag, i32 %cta_group_flag, i32 %collector_usage_a_op_flag)2709 2710  ; .scale_d variants2711  declare void @llvm.nvvm.tcgen05.mma.shared.scale_d(ptr addrspace(6) %d, i64 %adesc, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, i64 %scale_d_imm, i32 %cta_group_flag, i32 %kind_flag, i32 %collector_usage_a_op_flag)2712  declare void @llvm.nvvm.tcgen05.mma.tensor.scale_d<.ashift>(ptr addrspace(6) %d, ptr addrspace(6) %atensor, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, i64 %scale_d_imm, i32 %cta_group_flag, i32 %kind_flag, i32 %collector_usage_a_op_flag)2713 2714  ; sp.scale_d variants2715  declare void @llvm.nvvm.tcgen05.mma.sp.shared.scale_d(ptr addrspace(6) %d, i64 %adesc, i64 %bdesc, i32 %idesc, ptr addrspace(6) %spmetadata, i1 %enable_inp_d, i64 %scale_d_imm, i32 %cta_group_flag, i32 %collector_usage_op_flag)2716  declare void @llvm.nvvm.tcgen05.mma.sp.tensor.scale_d<.ashift>(ptr addrspace(6) %d, ptr addrspace(6) %atensor, i64 %bdesc, i32 %idesc, ptr addrspace(6) %spmetadata, i1 %enable_inp_d, i64 %scale_d_imm, i32 %cta_group, i32 %collector_usage_a_op_flag)2717 2718Overview:2719"""""""""2720 2721`nvvm.tcgen05.mma` is an asynchronous intrinsic which initiates an `M x N x K` matrix2722multiply and accumulate operation, `D = A * B + D` where the `A` matrix is `M x K`,2723the `B` matrix is `K x N`, and the `D` matrix is `M x N`. The operation of the form2724`D = A*B` is issued when the input predicate argument `%enable_inp_d` is false.2725The optional immediate argument `%scale_d_imm` can be specified to scale the input2726matrix `D` as follows: `D = A * B + D * (2 ^ - %scale_d_imm)`. The valid range of2727values for argument `%scale_d_imm` is `[0, 15]`. The 32-bit register operand idesc2728is the instruction descriptor as described in `Instruction descriptor <https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-instruction-descriptor>`__2729 2730`nvvm.tcgen05.mma` has single thread semantics, unlike the collective instructions2731`nvvm.mma.sync` or the PTX `wgmma.mma_async` instruction. So, a single thread issuing2732the `nvvm.tcgen05.mma` will result in the initiation of the whole matrix and accumulate2733operation2734 2735When `.sp` is specifed, the dimension of A matrix is `M x (K/2)` and requires2736specifying an additional `%spmetadata` argument2737 2738`.ashift` shifts the rows of the A matrix down by one row, except for the last row2739in the Tensor Memory. `.ashift` is only allowed with M = 128 or M = 256.2740 2741The `%collector_usage_a_op_flag` flag specifies the usage of collector buffer for2742matrix `A`. It is illegal to specify either of `USE` or `FILL` for `%collector_usage_a_op_flag`2743along with `.ashift`2744 2745For more information, refer to the2746`PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-mma-instructions-mma>`__2747 2748The following tables describe the possible values of the flag arguments2749 2750`%kind_flag` flag:2751 2752============= ==========2753  `kind_flag`   value2754============= ==========2755     F16          02756     TF32         12757     F8F6F4       22758     I8           32759============= ==========2760 2761`%cta_group_flag` flag:2762 2763================= ==========2764 `cta_group_flag`    value2765================= ==========2766     CG1               12767     CG2               22768================= ==========2769 2770`%collector_usage_a_op_flag` flag:2771 2772============================= ==========2773 `collector_usage_a_op_flag`    value2774============================= ==========2775           DISCARD                 02776           LASTUSE                 12777           USE                     22778           FILL                    32779============================= ==========2780 2781'``llvm.nvvm.tcgen05.mma.block_scale*``'2782^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^2783 2784Syntax:2785"""""""2786 2787.. code-block:: llvm2788 2789  ; mxf8f6f42790  declare void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale(ptr addrspace(6) %d, i64 %adesc, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, i32 cta_group_flag, i32 %collector_usage_a_op_flag)2791  declare void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale(ptr addrspace(6) %d, ptr addrspace(6) %atensor, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, i32 cta_group_flag, i32 %collector_usage_a_op_flag)2792  declare void @llvm.nvvm.tcgen05.mma.shared.mxf8f6f4.block_scale.block32(ptr addrspace(6) %d, i64 %adesc, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, i32 cta_group_flag, i32 %collector_usage_a_op_flag)2793  declare void @llvm.nvvm.tcgen05.mma.tensor.mxf8f6f4.block_scale.block32(ptr addrspace(6) %d, ptr addrspace(6) %atensor, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, i32 cta_group_flag, i32 %collector_usage_a_op_flag)2794  declare void @llvm.nvvm.tcgen05.mma.sp.shared.mxf8f6f4.block_scale(ptr addrspace(6) %d, i64 %adesc, i64 %bdesc, i32 %idesc, ptr addrspace(6) %spmetadata, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, i32 cta_group_flag, i32 %collector_usage_a_op_flag)2795  declare void @llvm.nvvm.tcgen05.mma.sp.tensor.mxf8f6f4.block_scale(ptr addrspace(6) %d, ptr addrspace(6) %atensor, i64 %bdesc, i32 %idesc, ptr addrspace(6) %spmetadata, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, i32 cta_group_flag, i32 %collector_usage_a_op_flag)2796  declare void @llvm.nvvm.tcgen05.mma.sp.shared.mxf8f6f4.block_scale.block32(ptr addrspace(6) %d, i64 %adesc, i64 %bdesc, i32 %idesc, ptr addrspace(6) %spmetadata, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, i32 cta_group_flag, i32 %collector_usage_a_op_flag)2797  declare void @llvm.nvvm.tcgen05.mma.sp.tensor.mxf8f6f4.block_scale.block32(ptr addrspace(6) %d, ptr addrspace(6) %atensor, i64 %bdesc, i32 %idesc, ptr addrspace(6) %spmetadata, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, i32 cta_group_flag, i32 %collector_usage_a_op_flag)2798 2799  ; mxf42800  declare void @llvm.nvvm.tcgen05.mma.shared.mxf4.block_scale(ptr addrspace(6) %d, i64 %adesc, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, i32 cta_group_flag, i32 %collector_usage_a_op_flag)2801  declare void @llvm.nvvm.tcgen05.mma.tensor.mxf4.block_scale(ptr addrspace(6) %d, ptr addrspace(6) %atensor, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, i32 cta_group_flag, i32 %collector_usage_a_op_flag)2802  declare void @llvm.nvvm.tcgen05.mma.shared.mxf4.block_scale.block32(ptr addrspace(6) %d, i64 %adesc, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, i32 cta_group_flag, i32 %collector_usage_a_op_flag)2803  declare void @llvm.nvvm.tcgen05.mma.tensor.mxf4.block_scale.block32(ptr addrspace(6) %d, ptr addrspace(6) %atensor, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, i32 cta_group_flag, i32 %collector_usage_a_op_flag)2804  declare void @llvm.nvvm.tcgen05.mma.sp.shared.mxf4.block_scale(ptr addrspace(6) %d, i64 %adesc, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %spmetadata, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, i32 cta_group_flag, i32 %collector_usage_a_op_flag)2805  declare void @llvm.nvvm.tcgen05.mma.sp.tensor.mxf4.block_scale(ptr addrspace(6) %d, ptr addrspace(6) %atensor, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %spmetadata, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, i32 cta_group_flag, i32 %collector_usage_a_op_flag)2806  declare void @llvm.nvvm.tcgen05.mma.sp.shared.mxf4.block_scale.block32(ptr addrspace(6) %d, i64 %adesc, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %spmetadata, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, i32 cta_group_flag, i32 %collector_usage_a_op_flag)2807  declare void @llvm.nvvm.tcgen05.mma.sp.tensor.mxf4.block_scale.block32(ptr addrspace(6) %d, ptr addrspace(6) %atensor, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %spmetadata, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, i32 cta_group_flag, i32 %collector_usage_a_op_flag)2808 2809  ; mxf4nvf42810  declare void @llvm.nvvm.tcgen05.mma.shared.mxf4nvf4.block_scale.block16(ptr addrspace(6) %d, i64 %adesc, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, i32 cta_group_flag, i32 %collector_usage_a_op_flag)2811  declare void @llvm.nvvm.tcgen05.mma.tensor.mxf4nvf4.block_scale.block16(ptr addrspace(6) %d, ptr addrspace(6) %atensor, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, i32 cta_group_flag, i32 %collector_usage_a_op_flag)2812  declare void @llvm.nvvm.tcgen05.mma.shared.mxf4nvf4.block_scale.block32(ptr addrspace(6) %d, i64 %adesc, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, i32 cta_group_flag, i32 %collector_usage_a_op_flag)2813  declare void @llvm.nvvm.tcgen05.mma.tensor.mxf4nvf4.block_scale.block32(ptr addrspace(6) %d, ptr addrspace(6) %atensor, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, i32 cta_group_flag, i32 %collector_usage_a_op_flag)2814  declare void @llvm.nvvm.tcgen05.mma.sp.shared.mxf4nvf4.block_scale.block16(ptr addrspace(6) %d, i64 %adesc, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %spmetadata, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, i32 cta_group_flag, i32 %collector_usage_a_op_flag)2815  declare void @llvm.nvvm.tcgen05.mma.sp.tensor.mxf4nvf4.block_scale.block16(ptr addrspace(6) %d, ptr addrspace(6) %atensor, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %spmetadata, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, i32 cta_group_flag, i32 %collector_usage_a_op_flag)2816  declare void @llvm.nvvm.tcgen05.mma.sp.shared.mxf4nvf4.block_scale.block32(ptr addrspace(6) %d, i64 %adesc, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %spmetadata, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, i32 cta_group_flag, i32 %collector_usage_a_op_flag)2817  declare void @llvm.nvvm.tcgen05.mma.sp.tensor.mxf4nvf4.block_scale.block32(ptr addrspace(6) %d, ptr addrspace(6) %atensor, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %spmetadata, ptr addrspace(6) %scale_a, ptr addrspace(6) %scale_b, i32 cta_group_flag, i32 %collector_usage_a_op_flag)2818 2819Overview:2820"""""""""2821`nvvm.tcgen05.mma.block_scale` is an asynchronous intrinsic which initiates an `M x N x K` matrix multiply and accumulate operation, `D = (A * scale_a)  * (B * scale_b) + D` where the `A` matrix is `M x K`, the `B` matrix is `K x N`, and the `D` matrix is `M x N`. The matrices `A` and `B` are scaled with `%scale_A` and `%scale_B` matrices respectively before performing the matrix multiply and accumulate operation. The operation of the form `D = A*B` is issued when the input predicate argument `%enable_inp_d` is false. The 32-bit register operand idesc is the instruction descriptor as described in `Instruction descriptor <https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-instruction-descriptor>`__2822 2823`nvvm.tcgen05.mma.block_scale` has single thread semantics, unlike the collective instructions `nvvm.mma.sync` or the PTX `wgmma.mma_async` instruction. So, a single thread issuing the `nvvm.tcgen05.mma.block_scale` will result in the initiation of the whole matrix multiply and accumulate operation2824 2825When `.sp` is specified, the dimension of A matrix is `M x (K / 2)` and requires specifying an additional `%spmetadata` argument2826 2827The `%collector_usage_a_op_flag` flag specifies the usage of collector buffer for matrix `A`2828 2829For more information, refer to the2830`PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-mma-instructions-mma>`__2831 2832The following tables describe the possible values of the flag arguments2833 2834`%cta_group`:2835 2836============= ==========2837 `cta_group`    value2838============= ==========2839     CG1          12840     CG2          22841============= ==========2842 2843`%collector_usage_a_op_flag`:2844 2845============================= ==========2846 `collector_usage_a_op_flag`    value2847============================= ==========2848     DISCARD                      02849     LASTUSE                      12850     USE                          22851     FILL                         32852============================= ==========2853 2854'``llvm.nvvm.tcgen05.mma.disable_output_lane*``'2855^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^2856 2857Syntax:2858"""""""2859 2860.. code-block:: llvm2861 2862  declare void @llvm.nvvm.tcgen05.mma.shared.disable_output_lane.cg1(ptr addrspace(6) %d, i64 %adesc, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, <4 x i32> %disable_output_lane_v4, i32 %kind_flag, i32 %collector_usage_a_op_flag)2863  declare void @llvm.nvvm.tcgen05.mma.shared.disable_output_lane.cg2(ptr addrspace(6) %d, i64 %adesc, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, <8 x i32> %disable_output_lane_v8, i32 %kind_flag, i32 %collector_usage_a_op_flag)2864  declare void @llvm.nvvm.tcgen05.mma.tensor.disable_output_lane.cg1<.ashift>(ptr addrspace(6) %d, ptr addrspace(6) %atensor, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, <4 x i32> %disable_output_lane_v4, i32 %kind_flag, i32 %collector_usage_a_op_flag)2865  declare void @llvm.nvvm.tcgen05.mma.tensor.disable_output_lane.cg2<.ashift>(ptr addrspace(6) %d, ptr addrspace(6) %atensor, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, <8 x i32> %disable_output_lane_v8, i32 %kind_flag, i32 %collector_usage_a_op_flag)2866 2867  ; .sp variants2868  declare void @llvm.nvvm.tcgen05.mma.sp.shared.disable_output_lane.cg1(ptr addrspace(6) %d, i64 %adesc, i64 %bdesc, i32 %idesc, ptr addrspace(6) %spmetadata, i1 %enable_inp_d, <4 x i32> %disable_output_lane_v4, i32 %kind_flag, i32 %collector_usage_op_flag)2869  declare void @llvm.nvvm.tcgen05.mma.sp.shared.disable_output_lane.cg2(ptr addrspace(6) %d, i64 %adesc, i64 %bdesc, i32 %idesc, ptr addrspace(6) %spmetadata, i1 %enable_inp_d, <8 x i32> %disable_output_lane_v8, i32 %kind_flag, i32 %collector_usage_op_flag)2870  declare void @llvm.nvvm.tcgen05.mma.sp.tensor.disable_output_lane.cg1<.ashift>(ptr addrspace(6) %d, ptr addrspace(6) %atensor, i64 %bdesc, i32 %idesc, ptr addrspace(6) %spmetadata, i1 %enable_inp_d, <4 x i32> %disable_output_lane_v4, i32 %kind_flag, i32 %collector_usage_a_op_flag)2871  declare void @llvm.nvvm.tcgen05.mma.sp.tensor.disable_output_lane.cg2<.ashift>(ptr addrspace(6) %d, ptr addrspace(6) %atensor, i64 %bdesc, i32 %idesc, ptr addrspace(6) %spmetadata, i1 %enable_inp_d, <8 x i32> %disable_output_lane_v8, i32 %kind_flag, i32 %collector_usage_a_op_flag)2872 2873  ; .scale_d variants2874  declare void @llvm.nvvm.tcgen05.mma.shared.scale_d.disable_output_lane.cg1(ptr addrspace(6) %d, i64 %adesc, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, i64 %scale_d_imm, <4 x i32> %disable_output_lane_v4, i32 %kind_flag, i32 %collector_usage_a_op_flag)2875  declare void @llvm.nvvm.tcgen05.mma.shared.scale_d.disable_output_lane.cg2(ptr addrspace(6) %d, i64 %adesc, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, i64 %scale_d_imm, <8 x i32> %disable_output_lane_v8, i32 %kind_flag, i32 %collector_usage_a_op_flag)2876  declare void @llvm.nvvm.tcgen05.mma.tensor.scale_d.disable_output_lane.cg1<.ashift>(ptr addrspace(6) %d, ptr addrspace(6) %atensor, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, i64 %scale_d_imm, <4 x i32> %disable_output_lane_v4, i32 %kind_flag, i32 %collector_usage_a_op_flag)2877  declare void @llvm.nvvm.tcgen05.mma.tensor.scale_d.disable_output_lane.cg2<.ashift>(ptr addrspace(6) %d, ptr addrspace(6) %atensor, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, i64 %scale_d_imm, <8 x i32> %disable_output_lane_v8, i32 %kind_flag, i32 %collector_usage_a_op_flag)2878 2879  ; .sp.scale_d variants2880  declare void @llvm.nvvm.tcgen05.mma.sp.shared.scale_d.disable_output_lane.cg1(ptr addrspace(6) %d, i64 %adesc, i64 %bdesc, i32 %idesc, ptr addrspace(6) %spmetadata, i1 %enable_inp_d, i64 %scale_d_imm, <4 x i32> %disable_output_lane_v4, i32 %kind_flag, i32 %collector_usage_op_flag)2881  declare void @llvm.nvvm.tcgen05.mma.sp.shared.scale_d.disable_output_lane.cg2(ptr addrspace(6) %d, i64 %adesc, i64 %bdesc, i32 %idesc, ptr addrspace(6) %spmetadata, i1 %enable_inp_d, i64 %scale_d_imm, <8 x i32> %disable_output_lane_v8, i32 %kind_flag, i32 %collector_usage_op_flag)2882  declare void @llvm.nvvm.tcgen05.mma.sp.tensor.scale_d.disable_output_lane.cg1<.ashift>(ptr addrspace(6) %d, ptr addrspace(6) %atensor, i64 %bdesc, i32 %idesc, ptr addrspace(6) %spmetadata, i1 %enable_inp_d, i64 %scale_d_imm, <4 x i32> %disable_output_lane_v4, i32 %kind_flag, i32 %collector_usage_a_op_flag)2883  declare void @llvm.nvvm.tcgen05.mma.sp.tensor.scale_d.disable_output_lane.cg2<.ashift>(ptr addrspace(6) %d, ptr addrspace(6) %atensor, i64 %bdesc, i32 %idesc, ptr addrspace(6) %spmetadata, i1 %enable_inp_d, i64 %scale_d_imm, <8 x i32> %disable_output_lane_v8, i32 %kind_flag, i32 %collector_usage_a_op_flag)2884 2885Overview:2886"""""""""2887 2888`nvvm.tcgen05.mma.disable_output_lane` is an asynchronous intrinsic which initiates an `M x N x K` matrix multiply and accumulate operation, `D = A * B + D` where the `A` matrix is `M x K`, the `B` matrix is `K x N`, and the `D` matrix is `M x N`. The operation of the form `D = A*B` is issued when the input predicate argument `%enable_inp_d` is false. The optional immediate argument `%scale_d_imm` can be specified to scale the input matrix `D` as follows: `D = A*B+D * (2 ^ - %scale_d_imm)`. The valid range of values for argument `%scale_d_imm` is `[0, 15]`. The 32-bit register operand idesc is the instruction descriptor as described in `Instruction descriptor <https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-instruction-descriptor>`__2889 2890The vector operand `%disable_output_lane` specifies the lane(s) in the Tensor Memory that should be not be updated with the resultant matrix `D`. Elements of the vector operand `%disable_output_lane` forms a mask where each bit corresponds to a lane of the Tensor Memory, with least significant bit of the first element of the vector corresponding to the `lane 0` of the Tensor Memory. If a bit in the mask is 1, then the corresponding lane in the Tensor Memory for the resultant matrix `D` will not be updated2891 2892`nvvm.tcgen05.mma.disable_output_lane` has single thread semantics, unlike the collective instructions `nvvm.mma.sync` or the PTX `wgmma.mma_async` instruction. So, a single thread issuing the `nvvm.tcgen05.mma.disable_output_lane` will result in the initiation of the whole matrix multiply and accumulate operation2893 2894When `.sp` is specifed, the dimension of A matrix is `M x (K / 2)` and requires specifiying an additional `%spmetadata` argument2895 2896`.ashift` shifts the rows of the A matrix down by one row, except for the last row in the Tensor Memory. `.ashift` is only allowed with M = 128 or M = 256.2897 2898The `%collector_usage_a_op_flag` flag specifies the usage of collector buffer for matrix `A`. It is illegal to specify either of `USE` or `FILL` for `%collector_usage_a_op_flag` along with `.ashift`2899 2900For more information, refer to the `PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-mma-instructions-mma>`__2901 2902The following tables describes the possible values of the flag arguments2903 2904`%kind_flag`:2905 2906============= ==========2907 `kind_flag`    value2908============= ==========2909     F16          02910     TF32         12911     F8F6F4       22912     I8           32913============= ==========2914 2915`%cta_group_flag`:2916 2917================= ==========2918 `cta_group_flag`    value2919================= ==========2920        CG1           12921        CG2           22922================= ==========2923 2924`%collector_usage_a_op_flag`:2925 2926============================= ==========2927 `collector_usage_a_op_flag`    value2928============================= ==========2929     DISCARD                      02930     LASTUSE                      12931     USE                          22932     FILL                         32933============================= ==========2934 2935 2936'``llvm.nvvm.tcgen05.mma.ws*``'2937^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^2938 2939Syntax:2940"""""""2941 2942.. code-block:: llvm2943 2944  // tcgen05.mma.ws2945  declare void @llvm.nvvm.tcgen05.mma.ws.shared(ptr addrspace(6) %d, i64 %adesc, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, i32 %kind_flag, i32 %collector_usage_b_buffer_flag, i32 %collector_usage_b_op_flag)2946  declare void @llvm.nvvm.tcgen05.mma.ws.tensor(ptr addrspace(6) %d, ptr addrspace(6) %atensor, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, i32 %kind_flag, i32 %collector_usage_b_buffer_flag, i32 %collector_usage_b_op_flag)2947  declare void @llvm.nvvm.tcgen05.mma.ws.shared.zero_col_mask(ptr addrspace(6) %d, i64 %adesc, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, i64 %zero_col_mask, i32 %kind_flag, i32 %collector_usage_b_buffer_flag, i32 %collector_usage_b_op_flag)2948  declare void @llvm.nvvm.tcgen05.mma.ws.shared.zero_col_mask(ptr addrspace(6) %d, ptr addrspace(6) %atensor, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, i64 %zero_col_mask, i32 %kind_flag, i32 %collector_usage_b_buffer_flag, i32 %collector_usage_b_op_flag)2949  declare void @llvm.nvvm.tcgen05.mma.ws.tensor.zero_col_mask(ptr addrspace(6) %d, ptr addrspace(6) %atensor, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, i64 %zero_col_mask, i32 %kind_flag, i32 %collector_usage_b_buffer_flag, i32 %collector_usage_b_op_flag)2950 2951  ; .sp variants2952  declare void @llvm.nvvm.tcgen05.mma.ws.sp.shared(ptr addrspace(6) %d, i64 %adesc, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %spmetadata, i32 %kind_flag, i32 %collector_usage_b_buffer_flag, i32 %collector_usage_b_op_flag)2953  declare void @llvm.nvvm.tcgen05.mma.ws.sp.tensor(ptr addrspace(6) %d, ptr addrspace(6) %atensor, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %spmetadata, i32 %kind_flag, i32 %collector_usage_b_buffer_flag, i32 %collector_usage_b_op_flag)2954  declare void @llvm.nvvm.tcgen05.mma.ws.sp.shared.zero_col_mask(ptr addrspace(6) %d, i64 %adesc, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %spmetadata, i64 %zero_col_mask, i32 %kind_flag, i32 %collector_usage_b_buffer_flag, i32 %collector_usage_b_op_flag)2955  declare void @llvm.nvvm.tcgen05.mma.ws.sp.tensor.zero_col_mask(ptr addrspace(6) %d, ptr addrspace(6) %atensor, i64 %bdesc, i32 %idesc, i1 %enable_inp_d, ptr addrspace(6) %spmetadata, i64 %zero_col_mask, i32 %kind_flag, i32 %collector_usage_b_buffer_flag, i32 %collector_usage_b_op_flag)2956 2957Overview:2958"""""""""2959 2960`nvvm.tcgen05.mma.ws` is an asynchronous intrinsic which initiates an `M x N x K` weight stationary convolution matrix multiply and accumulate operation, `D = A * B + D` where the `A` matrix is `M x K`, the `B` matrix is `K x N`, and the `D` matrix is `M x N`. The operation of the form `D = A*B` is issued when the input predicate argument `%enable_inp_d` is false. The optional immediate argument `%scale_d_imm` can be specified to scale the input matrix `D` as follows: `D = A*B+D * (2 ^ - %scale_d_imm)`. The valid range of values for argument `%scale_d_imm` is `[0, 15]`. The 32-bit register operand idesc is the instruction descriptor as described in `Instruction descriptor <https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-instruction-descriptor>`__2961 2962`nvvm.tcgen05.mma` has single thread semantics, unlike the collective instructions `nvvm.mma.sync` or the PTX `wgmma.mma_async` instruction. So, a single thread issuing the `nvvm.tcgen05.mma` will result in the initiation of the whole matrix multiply and accumulate operation2963 2964When `.sp` is specifed, the dimension of A matrix is `M x (K / 2)` and requires specifiying an additional `%spmetadata` argument2965 2966The operand `%zero_col_mask` is a 64-bit register which specifies the `Zero-Column Mask Descriptor <https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-zero-column-mask-descriptor>`__. The zero-column mask descriptor is used to generate a mask that specifies which columns of `B` matrix will have zero value for the matrix multiply and accumulate operation regardless of the values present in the shared memory.2967 2968The `%collector_usage_b_buffer_flag` and `%collector_usage_b_op_flag` together flag specifies the usage of collector buffer for Matrix `B`2969 2970For more information, refer to the2971`PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-mma-instructions-mma-ws>`__2972 2973The following tables describes the possible values of the flag arguments2974 2975`%kind_flag`:2976 2977============= ==========2978 `kind_flag`    value2979============= ==========2980     F16          02981     TF32         12982     F8F6F4       22983     I8           32984============= ==========2985 2986`%collector_usage_b_buffer_flag`:2987 2988================================ ==========2989 `collector_usage_b_buffer_flag`   value2990================================ ==========2991              B0                     02992              B1                     12993              B2                     22994              B3                     32995================================ ==========2996 2997`%collector_usage_b_op_flag`:2998 2999============================= ==========3000 `collector_usage_b_op_flag`    value3001============================= ==========3002     DISCARD                      03003     LASTUSE                      13004     USE                          23005     FILL                         33006============================= ==========3007 3008Store Intrinsics3009----------------3010 3011'``llvm.nvvm.st.bulk.*``'3012^^^^^^^^^^^^^^^^^^^^^^^^^3013 3014Syntax:3015"""""""3016 3017.. code-block:: llvm3018 3019  declare void @llvm.nvvm.st.bulk(ptr addrspace(1) %dst, i64 %size, i64 immarg %initval)3020  declare void @llvm.nvvm.st.bulk.shared.cta(ptr addrspace(3) %dst, i64 %size, i64 immarg %initval)3021 3022Overview:3023"""""""""3024 3025The '``@llvm.nvvm.st.bulk.*``' intrinsics initialize a region of shared memory 3026starting from the location specified by the destination address operand `%dst`.3027 3028The integer operand `%size` specifies the amount of memory to be initialized in 3029terms of number of bytes and must be a multiple of 8. Otherwise, the behavior 3030is undefined.3031 3032The integer immediate operand `%initval` specifies the initialization value for 3033the memory locations. The only numeric value allowed is 0.3034 3035The ``@llvm.nvvm.st.bulk.shared.cta`` and ``@llvm.nvvm.st.bulk`` intrinsics are 3036similar but the latter uses generic addressing (see `Generic Addressing <https://docs.nvidia.com/cuda/parallel-thread-execution/#generic-addressing>`__).3037 3038For more information, refer `PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-st-bulk>`__.3039 3040 3041clusterlaunchcontrol Intrinsics3042-------------------------------3043 3044'``llvm.nvvm.clusterlaunchcontrol.try_cancel*``' Intrinsics3045^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^3046 3047Syntax:3048"""""""3049 3050.. code-block:: llvm3051 3052  declare void @llvm.nvvm.clusterlaunchcontrol.try_cancel.async.shared(ptr addrspace(3) %addr, ptr addrspace(3) %mbar)3053  declare void @llvm.nvvm.clusterlaunchcontrol.try_cancel.async.multicast.shared(ptr addrspace(3) %addr, ptr addrspace(3) %mbar)3054 3055Overview:3056"""""""""3057 3058The ``clusterlaunchcontrol.try_cancel`` intrinsics requests atomically cancelling3059the launch of a cluster that has not started running yet. It asynchronously non-atomically writes3060a 16-byte opaque response to shared memory, pointed to by 16-byte-aligned ``addr`` indicating whether the3061operation succeeded or failed. ``addr`` and 8-byte-aligned ``mbar`` must refer to ``shared::cta``3062otherwise the behavior is undefined. The completion of the asynchronous operation3063is tracked using the mbarrier completion mechanism at ``.cluster`` scope referenced3064by the shared memory pointer, ``mbar``. On success, the opaque response contains3065the CTA id of the first CTA of the canceled cluster; no other successful response3066from other ``clusterlaunchcontrol.try_cancel`` operations from the same grid will3067contain that id.3068 3069The ``multicast`` variant specifies that the response is asynchronously non-atomically written to3070the corresponding shared memory location of each CTA in the requesting cluster.3071The completion of the write of each local response is tracked by independent3072mbarriers at the corresponding shared memory location of each CTA in the3073cluster.3074 3075For more information, refer `PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/?a#parallel-synchronization-and-communication-instructions-clusterlaunchcontrol-try-cancel>`__.3076 3077'``llvm.nvvm.clusterlaunchcontrol.query_cancel.is_canceled``' Intrinsic3078^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^3079 3080Syntax:3081"""""""3082 3083.. code-block:: llvm3084 3085  declare i1 @llvm.nvvm.clusterlaunchcontrol.query_cancel.is_canceled(i128 %try_cancel_response)3086 3087Overview:3088"""""""""3089 3090The ``llvm.nvvm.clusterlaunchcontrol.query_cancel.is_canceled`` intrinsic decodes the opaque response written by the3091``llvm.nvvm.clusterlaunchcontrol.try_cancel`` operation.3092 3093The intrinsic returns ``0`` (false) if the request failed. If the request succeeded,3094it returns ``1`` (true). A true result indicates that:3095 3096- the thread block cluster whose first CTA id matches that of the response3097  handle will not run, and3098- no other successful response of another ``try_cancel`` request in the grid will contain3099  the first CTA id of that cluster3100 3101For more information, refer `PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/?a#parallel-synchronization-and-communication-instructions-clusterlaunchcontrol-query-cancel>`__.3102 3103 3104'``llvm.nvvm.clusterlaunchcontrol.query_cancel.get_first_ctaid.*``' Intrinsics3105^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^3106 3107Syntax:3108"""""""3109 3110.. code-block:: llvm3111 3112  declare i32 @llvm.nvvm.clusterlaunchcontrol.query_cancel.get_first_ctaid.x(i128 %try_cancel_response)3113  declare i32 @llvm.nvvm.clusterlaunchcontrol.query_cancel.get_first_ctaid.y(i128 %try_cancel_response)3114  declare i32 @llvm.nvvm.clusterlaunchcontrol.query_cancel.get_first_ctaid.z(i128 %try_cancel_response)3115 3116Overview:3117"""""""""3118 3119The ``clusterlaunchcontrol.query_cancel.get_first_ctaid.*`` intrinsic can be3120used to decode the successful opaque response written by the3121``llvm.nvvm.clusterlaunchcontrol.try_cancel`` operation.3122 3123If the request succeeded:3124 3125- ``llvm.nvvm.clusterlaunchcontrol.query_cancel.get_first_ctaid.{x,y,z}`` returns3126  the coordinate of the first CTA in the canceled cluster, either x, y, or z.3127 3128If the request failed, the behavior of these intrinsics is undefined.3129 3130For more information, refer `PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/?a#parallel-synchronization-and-communication-instructions-clusterlaunchcontrol-query-cancel>`__.3131 3132Perf Monitor Event Intrinsics3133-----------------------------3134 3135'``llvm.nvvm.pm.event.mask``' Intrinsic3136^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^3137 3138Syntax:3139"""""""3140 3141.. code-block:: llvm3142 3143    declare void @llvm.nvvm.pm.event.mask(i16 immarg %mask_val)3144 3145Overview:3146"""""""""3147 3148The '``llvm.nvvm.pm.event.mask``' intrinsic triggers one or more3149performance monitor events. Each bit in the 16-bit immediate operand3150``%mask_val`` controls an event.3151 3152For more information on the pmevent instructions, refer to the PTX ISA3153`<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#miscellaneous-instructions-pmevent>`_.3154 3155Other Intrinsics3156----------------3157 3158For the full set of NVPTX intrinsics, please see the3159``include/llvm/IR/IntrinsicsNVVM.td`` file in the LLVM source tree.3160 3161 3162.. _libdevice:3163 3164Linking with Libdevice3165======================3166 3167The CUDA Toolkit comes with an LLVM bitcode library called ``libdevice`` that3168implements many common mathematical functions. This library can be used as a3169high-performance math library for any compilers using the LLVM NVPTX target.3170The library can be found under ``nvvm/libdevice/`` in the CUDA Toolkit and3171there is a separate version for each compute architecture.3172 3173For a list of all math functions implemented in libdevice, see3174`libdevice Users Guide <http://docs.nvidia.com/cuda/libdevice-users-guide/index.html>`_.3175 3176To accommodate various math-related compiler flags that can affect code3177generation of libdevice code, the library code depends on a special LLVM IR3178pass (``NVVMReflect``) to handle conditional compilation within LLVM IR. This3179pass looks for calls to the ``@__nvvm_reflect`` function and replaces them3180with constants based on the defined reflection parameters. Such conditional3181code often follows a pattern:3182 3183.. code-block:: c++3184 3185  float my_function(float a) {3186    if (__nvvm_reflect("FASTMATH"))3187      return my_function_fast(a);3188    else3189      return my_function_precise(a);3190  }3191 3192The default value for all unspecified reflection parameters is zero.3193 3194The ``NVVMReflect`` pass should be executed early in the optimization3195pipeline, immediately after the link stage. The ``internalize`` pass is also3196recommended to remove unused math functions from the resulting PTX. For an3197input IR module ``module.bc``, the following compilation flow is recommended:3198 3199The ``NVVMReflect`` pass will attempt to remove dead code even without3200optimizations. This allows potentially incompatible instructions to be avoided3201at all optimizations levels by using the ``__CUDA_ARCH`` argument.3202 32031. Save list of external functions in ``module.bc``32042. Link ``module.bc`` with ``libdevice.compute_XX.YY.bc``32053. Internalize all functions not in list from (1)32064. Eliminate all unused internal functions32075. Run ``NVVMReflect`` pass32086. Run standard optimization pipeline3209 3210.. note::3211 3212  ``linkonce`` and ``linkonce_odr`` linkage types are not suitable for the3213  libdevice functions. It is possible to link two IR modules that have been3214  linked against libdevice using different reflection variables.3215 3216Since the ``NVVMReflect`` pass replaces conditionals with constants, it will3217often leave behind dead code of the form:3218 3219.. code-block:: llvm3220 3221  entry:3222    ..3223    br i1 true, label %foo, label %bar3224  foo:3225    ..3226  bar:3227    ; Dead code3228    ..3229 3230Therefore, it is recommended that ``NVVMReflect`` is executed early in the3231optimization pipeline before dead-code elimination.3232 3233The NVPTX TargetMachine knows how to schedule ``NVVMReflect`` at the beginning3234of your pass manager; just use the following code when setting up your pass3235manager and the PassBuilder will use ``registerPassBuilderCallbacks`` to let3236NVPTXTargetMachine::registerPassBuilderCallbacks add the pass to the3237pass manager:3238 3239.. code-block:: c++3240 3241    std::unique_ptr<TargetMachine> TM = ...;3242    PassBuilder PB(TM);3243    ModulePassManager MPM;3244    PB.parsePassPipeline(MPM, ...);3245 3246Reflection Parameters3247---------------------3248 3249The libdevice library currently uses the following reflection parameters to3250control code generation:3251 3252==================== ======================================================3253Flag                 Description3254==================== ======================================================3255``__CUDA_FTZ=[0,1]`` Use optimized code paths that flush subnormals to zero3256==================== ======================================================3257 3258The value of this flag is determined by the "nvvm-reflect-ftz" module flag.3259The following sets the ftz flag to 1.3260 3261.. code-block:: llvm3262 3263    !llvm.module.flags = !{!0}3264    !0 = !{i32 4, !"nvvm-reflect-ftz", i32 1}3265 3266(``i32 4`` indicates that the value set here overrides the value in another3267module we link with.  See the `LangRef <LangRef.html#module-flags-metadata>`3268for details.)3269 3270Executing PTX3271=============3272 3273The most common way to execute PTX assembly on a GPU device is to use the CUDA3274Driver API. This API is a low-level interface to the GPU driver and allows for3275JIT compilation of PTX code to native GPU machine code.3276 3277Initializing the Driver API:3278 3279.. code-block:: c++3280 3281    CUdevice device;3282    CUcontext context;3283 3284    // Initialize the driver API3285    cuInit(0);3286    // Get a handle to the first compute device3287    cuDeviceGet(&device, 0);3288    // Create a compute device context3289    cuCtxCreate(&context, 0, device);3290 3291JIT compiling a PTX string to a device binary:3292 3293.. code-block:: c++3294 3295    CUmodule module;3296    CUfunction function;3297 3298    // JIT compile a null-terminated PTX string3299    cuModuleLoadData(&module, (void*)PTXString);3300 3301    // Get a handle to the "myfunction" kernel function3302    cuModuleGetFunction(&function, module, "myfunction");3303 3304For full examples of executing PTX assembly, please see the `CUDA Samples3305<https://developer.nvidia.com/cuda-downloads>`_ distribution.3306 3307 3308Common Issues3309=============3310 3311ptxas complains of undefined function: __nvvm_reflect3312-----------------------------------------------------3313 3314When linking with libdevice, the ``NVVMReflect`` pass must be used. See3315:ref:`libdevice` for more information.3316 3317 3318Tutorial: A Simple Compute Kernel3319=================================3320 3321To start, let us take a look at a simple compute kernel written directly in3322LLVM IR. The kernel implements vector addition, where each thread computes one3323element of the output vector C from the input vectors A and B.  To make this3324easier, we also assume that only a single CTA (thread block) will be launched,3325and that it will be one dimensional.3326 3327 3328The Kernel3329----------3330 3331.. code-block:: llvm3332 3333  target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"3334  target triple = "nvptx64-nvidia-cuda"3335 3336  ; Intrinsic to read X component of thread ID3337  declare i32 @llvm.nvvm.read.ptx.sreg.tid.x() readnone nounwind3338 3339  define ptx_kernel void @kernel(ptr addrspace(1) %A,3340                                 ptr addrspace(1) %B,3341                                 ptr addrspace(1) %C) {3342  entry:3343    ; What is my ID?3344    %id = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x() readnone nounwind3345 3346    ; Compute pointers into A, B, and C3347    %ptrA = getelementptr float, ptr addrspace(1) %A, i32 %id3348    %ptrB = getelementptr float, ptr addrspace(1) %B, i32 %id3349    %ptrC = getelementptr float, ptr addrspace(1) %C, i32 %id3350 3351    ; Read A, B3352    %valA = load float, ptr addrspace(1) %ptrA, align 43353    %valB = load float, ptr addrspace(1) %ptrB, align 43354 3355    ; Compute C = A + B3356    %valC = fadd float %valA, %valB3357 3358    ; Store back to C3359    store float %valC, ptr addrspace(1) %ptrC, align 43360 3361    ret void3362  }3363 3364 3365We can use the LLVM ``llc`` tool to directly run the NVPTX code generator:3366 3367.. code-block:: text3368 3369  # llc -mcpu=sm_20 kernel.ll -o kernel.ptx3370 3371 3372.. note::3373 3374  If you want to generate 32-bit code, change ``p:64:64:64`` to ``p:32:32:32``3375  in the module data layout string and use ``nvptx-nvidia-cuda`` as the3376  target triple.3377 3378 3379The output we get from ``llc`` (as of LLVM 3.4):3380 3381.. code-block:: text3382 3383  //3384  // Generated by LLVM NVPTX Back-End3385  //3386 3387  .version 3.13388  .target sm_203389  .address_size 643390 3391    // .globl kernel3392                                          // @kernel3393  .visible .entry kernel(3394    .param .u64 kernel_param_0,3395    .param .u64 kernel_param_1,3396    .param .u64 kernel_param_23397  )3398  {3399    .reg .f32   %f<4>;3400    .reg .s32   %r<2>;3401    .reg .s64   %rl<8>;3402 3403  // %bb.0:                                // %entry3404    ld.param.u64    %rl1, [kernel_param_0];3405    mov.u32         %r1, %tid.x;3406    mul.wide.s32    %rl2, %r1, 4;3407    add.s64         %rl3, %rl1, %rl2;3408    ld.param.u64    %rl4, [kernel_param_1];3409    add.s64         %rl5, %rl4, %rl2;3410    ld.param.u64    %rl6, [kernel_param_2];3411    add.s64         %rl7, %rl6, %rl2;3412    ld.global.f32   %f1, [%rl3];3413    ld.global.f32   %f2, [%rl5];3414    add.f32         %f3, %f1, %f2;3415    st.global.f32   [%rl7], %f3;3416    ret;3417  }3418 3419 3420Dissecting the Kernel3421---------------------3422 3423Now let us dissect the LLVM IR that makes up this kernel.3424 3425Data Layout3426^^^^^^^^^^^3427 3428The data layout string determines the size in bits of common data types, their3429ABI alignment, and their storage size.  For NVPTX, you should use one of the3430following:3431 343232-bit PTX:3433 3434.. code-block:: llvm3435 3436  target datalayout = "e-p:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"3437 343864-bit PTX:3439 3440.. code-block:: llvm3441 3442  target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"3443 3444 3445Target Intrinsics3446^^^^^^^^^^^^^^^^^3447 3448In this example, we use the ``@llvm.nvvm.read.ptx.sreg.tid.x`` intrinsic to3449read the X component of the current thread's ID, which corresponds to a read3450of register ``%tid.x`` in PTX. The NVPTX back-end supports a large set of3451intrinsics.  A short list is shown below; please see3452``include/llvm/IR/IntrinsicsNVVM.td`` for the full list.3453 3454 3455================================================ ====================3456Intrinsic                                        CUDA Equivalent3457================================================ ====================3458``i32 @llvm.nvvm.read.ptx.sreg.tid.{x,y,z}``     threadIdx.{x,y,z}3459``i32 @llvm.nvvm.read.ptx.sreg.ctaid.{x,y,z}``   blockIdx.{x,y,z}3460``i32 @llvm.nvvm.read.ptx.sreg.ntid.{x,y,z}``    blockDim.{x,y,z}3461``i32 @llvm.nvvm.read.ptx.sreg.nctaid.{x,y,z}``  gridDim.{x,y,z}3462``void @llvm.nvvm.barrier0()``                   __syncthreads()3463================================================ ====================3464 3465 3466Address Spaces3467^^^^^^^^^^^^^^3468 3469You may have noticed that all of the pointer types in the LLVM IR example had3470an explicit address space specifier. What is address space 1? NVIDIA GPU3471devices (generally) have four types of memory:3472 3473- Global: Large, off-chip memory3474- Shared: Small, on-chip memory shared among all threads in a CTA3475- Local: Per-thread, private memory3476- Constant: Read-only memory shared across all threads3477 3478These different types of memory are represented in LLVM IR as address spaces.3479There is also a fifth address space used by the NVPTX code generator that3480corresponds to the "generic" address space.  This address space can represent3481addresses in any other address space (with a few exceptions).  This allows3482users to write IR functions that can load/store memory using the same3483instructions. Intrinsics are provided to convert pointers between the generic3484and non-generic address spaces.3485 3486See :ref:`address_spaces` and :ref:`nvptx_intrinsics` for more information.3487 3488 3489Running the Kernel3490------------------3491 3492Generating PTX from LLVM IR is all well and good, but how do we execute it on3493a real GPU device? The CUDA Driver API provides a convenient mechanism for3494loading and JIT compiling PTX to a native GPU device, and launching a kernel.3495The API is similar to OpenCL.  A simple example showing how to load and3496execute our vector addition code is shown below. Note that for brevity this3497code does not perform much error checking!3498 3499.. note::3500 3501  You can also use the ``ptxas`` tool provided by the CUDA Toolkit to offline3502  compile PTX to machine code (SASS) for a specific GPU architecture. Such3503  binaries can be loaded by the CUDA Driver API in the same way as PTX. This3504  can be useful for reducing startup time by precompiling the PTX kernels.3505 3506 3507.. code-block:: c++3508 3509  #include <iostream>3510  #include <fstream>3511  #include <cassert>3512  #include "cuda.h"3513 3514 3515  void checkCudaErrors(CUresult err) {3516    assert(err == CUDA_SUCCESS);3517  }3518 3519  /// main - Program entry point3520  int main(int argc, char **argv) {3521    CUdevice    device;3522    CUmodule    cudaModule;3523    CUcontext   context;3524    CUfunction  function;3525    CUlinkState linker;3526    int         devCount;3527 3528    // CUDA initialization3529    checkCudaErrors(cuInit(0));3530    checkCudaErrors(cuDeviceGetCount(&devCount));3531    checkCudaErrors(cuDeviceGet(&device, 0));3532 3533    char name[128];3534    checkCudaErrors(cuDeviceGetName(name, 128, device));3535    std::cout << "Using CUDA Device [0]: " << name << "\n";3536 3537    int devMajor, devMinor;3538    checkCudaErrors(cuDeviceComputeCapability(&devMajor, &devMinor, device));3539    std::cout << "Device Compute Capability: "3540              << devMajor << "." << devMinor << "\n";3541    if (devMajor < 2) {3542      std::cerr << "ERROR: Device 0 is not SM 2.0 or greater\n";3543      return 1;3544    }3545 3546    std::ifstream t("kernel.ptx");3547    if (!t.is_open()) {3548      std::cerr << "kernel.ptx not found\n";3549      return 1;3550    }3551    std::string str((std::istreambuf_iterator<char>(t)),3552                      std::istreambuf_iterator<char>());3553 3554    // Create driver context3555    checkCudaErrors(cuCtxCreate(&context, 0, device));3556 3557    // Create module for object3558    checkCudaErrors(cuModuleLoadDataEx(&cudaModule, str.c_str(), 0, 0, 0));3559 3560    // Get kernel function3561    checkCudaErrors(cuModuleGetFunction(&function, cudaModule, "kernel"));3562 3563    // Device data3564    CUdeviceptr devBufferA;3565    CUdeviceptr devBufferB;3566    CUdeviceptr devBufferC;3567 3568    checkCudaErrors(cuMemAlloc(&devBufferA, sizeof(float)*16));3569    checkCudaErrors(cuMemAlloc(&devBufferB, sizeof(float)*16));3570    checkCudaErrors(cuMemAlloc(&devBufferC, sizeof(float)*16));3571 3572    float* hostA = new float[16];3573    float* hostB = new float[16];3574    float* hostC = new float[16];3575 3576    // Populate input3577    for (unsigned i = 0; i != 16; ++i) {3578      hostA[i] = (float)i;3579      hostB[i] = (float)(2*i);3580      hostC[i] = 0.0f;3581    }3582 3583    checkCudaErrors(cuMemcpyHtoD(devBufferA, &hostA[0], sizeof(float)*16));3584    checkCudaErrors(cuMemcpyHtoD(devBufferB, &hostB[0], sizeof(float)*16));3585 3586 3587    unsigned blockSizeX = 16;3588    unsigned blockSizeY = 1;3589    unsigned blockSizeZ = 1;3590    unsigned gridSizeX  = 1;3591    unsigned gridSizeY  = 1;3592    unsigned gridSizeZ  = 1;3593 3594    // Kernel parameters3595    void *KernelParams[] = { &devBufferA, &devBufferB, &devBufferC };3596 3597    std::cout << "Launching kernel\n";3598 3599    // Kernel launch3600    checkCudaErrors(cuLaunchKernel(function, gridSizeX, gridSizeY, gridSizeZ,3601                                   blockSizeX, blockSizeY, blockSizeZ,3602                                   0, NULL, KernelParams, NULL));3603 3604    // Retrieve device data3605    checkCudaErrors(cuMemcpyDtoH(&hostC[0], devBufferC, sizeof(float)*16));3606 3607 3608    std::cout << "Results:\n";3609    for (unsigned i = 0; i != 16; ++i) {3610      std::cout << hostA[i] << " + " << hostB[i] << " = " << hostC[i] << "\n";3611    }3612 3613 3614    // Clean up after ourselves3615    delete [] hostA;3616    delete [] hostB;3617    delete [] hostC;3618 3619    // Clean-up3620    checkCudaErrors(cuMemFree(devBufferA));3621    checkCudaErrors(cuMemFree(devBufferB));3622    checkCudaErrors(cuMemFree(devBufferC));3623    checkCudaErrors(cuModuleUnload(cudaModule));3624    checkCudaErrors(cuCtxDestroy(context));3625 3626    return 0;3627  }3628 3629 3630You will need to link with the CUDA driver and specify the path to cuda.h.3631 3632.. code-block:: text3633 3634  # clang++ sample.cpp -o sample -O2 -g -I/usr/local/cuda-5.5/include -lcuda3635 3636We don't need to specify a path to ``libcuda.so`` since this is installed in a3637system location by the driver, not the CUDA toolkit.3638 3639If everything goes as planned, you should see the following output when3640running the compiled program:3641 3642.. code-block:: text3643 3644  Using CUDA Device [0]: GeForce GTX 6803645  Device Compute Capability: 3.03646  Launching kernel3647  Results:3648  0 + 0 = 03649  1 + 2 = 33650  2 + 4 = 63651  3 + 6 = 93652  4 + 8 = 123653  5 + 10 = 153654  6 + 12 = 183655  7 + 14 = 213656  8 + 16 = 243657  9 + 18 = 273658  10 + 20 = 303659  11 + 22 = 333660  12 + 24 = 363661  13 + 26 = 393662  14 + 28 = 423663  15 + 30 = 453664 3665.. note::3666 3667  You will likely see a different device identifier based on your hardware3668 3669 3670Tutorial: Linking with Libdevice3671================================3672 3673In this tutorial, we show a simple example of linking LLVM IR with the3674libdevice library. We will use the same kernel as the previous tutorial,3675except that we will compute ``C = pow(A, B)`` instead of ``C = A + B``.3676Libdevice provides an ``__nv_powf`` function that we will use.3677 3678.. code-block:: llvm3679 3680  target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"3681  target triple = "nvptx64-nvidia-cuda"3682 3683  ; Intrinsic to read X component of thread ID3684  declare i32 @llvm.nvvm.read.ptx.sreg.tid.x() readnone nounwind3685  ; libdevice function3686  declare float @__nv_powf(float, float)3687 3688  define ptx_kernel void @kernel(ptr addrspace(1) %A,3689                                 ptr addrspace(1) %B,3690                                 ptr addrspace(1) %C) {3691  entry:3692    ; What is my ID?3693    %id = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x() readnone nounwind3694 3695    ; Compute pointers into A, B, and C3696    %ptrA = getelementptr float, ptr addrspace(1) %A, i32 %id3697    %ptrB = getelementptr float, ptr addrspace(1) %B, i32 %id3698    %ptrC = getelementptr float, ptr addrspace(1) %C, i32 %id3699 3700    ; Read A, B3701    %valA = load float, ptr addrspace(1) %ptrA, align 43702    %valB = load float, ptr addrspace(1) %ptrB, align 43703 3704    ; Compute C = pow(A, B)3705    %valC = call float @__nv_powf(float %valA, float %valB)3706 3707    ; Store back to C3708    store float %valC, ptr addrspace(1) %ptrC, align 43709 3710    ret void3711  }3712 3713 3714To compile this kernel, we perform the following steps:3715 37161. Link with libdevice37172. Internalize all but the public kernel function37183. Run ``NVVMReflect`` and set ``__CUDA_FTZ`` to 037194. Optimize the linked module37205. Codegen the module3721 3722 3723These steps can be performed by the LLVM ``llvm-link``, ``opt``, and ``llc``3724tools. In a complete compiler, these steps can also be performed entirely3725programmatically by setting up an appropriate pass configuration (see3726:ref:`libdevice`).3727 3728.. code-block:: text3729 3730  # llvm-link t2.bc libdevice.compute_20.10.bc -o t2.linked.bc3731  # opt -internalize -internalize-public-api-list=kernel -nvvm-reflect-list=__CUDA_FTZ=0 -nvvm-reflect -O3 t2.linked.bc -o t2.opt.bc3732  # llc -mcpu=sm_20 t2.opt.bc -o t2.ptx3733 3734.. note::3735 3736  The ``-nvvm-reflect-list=_CUDA_FTZ=0`` is not strictly required, as any3737  undefined variables will default to zero. It is shown here for evaluation3738  purposes.3739 3740 3741This gives us the following PTX (excerpt):3742 3743.. code-block:: text3744 3745  //3746  // Generated by LLVM NVPTX Back-End3747  //3748 3749  .version 3.13750  .target sm_203751  .address_size 643752 3753    // .globl kernel3754                                          // @kernel3755  .visible .entry kernel(3756    .param .u64 kernel_param_0,3757    .param .u64 kernel_param_1,3758    .param .u64 kernel_param_23759  )3760  {3761    .reg .pred  %p<30>;3762    .reg .f32   %f<111>;3763    .reg .s32   %r<21>;3764    .reg .s64   %rl<8>;3765 3766  // %bb.0:                                // %entry3767    ld.param.u64  %rl2, [kernel_param_0];3768    mov.u32   %r3, %tid.x;3769    ld.param.u64  %rl3, [kernel_param_1];3770    mul.wide.s32  %rl4, %r3, 4;3771    add.s64   %rl5, %rl2, %rl4;3772    ld.param.u64  %rl6, [kernel_param_2];3773    add.s64   %rl7, %rl3, %rl4;3774    add.s64   %rl1, %rl6, %rl4;3775    ld.global.f32   %f1, [%rl5];3776    ld.global.f32   %f2, [%rl7];3777    setp.eq.f32 %p1, %f1, 0f3F800000;3778    setp.eq.f32 %p2, %f2, 0f00000000;3779    or.pred   %p3, %p1, %p2;3780    @%p3 bra  BB0_1;3781    bra.uni   BB0_2;3782  BB0_1:3783    mov.f32   %f110, 0f3F800000;3784    st.global.f32   [%rl1], %f110;3785    ret;3786  BB0_2:                                  // %__nv_isnanf.exit.i3787    abs.f32   %f4, %f1;3788    setp.gtu.f32  %p4, %f4, 0f7F800000;3789    @%p4 bra  BB0_4;3790  // %bb.3:                                // %__nv_isnanf.exit5.i3791    abs.f32   %f5, %f2;3792    setp.le.f32 %p5, %f5, 0f7F800000;3793    @%p5 bra  BB0_5;3794  BB0_4:                                  // %.critedge1.i3795    add.f32   %f110, %f1, %f2;3796    st.global.f32   [%rl1], %f110;3797    ret;3798  BB0_5:                                  // %__nv_isinff.exit.i3799 3800    ...3801 3802  BB0_26:                                 // %__nv_truncf.exit.i.i.i.i.i3803    mul.f32   %f90, %f107, 0f3FB8AA3B;3804    cvt.rzi.f32.f32 %f91, %f90;3805    mov.f32   %f92, 0fBF317200;3806    fma.rn.f32  %f93, %f91, %f92, %f107;3807    mov.f32   %f94, 0fB5BFBE8E;3808    fma.rn.f32  %f95, %f91, %f94, %f93;3809    mul.f32   %f89, %f95, 0f3FB8AA3B;3810    // inline asm3811    ex2.approx.ftz.f32 %f88,%f89;3812    // inline asm3813    add.f32   %f96, %f91, 0f00000000;3814    ex2.approx.f32  %f97, %f96;3815    mul.f32   %f98, %f88, %f97;3816    setp.lt.f32 %p15, %f107, 0fC2D20000;3817    selp.f32  %f99, 0f00000000, %f98, %p15;3818    setp.gt.f32 %p16, %f107, 0f42D20000;3819    selp.f32  %f110, 0f7F800000, %f99, %p16;3820    setp.eq.f32 %p17, %f110, 0f7F800000;3821    @%p17 bra   BB0_28;3822  // %bb.27:3823    fma.rn.f32  %f110, %f110, %f108, %f110;3824  BB0_28:                                 // %__internal_accurate_powf.exit.i3825    setp.lt.f32 %p18, %f1, 0f00000000;3826    setp.eq.f32 %p19, %f3, 0f3F800000;3827    and.pred    %p20, %p18, %p19;3828    @!%p20 bra  BB0_30;3829    bra.uni   BB0_29;3830  BB0_29:3831    mov.b32    %r9, %f110;3832    xor.b32   %r10, %r9, -2147483648;3833    mov.b32    %f110, %r10;3834  BB0_30:                                 // %__nv_powf.exit3835    st.global.f32   [%rl1], %f110;3836    ret;3837  }3838