Title: Contents — PTX ISA 9.1 documentation

URL Source: https://docs.nvidia.com/cuda/parallel-thread-execution/contents.html

Published Time: Thu, 04 Dec 2025 20:19:53 GMT

Markdown Content:
*   [](https://docs.nvidia.com/cuda/index.html) »
*   Contents

*   v9.1 |[PDF](https://docs.nvidia.com/cuda/pdf/ptx_isa_9.1.pdf)|[Archive](https://developer.nvidia.com/cuda-toolkit-archive.md)

* * *

Contents[](https://docs.nvidia.com/cuda/parallel-thread-execution/contents.html#contents "Permalink to this headline")
=======================================================================================================================

*   [1. Introduction](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md)
    *   [1.1. Scalable Data-Parallel Computing using GPUs](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#scalable-data-parallel-computing-using-gpus)
    *   [1.2. Goals of PTX](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#goals-of-ptx)
    *   [1.3. PTX ISA Version 9.1](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#ptx-isa-version-9-1)
    *   [1.4. Document Structure](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#document-structure)

*   [2. Programming Model](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#programming-model)
    *   [2.1. A Highly Multithreaded Coprocessor](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#highly-multithreaded-coprocessor)
    *   [2.2. Thread Hierarchy](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#thread-hierarchy)
        *   [2.2.1. Cooperative Thread Arrays](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#cooperative-thread-arrays)
        *   [2.2.2. Cluster of Cooperative Thread Arrays](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#cluster-of-cooperative-thread-arrays)
        *   [2.2.3. Grid of Clusters](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#grid-of-clusters)

    *   [2.3. Memory Hierarchy](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#memory-hierarchy)

*   [3. PTX Machine Model](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#ptx-machine-model)
    *   [3.1. A Set of SIMT Multiprocessors](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#set-of-simt-multiprocessors)
    *   [3.2. Independent Thread Scheduling](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#independent-thread-scheduling)
    *   [3.3. On-chip Shared Memory](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#on-chip-shared-memory)

*   [4. Syntax](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#syntax)
    *   [4.1. Source Format](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#source-format)
    *   [4.2. Comments](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#comments)
    *   [4.3. Statements](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#statements)
        *   [4.3.1. Directive Statements](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#directive-statements)
        *   [4.3.2. Instruction Statements](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#instruction-statements)

    *   [4.4. Identifiers](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#identifiers)
    *   [4.5. Constants](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#constants)
        *   [4.5.1. Integer Constants](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#integer-constants)
        *   [4.5.2. Floating-Point Constants](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#floating-point-constants)
        *   [4.5.3. Predicate Constants](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#predicate-constants)
        *   [4.5.4. Constant Expressions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#constant-expressions)
        *   [4.5.5. Integer Constant Expression Evaluation](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#integer-constant-expression-evaluation)
        *   [4.5.6. Summary of Constant Expression Evaluation Rules](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#summary-of-constant-expression-evaluation-rules)

*   [5. State Spaces, Types, and Variables](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#state-spaces-types-and-variables)
    *   [5.1. State Spaces](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#state-spaces)
        *   [5.1.1. Register State Space](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#register-state-space)
        *   [5.1.2. Special Register State Space](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#special-register-state-space)
        *   [5.1.3. Constant State Space](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#constant-state-space)
            *   [5.1.3.1. Banked Constant State Space (deprecated)](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#banked-constant-state-space-deprecated)

        *   [5.1.4. Global State Space](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#global-state-space)
        *   [5.1.5. Local State Space](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#local-state-space)
        *   [5.1.6. Parameter State Space](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parameter-state-space)
            *   [5.1.6.1. Kernel Function Parameters](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#kernel-function-parameters)
            *   [5.1.6.2. Kernel Function Parameter Attributes](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#kernel-function-parameter-attributes)
            *   [5.1.6.3. Kernel Parameter Attribute: `.ptr`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#kernel-parameter-attribute-ptr)
            *   [5.1.6.4. Device Function Parameters](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#device-function-parameters)

        *   [5.1.7. Shared State Space](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#shared-state-space)
        *   [5.1.8. Texture State Space (deprecated)](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#texture-state-space-deprecated)

    *   [5.2. Types](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#types)
        *   [5.2.1. Fundamental Types](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#fundamental-types)
        *   [5.2.2. Restricted Use of Sub-Word Sizes](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#restricted-use-of-sub-word-sizes)
        *   [5.2.3. Alternate Floating-Point Data Formats](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#alternate-floating-point-data-formats)
        *   [5.2.4. Fixed-point Data format](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#fixed-point-data-formats)
        *   [5.2.5. Packed Data Types](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#packed-data-types)
            *   [5.2.5.1. Packed Floating Point Data Types](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#packed-floating-point-data-types)
            *   [5.2.5.2. Packed Integer Data Types](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#packed-integer-data-types)
            *   [5.2.5.3. Packed Fixed-Point Data Types](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#packed-fixed-point-data-types)

    *   [5.3. Texture Sampler and Surface Types](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#texture-sampler-and-surface-types)
        *   [5.3.1. Texture and Surface Properties](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#texture-surface-properties)
        *   [5.3.2. Sampler Properties](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#sampler-properties)
        *   [5.3.3. Channel Data Type and Channel Order Fields](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#channel-data-type-and-channel-order-fields)

    *   [5.4. Variables](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#variables)
        *   [5.4.1. Variable Declarations](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#variable-declarations)
        *   [5.4.2. Vectors](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#vectors)
        *   [5.4.3. Array Declarations](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#array-declarations)
        *   [5.4.4. Initializers](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#initializers)
        *   [5.4.5. Alignment](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#alignment)
        *   [5.4.6. Parameterized Variable Names](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parameterized-variable-names)
        *   [5.4.7. Variable Attributes](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#variable-attributes)
        *   [5.4.8. Variable and Function Attribute Directive: `.attribute`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#variable-and-function-attribute-directive-attribute)

    *   [5.5. Tensors](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tensors)
        *   [5.5.1. Tensor Dimension, size and format](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tensor-dimension-size-format)
            *   [5.5.1.1. Sub-byte Types](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tensor-dimension-size-format-sub-bytes)
                *   [5.5.1.1.1. Padding and alignment of the sub-byte types](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tensor-dimension-size-format-sub-bytes-padding-align)

        *   [5.5.2. Tensor Access Modes](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tensor-access-modes)
        *   [5.5.3. Tiled Mode](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tensor-tiled-mode)
            *   [5.5.3.1. Bounding Box](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tensor-tiled-mode-bounding-box)
            *   [5.5.3.2. Traversal-Stride](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tensor-tiled-mode-traversal-stride)
            *   [5.5.3.3. Out of Boundary Access](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tensor-tiled-mode-oob-access)
            *   [5.5.3.4. `.tile::scatter4` and `.tile::gather4` modes](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tensor-tiled-scatter4-gather4-modes)
                *   [5.5.3.4.1. Bounding Box](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tensor-tiled-scatter4-gather4-modes-bounding-box)

        *   [5.5.4. `im2col` mode](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tensor-im2col-mode)
            *   [5.5.4.1. Bounding Box](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tensor-im2col-mode-bounding-box)
            *   [5.5.4.2. Traversal Stride](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tensor-im2col-mode-traversal-stride)
            *   [5.5.4.3. Out of Boundary Access](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tensor-im2col-mode-oob-access)

        *   [5.5.5. `im2col::w` and `im2col::w::128` modes](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tensor-im2col-w-w128-modes)
            *   [5.5.5.1. Bounding Box](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tensor-im2col-w-w128-modes-bounding-box)
            *   [5.5.5.2. Traversal Stride](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tensor-im2col-w-w128-modes-traversal-stride)
            *   [5.5.5.3. `wHalo`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tensor-im2col-w-w128-modes-whalo)
            *   [5.5.5.4. `wOffset`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tensor-im2col-w-w128-modes-woffset)

        *   [5.5.6. Interleave layout](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tensor-interleaved-layout)
        *   [5.5.7. Swizzling Modes](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tensor-swizzling-modes)
        *   [5.5.8. Tensor-map](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tensor-tensormap)

*   [6. Instruction Operands](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#instruction-operands)
    *   [6.1. Operand Type Information](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#operand-type-information)
    *   [6.2. Source Operands](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#source-operands)
    *   [6.3. Destination Operands](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#destination-operands)
    *   [6.4. Using Addresses, Arrays, and Vectors](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#using-addresses-arrays-and-vectors)
        *   [6.4.1. Addresses as Operands](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#addresses-as-operands)
            *   [6.4.1.1. Generic Addressing](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#generic-addressing)

        *   [6.4.2. Arrays as Operands](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#arrays-as-operands)
        *   [6.4.3. Vectors as Operands](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#vectors-as-operands)
        *   [6.4.4. Labels and Function Names as Operands](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#labels-and-function-names-as-operands)

    *   [6.5. Type Conversion](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#type-conversion)
        *   [6.5.1. Scalar Conversions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#scalar-conversions)
        *   [6.5.2. Rounding Modifiers](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#rounding-modifiers)

    *   [6.6. Operand Costs](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#operand-costs)

*   [7. Abstracting the ABI](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#abstracting-abi)
    *   [7.1. Function Declarations and Definitions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#function-declarations-and-definitions)
        *   [7.1.1. Changes from PTX ISA Version 1.x](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-from-ptx-isa-version-1-x)

    *   [7.2. Variadic Functions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#variadic-functions)
    *   [7.3. Alloca](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#alloca)

*   [8. Memory Consistency Model](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#memory-consistency-model)
    *   [8.1. Scope and applicability of the model](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#scope-and-applicability)
        *   [8.1.1. Limitations on atomicity at system scope](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#limitations-system-scope-atomicity)

    *   [8.2. Memory operations](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#memory-operations)
        *   [8.2.1. Overlap](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#overlap)
        *   [8.2.2. Aliases](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#aliases)
        *   [8.2.3. Multimem Addresses](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#multimem-addresses)
        *   [8.2.4. Memory Operations on Vector Data Types](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#memory-operations-on-vector-data-types)
        *   [8.2.5. Memory Operations on Packed Data Types](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#memory-operations-on-packed-data-types)
        *   [8.2.6. Initialization](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#initialization)

    *   [8.3. State spaces](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#memory-consistency-state-spaces)
    *   [8.4. Operation types](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#operation-types)
        *   [8.4.1. mmio Operation](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#mmio-operation)
        *   [8.4.2. volatile Operation](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#volatile-operation)

    *   [8.5. Scope](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#scope)
    *   [8.6. Proxies](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#proxies)
    *   [8.7. Morally strong operations](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#morally-strong-operations)
        *   [8.7.1. Conflict and Data-races](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#conflict-and-data-races)
        *   [8.7.2. Limitations on Mixed-size Data-races](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#mixed-size-limitations)

    *   [8.8. Release and Acquire Patterns](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#release-acquire-patterns)
    *   [8.9. Ordering of memory operations](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#ordering-memory-operations)
        *   [8.9.1. Program Order](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#program-order)
            *   [8.9.1.1. Asynchronous Operations](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#program-order-async-operations)

        *   [8.9.2. Observation Order](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#observation-order)
        *   [8.9.3. Fence-SC Order](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#fence-sc-order)
        *   [8.9.4. Memory synchronization](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#memory-synchronization)
        *   [8.9.5. Causality Order](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#causality-order)
        *   [8.9.6. Coherence Order](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#coherence-order)
        *   [8.9.7. Communication Order](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#communication-order)

    *   [8.10. Axioms](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#axioms)
        *   [8.10.1. Coherence](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#coherence-axiom)
        *   [8.10.2. Fence-SC](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#fence-sc-axiom)
        *   [8.10.3. Atomicity](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#atomicity-axiom)
        *   [8.10.4. No Thin Air](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#no-thin-air-axiom)
        *   [8.10.5. Sequential Consistency Per Location](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#sc-per-loc-axiom)
        *   [8.10.6. Causality](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#causality-axiom)

    *   [8.11. Special Cases](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#special-cases)
        *   [8.11.1. Reductions do not form Acquire Patterns](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#red-read)

*   [9. Instruction Set](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#instruction-set)
    *   [9.1. Format and Semantics of Instruction Descriptions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#format-and-semantics-of-instruction-descriptions)
    *   [9.2. PTX Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#ptx-instructions)
    *   [9.3. Predicated Execution](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#predicated-execution)
        *   [9.3.1. Comparisons](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#comparisons)
            *   [9.3.1.1. Integer and Bit-Size Comparisons](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#integer-and-bit-size-comparisons)
            *   [9.3.1.2. Floating Point Comparisons](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#floating-point-comparisons)

        *   [9.3.2. Manipulating Predicates](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#manipulating-predicates)

    *   [9.4. Type Information for Instructions and Operands](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#type-information-for-instructions-and-operands)
        *   [9.4.1. Operand Size Exceeding Instruction-Type Size](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#operand-size-exceeding-instruction-type-size)

    *   [9.5. Divergence of Threads in Control Constructs](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#divergence-of-threads-in-control-constructs)
    *   [9.6. Semantics](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#semantics)
        *   [9.6.1. Machine-Specific Semantics of 16-bit Code](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#machine-specific-semantics-of-16-bit-code)

    *   [9.7. Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#instructions)
        *   [9.7.1. Integer Arithmetic Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#integer-arithmetic-instructions)
            *   [9.7.1.1. Integer Arithmetic Instructions: `add`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-add)
            *   [9.7.1.2. Integer Arithmetic Instructions: `sub`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-sub)
            *   [9.7.1.3. Integer Arithmetic Instructions: `mul`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-mul)
            *   [9.7.1.4. Integer Arithmetic Instructions: `mad`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-mad)
            *   [9.7.1.5. Integer Arithmetic Instructions: `mul24`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-mul24)
            *   [9.7.1.6. Integer Arithmetic Instructions: `mad24`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-mad24)
            *   [9.7.1.7. Integer Arithmetic Instructions: `sad`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-sad)
            *   [9.7.1.8. Integer Arithmetic Instructions: `div`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-div)
            *   [9.7.1.9. Integer Arithmetic Instructions: `rem`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-rem)
            *   [9.7.1.10. Integer Arithmetic Instructions: `abs`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-abs)
            *   [9.7.1.11. Integer Arithmetic Instructions: `neg`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-neg)
            *   [9.7.1.12. Integer Arithmetic Instructions: `min`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-min)
            *   [9.7.1.13. Integer Arithmetic Instructions: `max`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-max)
            *   [9.7.1.14. Integer Arithmetic Instructions: `popc`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-popc)
            *   [9.7.1.15. Integer Arithmetic Instructions: `clz`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-clz)
            *   [9.7.1.16. Integer Arithmetic Instructions: `bfind`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-bfind)
            *   [9.7.1.17. Integer Arithmetic Instructions: `fns`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-fns)
            *   [9.7.1.18. Integer Arithmetic Instructions: `brev`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-brev)
            *   [9.7.1.19. Integer Arithmetic Instructions: `bfe`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-bfe)
            *   [9.7.1.20. Integer Arithmetic Instructions: `bfi`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-bfi)
            *   [9.7.1.21. Integer Arithmetic Instructions: `szext`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-szext)
            *   [9.7.1.22. Integer Arithmetic Instructions: `bmsk`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-bmsk)
            *   [9.7.1.23. Integer Arithmetic Instructions: `dp4a`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-dp4a)
            *   [9.7.1.24. Integer Arithmetic Instructions: `dp2a`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-dp2a)

        *   [9.7.2. Extended-Precision Integer Arithmetic Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#extended-precision-integer-arithmetic-instructions)
            *   [9.7.2.1. Extended-Precision Arithmetic Instructions: `add.cc`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#extended-precision-arithmetic-instructions-add-cc)
            *   [9.7.2.2. Extended-Precision Arithmetic Instructions: `addc`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#extended-precision-arithmetic-instructions-addc)
            *   [9.7.2.3. Extended-Precision Arithmetic Instructions: `sub.cc`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#extended-precision-arithmetic-instructions-sub-cc)
            *   [9.7.2.4. Extended-Precision Arithmetic Instructions: `subc`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#extended-precision-arithmetic-instructions-subc)
            *   [9.7.2.5. Extended-Precision Arithmetic Instructions: `mad.cc`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#extended-precision-arithmetic-instructions-mad-cc)
            *   [9.7.2.6. Extended-Precision Arithmetic Instructions: `madc`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#extended-precision-arithmetic-instructions-madc)

        *   [9.7.3. Floating-Point Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#floating-point-instructions)
            *   [9.7.3.1. Floating Point Instructions: `testp`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#floating-point-instructions-testp)
            *   [9.7.3.2. Floating Point Instructions: `copysign`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#floating-point-instructions-copysign)
            *   [9.7.3.3. Floating Point Instructions: `add`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#floating-point-instructions-add)
            *   [9.7.3.4. Floating Point Instructions: `sub`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#floating-point-instructions-sub)
            *   [9.7.3.5. Floating Point Instructions: `mul`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#floating-point-instructions-mul)
            *   [9.7.3.6. Floating Point Instructions: `fma`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#floating-point-instructions-fma)
            *   [9.7.3.7. Floating Point Instructions: `mad`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#floating-point-instructions-mad)
            *   [9.7.3.8. Floating Point Instructions: `div`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#floating-point-instructions-div)
            *   [9.7.3.9. Floating Point Instructions: `abs`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#floating-point-instructions-abs)
            *   [9.7.3.10. Floating Point Instructions: `neg`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#floating-point-instructions-neg)
            *   [9.7.3.11. Floating Point Instructions: `min`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#floating-point-instructions-min)
            *   [9.7.3.12. Floating Point Instructions: `max`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#floating-point-instructions-max)
            *   [9.7.3.13. Floating Point Instructions: `rcp`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#floating-point-instructions-rcp)
            *   [9.7.3.14. Floating Point Instructions: `rcp.approx.ftz.f64`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#floating-point-instructions-rcp-approx-ftz-f64)
            *   [9.7.3.15. Floating Point Instructions: `sqrt`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#floating-point-instructions-sqrt)
            *   [9.7.3.16. Floating Point Instructions: `rsqrt`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#floating-point-instructions-rsqrt)
            *   [9.7.3.17. Floating Point Instructions: `rsqrt.approx.ftz.f64`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#floating-point-instructions-rsqrt-approx-ftz-f64)
            *   [9.7.3.18. Floating Point Instructions: `sin`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#floating-point-instructions-sin)
            *   [9.7.3.19. Floating Point Instructions: `cos`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#floating-point-instructions-cos)
            *   [9.7.3.20. Floating Point Instructions: `lg2`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#floating-point-instructions-lg2)
            *   [9.7.3.21. Floating Point Instructions: `ex2`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#floating-point-instructions-ex2)
            *   [9.7.3.22. Floating Point Instructions: `tanh`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#floating-point-instructions-tanh)

        *   [9.7.4. Half Precision Floating-Point Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#half-precision-floating-point-instructions)
            *   [9.7.4.1. Half Precision Floating Point Instructions: `add`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#half-precision-floating-point-instructions-add)
            *   [9.7.4.2. Half Precision Floating Point Instructions: `sub`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#half-precision-floating-point-instructions-sub)
            *   [9.7.4.3. Half Precision Floating Point Instructions: `mul`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#half-precision-floating-point-instructions-mul)
            *   [9.7.4.4. Half Precision Floating Point Instructions: `fma`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#half-precision-floating-point-instructions-fma)
            *   [9.7.4.5. Half Precision Floating Point Instructions: `neg`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#half-precision-floating-point-instructions-neg)
            *   [9.7.4.6. Half Precision Floating Point Instructions: `abs`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#half-precision-floating-point-instructions-abs)
            *   [9.7.4.7. Half Precision Floating Point Instructions: `min`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#half-precision-floating-point-instructions-min)
            *   [9.7.4.8. Half Precision Floating Point Instructions: `max`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#half-precision-floating-point-instructions-max)
            *   [9.7.4.9. Half Precision Floating Point Instructions: `tanh`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#half-precision-floating-point-instructions-tanh)
            *   [9.7.4.10. Half Precision Floating Point Instructions: `ex2`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#half-precision-floating-point-instructions-ex2)

        *   [9.7.5. Mixed Precision Floating-Point Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#mixed-precision-floating-point-instructions)
            *   [9.7.5.1. Mixed Precision Floating Point Instructions: `add`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#mixed-precision-floating-point-instructions-add)
            *   [9.7.5.2. Mixed Precision Floating Point Instructions: `sub`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#mixed-precision-floating-point-instructions-sub)
            *   [9.7.5.3. Mixed Precision Floating Point Instructions: `fma`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#mixed-precision-floating-point-instructions-fma)

        *   [9.7.6. Comparison and Selection Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#comparison-and-selection-instructions)
            *   [9.7.6.1. Comparison and Selection Instructions: `set`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#comparison-and-selection-instructions-set)
            *   [9.7.6.2. Comparison and Selection Instructions: `setp`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#comparison-and-selection-instructions-setp)
            *   [9.7.6.3. Comparison and Selection Instructions: `selp`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#comparison-and-selection-instructions-selp)
            *   [9.7.6.4. Comparison and Selection Instructions: `slct`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#comparison-and-selection-instructions-slct)

        *   [9.7.7. Half Precision Comparison Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#half-precision-comparison-instructions)
            *   [9.7.7.1. Half Precision Comparison Instructions: `set`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#half-precision-comparison-instructions-set)
            *   [9.7.7.2. Half Precision Comparison Instructions: `setp`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#half-precision-comparison-instructions-setp)

        *   [9.7.8. Logic and Shift Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#logic-and-shift-instructions)
            *   [9.7.8.1. Logic and Shift Instructions: `and`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#logic-and-shift-instructions-and)
            *   [9.7.8.2. Logic and Shift Instructions: `or`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#logic-and-shift-instructions-or)
            *   [9.7.8.3. Logic and Shift Instructions: `xor`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#logic-and-shift-instructions-xor)
            *   [9.7.8.4. Logic and Shift Instructions: `not`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#logic-and-shift-instructions-not)
            *   [9.7.8.5. Logic and Shift Instructions: `cnot`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#logic-and-shift-instructions-cnot)
            *   [9.7.8.6. Logic and Shift Instructions: `lop3`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#logic-and-shift-instructions-lop3)
            *   [9.7.8.7. Logic and Shift Instructions: `shf`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#logic-and-shift-instructions-shf)
            *   [9.7.8.8. Logic and Shift Instructions: `shl`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#logic-and-shift-instructions-shl)
            *   [9.7.8.9. Logic and Shift Instructions: `shr`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#logic-and-shift-instructions-shr)

        *   [9.7.9. Data Movement and Conversion Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions)
            *   [9.7.9.1. Cache Operators](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#cache-operators)
            *   [9.7.9.2. Cache Eviction Priority Hints](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#cache-eviction-priority-hints)
            *   [9.7.9.3. Data Movement and Conversion Instructions: `mov`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-mov)
            *   [9.7.9.4. Data Movement and Conversion Instructions: `mov`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-mov-2)
            *   [9.7.9.5. Data Movement and Conversion Instructions: `shfl` (deprecated)](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-shfl)
            *   [9.7.9.6. Data Movement and Conversion Instructions: `shfl.sync`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-shfl-sync)
            *   [9.7.9.7. Data Movement and Conversion Instructions: `prmt`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-prmt)
            *   [9.7.9.8. Data Movement and Conversion Instructions: `ld`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-ld)
            *   [9.7.9.9. Data Movement and Conversion Instructions: `ld.global.nc`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-ld-global-nc)
            *   [9.7.9.10. Data Movement and Conversion Instructions: `ldu`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-ldu)
            *   [9.7.9.11. Data Movement and Conversion Instructions: `st`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-st)
            *   [9.7.9.12. Data Movement and Conversion Instructions: `st.async`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-st-async)
            *   [9.7.9.13. Data Movement and Conversion Instructions: `st.bulk`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-st-bulk)
            *   [9.7.9.14. Data Movement and Conversion Instructions: `multimem.ld_reduce`, `multimem.st`, `multimem.red`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-multimem)
            *   [9.7.9.15. Data Movement and Conversion Instructions: `prefetch`, `prefetchu`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-prefetch-prefetchu)
            *   [9.7.9.16. Data Movement and Conversion Instructions: `applypriority`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-applypriority)
            *   [9.7.9.17. Data Movement and Conversion Instructions: `discard`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-discard)
            *   [9.7.9.18. Data Movement and Conversion Instructions: `createpolicy`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-createpolicy)
            *   [9.7.9.19. Data Movement and Conversion Instructions: `isspacep`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-isspacep)
            *   [9.7.9.20. Data Movement and Conversion Instructions: `cvta`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cvta)
            *   [9.7.9.21. Data Movement and Conversion Instructions: `cvt`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cvt)
            *   [9.7.9.22. Data Movement and Conversion Instructions: `cvt.pack`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cvt-pack)
            *   [9.7.9.23. Data Movement and Conversion Instructions: `mapa`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-mapa)
            *   [9.7.9.24. Data Movement and Conversion Instructions: `getctarank`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-getctarank)
            *   [9.7.9.25. Data Movement and Conversion Instructions: Asynchronous copy](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-asynchronous-copy)
                *   [9.7.9.25.1. Completion Mechanisms for Asynchronous Copy Operations](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-asynchronous-copy-completion-mechanisms)
                    *   [9.7.9.25.1.1. Async-group mechanism](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-asynchronous-copy-completion-mechanisms-async-group)
                    *   [9.7.9.25.1.2. Mbarrier-based mechanism](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-asynchronous-copy-completion-mechanisms-mbarrier)

                *   [9.7.9.25.2. Async Proxy](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#async-proxy)
                *   [9.7.9.25.3. Data Movement and Conversion Instructions: Non-bulk copy](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-non-bulk-copy)
                    *   [9.7.9.25.3.1. Data Movement and Conversion Instructions: `cp.async`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async)
                    *   [9.7.9.25.3.2. Data Movement and Conversion Instructions: `cp.async.commit_group`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-commit-group)
                    *   [9.7.9.25.3.3. Data Movement and Conversion Instructions: `cp.async.wait_group` / `cp.async.wait_all`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-wait-group)

                *   [9.7.9.25.4. Data Movement and Conversion Instructions: Bulk copy](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-bulk-copy)
                    *   [9.7.9.25.4.1. Data Movement and Conversion Instructions: `cp.async.bulk`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk)
                    *   [9.7.9.25.4.2. Data Movement and Conversion Instructions: `cp.reduce.async.bulk`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-reduce-async-bulk)
                    *   [9.7.9.25.4.3. Data Movement and Conversion Instructions: `cp.async.bulk.prefetch`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk-prefetch)

            *   [9.7.9.26. Data Movement and Conversion Instructions: `multimem.cp.async.bulk`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-multimem-cp-async-bulk)
            *   [9.7.9.27. Data Movement and Conversion Instructions: `multimem.cp.reduce.async.bulk`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-multimem-cp-reduce-async-bulk)
                *   [9.7.9.27.1. Data Movement and Conversion Instructions: Tensor copy](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-tensor-copy)
                    *   [9.7.9.27.1.1. Restriction on Tensor Copy instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-tensor-copy-restrictions)
                    *   [9.7.9.27.1.2. Data Movement and Conversion Instructions: `cp.async.bulk.tensor`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk-tensor)
                    *   [9.7.9.27.1.3. Data Movement and Conversion Instructions: `cp.reduce.async.bulk.tensor`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-reduce-async-bulk-tensor)
                    *   [9.7.9.27.1.4. Data Movement and Conversion Instructions: `cp.async.bulk.prefetch.tensor`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk-prefetch-tensor)

                *   [9.7.9.27.2. Data Movement and Conversion Instructions: Bulk and Tensor copy completion instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-bulk-tensor-copy-completion)
                    *   [9.7.9.27.2.1. Data Movement and Conversion Instructions: `cp.async.bulk.commit_group`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk-commit-group)
                    *   [9.7.9.27.2.2. Data Movement and Conversion Instructions: `cp.async.bulk.wait_group`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk-wait-group)

            *   [9.7.9.28. Data Movement and Conversion Instructions: `tensormap.replace`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-tensormap-replace)

        *   [9.7.10. Texture Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#texture-instructions)
            *   [9.7.10.1. Texturing Modes](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#texturing-modes)
            *   [9.7.10.2. Mipmaps](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#mipmaps)
            *   [9.7.10.3. Texture Instructions: `tex`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#texture-instructions-tex)
            *   [9.7.10.4. Texture Instructions: `tld4`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#texture-instructions-tld4)
            *   [9.7.10.5. Texture Instructions: `txq`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#texture-instructions-txq)
            *   [9.7.10.6. Texture Instructions: `istypep`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#texture-instructions-istypep)

        *   [9.7.11. Surface Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#surface-instructions)
            *   [9.7.11.1. Surface Instructions: `suld`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#surface-instructions-suld)
            *   [9.7.11.2. Surface Instructions: `sust`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#surface-instructions-sust)
            *   [9.7.11.3. Surface Instructions: `sured`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#surface-instructions-sured)
            *   [9.7.11.4. Surface Instructions: `suq`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#surface-instructions-suq)

        *   [9.7.12. Control Flow Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#control-flow-instructions)
            *   [9.7.12.1. Control Flow Instructions: `{}`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#control-flow-instructions-curly-braces)
            *   [9.7.12.2. Control Flow Instructions: `@`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#control-flow-instructions-at)
            *   [9.7.12.3. Control Flow Instructions: `bra`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#control-flow-instructions-bra)
            *   [9.7.12.4. Control Flow Instructions: `brx.idx`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#control-flow-instructions-brx-idx)
            *   [9.7.12.5. Control Flow Instructions: `call`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#control-flow-instructions-call)
            *   [9.7.12.6. Control Flow Instructions: `ret`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#control-flow-instructions-ret)
            *   [9.7.12.7. Control Flow Instructions: `exit`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#control-flow-instructions-exit)

        *   [9.7.13. Parallel Synchronization and Communication Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions)
            *   [9.7.13.1. Parallel Synchronization and Communication Instructions: `bar`, `barrier`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-bar)
            *   [9.7.13.2. Parallel Synchronization and Communication Instructions: `bar.warp.sync`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-bar-warp-sync)
            *   [9.7.13.3. Parallel Synchronization and Communication Instructions: `barrier.cluster`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-barrier-cluster)
            *   [9.7.13.4. Parallel Synchronization and Communication Instructions: `membar` / `fence`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-membar)
            *   [9.7.13.5. Parallel Synchronization and Communication Instructions: `atom`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-atom)
            *   [9.7.13.6. Parallel Synchronization and Communication Instructions: `red`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-red)
            *   [9.7.13.7. Parallel Synchronization and Communication Instructions: `red.async`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-red-async)
            *   [9.7.13.8. Parallel Synchronization and Communication Instructions: `vote` (deprecated)](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-vote)
            *   [9.7.13.9. Parallel Synchronization and Communication Instructions: `vote.sync`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-vote-sync)
            *   [9.7.13.10. Parallel Synchronization and Communication Instructions: `match.sync`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-match-sync)
            *   [9.7.13.11. Parallel Synchronization and Communication Instructions: `activemask`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-activemask)
            *   [9.7.13.12. Parallel Synchronization and Communication Instructions: `redux.sync`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-redux-sync)
            *   [9.7.13.13. Parallel Synchronization and Communication Instructions: `griddepcontrol`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-griddepcontrol)
            *   [9.7.13.14. Parallel Synchronization and Communication Instructions: `elect.sync`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-elect-sync)
            *   [9.7.13.15. Parallel Synchronization and Communication Instructions: `mbarrier`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-mbarrier)
                *   [9.7.13.15.1. Size and alignment of mbarrier object](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-mbarrier-size-alignment)
                *   [9.7.13.15.2. Contents of the mbarrier object](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-mbarrier-contents)
                *   [9.7.13.15.3. Lifecycle of the mbarrier object](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-mbarrier-lifecycle)
                *   [9.7.13.15.4. Phase of the mbarrier object](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-mbarrier-phase)
                *   [9.7.13.15.5. Tracking asynchronous operations by the mbarrier object](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-mbarrier-tracking-async-operations)
                    *   [9.7.13.15.5.1. expect-tx operation](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-mbarrier-expect-tx-operation)
                    *   [9.7.13.15.5.2. complete-tx operation](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-mbarrier-complete-tx-operation)

                *   [9.7.13.15.6. Phase Completion of the mbarrier object](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-mbarrier-phase-completion)
                *   [9.7.13.15.7. Arrive-on operation on mbarrier object](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-mbarrier-arrive-on)
                *   [9.7.13.15.8. mbarrier support with shared memory](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-mbarrier-smem)
                *   [9.7.13.15.9. Parallel Synchronization and Communication Instructions: `mbarrier.init`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-mbarrier-init)
                *   [9.7.13.15.10. Parallel Synchronization and Communication Instructions: `mbarrier.inval`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-mbarrier-inval)
                *   [9.7.13.15.11. Parallel Synchronization and Communication Instructions: `mbarrier.expect_tx`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-mbarrier-expect-tx)
                *   [9.7.13.15.12. Parallel Synchronization and Communication Instructions: `mbarrier.complete_tx`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-mbarrier-complete-tx)
                *   [9.7.13.15.13. Parallel Synchronization and Communication Instructions: `mbarrier.arrive`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-mbarrier-arrive)
                *   [9.7.13.15.14. Parallel Synchronization and Communication Instructions: `mbarrier.arrive_drop`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-mbarrier-arrive-drop)
                *   [9.7.13.15.15. Parallel Synchronization and Communication Instructions: `cp.async.mbarrier.arrive`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-cp-async-mbarrier-arrive)
                *   [9.7.13.15.16. Parallel Synchronization and Communication Instructions: `mbarrier.test_wait` / `mbarrier.try_wait`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-mbarrier-test-wait-try-wait)
                *   [9.7.13.15.17. Parallel Synchronization and Communication Instructions: `mbarrier.pending_count`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-mbarrier-pending-count)

            *   [9.7.13.16. Parallel Synchronization and Communication Instructions: `tensormap.cp_fenceproxy`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-tensormap-cp-fenceproxy)
            *   [9.7.13.17. Parallel Synchronization and Communication Instructions: `clusterlaunchcontrol.try_cancel`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-clusterlaunchcontrol-try-cancel)
            *   [9.7.13.18. Parallel Synchronization and Communication Instructions: `clusterlaunchcontrol.query_cancel`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-clusterlaunchcontrol-query-cancel)

        *   [9.7.14. Warp Level Matrix Multiply-Accumulate Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-instructions)
            *   [9.7.14.1. Matrix Shape](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-shape)
            *   [9.7.14.2. Matrix Data-types](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-data-types)
            *   [9.7.14.3. Block Scaling for `mma.sync`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-block-scaling)
            *   [9.7.14.4. Matrix multiply-accumulate operation using `wmma` instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-instructions-wmma)
                *   [9.7.14.4.1. Matrix Fragments for WMMA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-fragment)
                *   [9.7.14.4.2. Matrix Storage for WMMA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-storage)
                *   [9.7.14.4.3. Warp-level Matrix Load Instruction: `wmma.load`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-instructions-wmma-ld)
                *   [9.7.14.4.4. Warp-level Matrix Store Instruction: `wmma.store`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-instructions-wmma-st)
                *   [9.7.14.4.5. Warp-level Matrix Multiply-and-Accumulate Instruction: `wmma.mma`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-instructions-wmma-mma)

            *   [9.7.14.5. Matrix multiply-accumulate operation using `mma` instruction](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-instructions-for-mma)
                *   [9.7.14.5.1. Matrix Fragments for `mma.m8n8k4` with `.f16` floating point type](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-fragment-mma-884-f16)
                *   [9.7.14.5.2. Matrix Fragments for `mma.m8n8k4` with `.f64` floating point type](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-fragment-mma-884-f64)
                *   [9.7.14.5.3. Matrix Fragments for `mma.m8n8k16`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-fragment-mma-8816)
                *   [9.7.14.5.4. Matrix Fragments for `mma.m8n8k32`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-fragment-mma-8832)
                *   [9.7.14.5.5. Matrix Fragments for `mma.m8n8k128`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-fragment-mma-88128)
                *   [9.7.14.5.6. Matrix Fragments for `mma.m16n8k4`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-fragment-mma-1684)
                *   [9.7.14.5.7. Matrix Fragments for `mma.m16n8k8`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-fragment-mma-1688)
                *   [9.7.14.5.8. Matrix Fragments for `mma.m16n8k16` with floating point type](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-fragment-mma-16816-float)
                *   [9.7.14.5.9. Matrix Fragments for `mma.m16n8k16` with integer type](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-fragment-mma-16816-i8-f8)
                *   [9.7.14.5.10. Matrix Fragments for `mma.m16n8k32`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-fragment-mma-16832)
                *   [9.7.14.5.11. Matrix Fragments for `mma.m16n8k64`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-fragment-mma-16864)
                *   [9.7.14.5.12. Matrix Fragments for `mma.m16n8k128`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-fragment-mma-168128)
                *   [9.7.14.5.13. Matrix Fragments for `mma.m16n8k256`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-fragment-mma-168256)
                *   [9.7.14.5.14. Multiply-and-Accumulate Instruction: `mma`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-instructions-mma)
                *   [9.7.14.5.15. Warp-level matrix load instruction: `ldmatrix`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-instructions-ldmatrix)
                *   [9.7.14.5.16. Warp-level matrix store instruction: `stmatrix`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-instructions-stmatrix)
                *   [9.7.14.5.17. Warp-level matrix transpose instruction: `movmatrix`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-instructions-movmatrix)

            *   [9.7.14.6. Matrix multiply-accumulate operation using `mma.sp` instruction with sparse matrix A](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-instructions-for-sparse-mma)
                *   [9.7.14.6.1. Sparse matrix storage](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-sparse-matrix-storage)
                *   [9.7.14.6.2. Matrix fragments for multiply-accumulate operation with sparse matrix A](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-fragments-for-sparse-mma)
                    *   [9.7.14.6.2.1. Matrix Fragments for sparse `mma.m16n8k16` with `.f16` and `.bf16` types](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-fragment-sparse-mma-16816-f16bf16)
                    *   [9.7.14.6.2.2. Matrix Fragments for sparse `mma.m16n8k32` with `.f16` and `.bf16` types](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-fragment-sparse-mma-16832-f16bf16)
                    *   [9.7.14.6.2.3. Matrix Fragments for sparse `mma.m16n8k16` with `.tf32` floating point type](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-fragment-sparse-mma-16816-tf32)
                    *   [9.7.14.6.2.4. Matrix Fragments for sparse `mma.m16n8k8` with `.tf32` floating point type](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-fragment-sparse-mma-1688-tf32)
                    *   [9.7.14.6.2.5. Matrix Fragments for sparse `mma.m16n8k32` with `.u8` / `.s8` integer type](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-fragment-sparse-mma-16832-u8s8)
                    *   [9.7.14.6.2.6. Matrix Fragments for sparse `mma.m16n8k64` with `.u8` / `.s8` / `.e4m3` / `.e5m2` type](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-fragment-sparse-mma-16864-u8s8-fp8)
                    *   [9.7.14.6.2.7. Matrix Fragments for sparse `mma.m16n8k64` with `.u4` / `.s4` integer type](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-fragment-sparse-mma-16864-u4s4)
                    *   [9.7.14.6.2.8. Matrix Fragments for sparse `mma.m16n8k128` with `.u4` / `.s4` integer type](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-fragment-sparse-mma-168128-u4s4)

                *   [9.7.14.6.3. Multiply-and-Accumulate Instruction: `mma.sp` / `mma.sp::ordered_metadata`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-instructions-sparse-mma)

        *   [9.7.15. Asynchronous Warpgroup Level Matrix Multiply-Accumulate Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#asynchronous-warpgroup-level-matrix-instructions)
            *   [9.7.15.1. Warpgroup](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#asynchronous-warpgroup-level-matrix-instructions-warpgroup)
            *   [9.7.15.2. Matrix Shape](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#asynchronous-warpgroup-level-matrix-shape)
            *   [9.7.15.3. Matrix Data-types](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#asynchronous-warpgroup-level-matrix-data-types)
            *   [9.7.15.4. Async Proxy](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#asynchronous-warpgroup-level-matrix-async-proxy)
            *   [9.7.15.5. Asynchronous Warpgroup Level Matrix Multiply-Accumulate Operation using `wgmma.mma_async` instruction](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-operation-wgmma-mma-async)
                *   [9.7.15.5.1. Register Fragments and Shared Memory Matrix Layouts](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#asynchronous-warpgroup-level-matrix-fragment)
                    *   [9.7.15.5.1.1. Register Fragments](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#asynchronous-warpgroup-level-matrix-register-fragment)
                        *   [9.7.15.5.1.1.1. Matrix Fragments for `wgmma.mma_async.m64nNk16`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-register-fragment-wgmma-64n16)
                        *   [9.7.15.5.1.1.2. Matrix Fragments for `wgmma.mma_async.m64nNk8`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-register-fragment-wgmma-64n8)
                        *   [9.7.15.5.1.1.3. Matrix Fragments for `wgmma.mma_async.m64nNk32`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-register-fragment-wgmma-64n32)
                        *   [9.7.15.5.1.1.4. Matrix Fragments for `wgmma.mma_async.m64nNk256`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-register-fragment-wgmma-64n256)

                    *   [9.7.15.5.1.2. Shared Memory Matrix Layout](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#asynchronous-warpgroup-level-matrix-shared-memory-layout)
                        *   [9.7.15.5.1.2.1. Major-ness supported by Strides](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#asynchronous-warpgroup-level-majorness-supported-by-strides)
                            *   [9.7.15.5.1.2.1.1. Leading Dimension Byte Offset](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#asynchronous-warpgroup-level-leading-dimension-byte-offset)
                            *   [9.7.15.5.1.2.1.2. Stride Dimension Byte Offset](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#asynchronous-warpgroup-level-stride-dimension-byte-offset)
                            *   [9.7.15.5.1.2.1.3. Canonical Layouts](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#asynchronous-warpgroup-level-canonical-layouts)

                        *   [9.7.15.5.1.2.2. Matrix Descriptor Format](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#asynchronous-warpgroup-level-matrix-shared-memory-layout-matrix-descriptor)

                *   [9.7.15.5.2. Asynchronous Multiply-and-Accumulate Instruction: `wgmma.mma_async`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-instructions-wgmma-mma)

            *   [9.7.15.6. Asynchronous Warpgroup Level Multiply-and-Accumulate Operation using `wgmma.mma_async.sp` instruction](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-instructions-for-sparse-wgmma)
                *   [9.7.15.6.1. Sparse matrix storage](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#asynchronous-warpgroup-level-sparse-matrix-storage)
                *   [9.7.15.6.2. Matrix fragments for warpgroup-level multiply-accumulate operation with sparse matrix A](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#asynchronous-warpgroup-level-matrix-fragments-for-sparse-wgmma)
                    *   [9.7.15.6.2.1. Matrix Fragments for sparse `wgmma.mma_async.m64nNk32`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-fragment-sparse-wgmma-64n32)
                    *   [9.7.15.6.2.2. Matrix Fragments for sparse `wgmma.mma_async.m64nNk16`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-fragment-sparse-wgmma-64n16)
                    *   [9.7.15.6.2.3. Matrix Fragments for sparse `wgmma.mma_async.m64nNk64`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-fragment-sparse-wgmma-64n64)

                *   [9.7.15.6.3. Asynchronous Multiply-and-Accumulate Instruction: `wgmma.mma_async.sp`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-instructions-wgmma-mma-sp)

            *   [9.7.15.7. Asynchronous `wgmma` Proxy Operations](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-wgmma-proxy-operations)
                *   [9.7.15.7.1. Asynchronous Multiply-and-Accumulate Instruction: `wgmma.fence`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-instructions-wgmma-fence)
                *   [9.7.15.7.2. Asynchronous Multiply-and-Accumulate Instruction: `wgmma.commit_group`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-instructions-wgmma-commit-group)
                *   [9.7.15.7.3. Asynchronous Multiply-and-Accumulate Instruction: `wgmma.wait_group`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-instructions-wgmma-wait-group)

        *   [9.7.16. TensorCore 5th Generation Family Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tensorcore-5th-generation-instructions)
            *   [9.7.16.1. Tensor Memory](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tensor-memory)
                *   [9.7.16.1.1. Tensor Memory Addressing](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tensor-memory-addressing)
                *   [9.7.16.1.2. Tensor Memory Allocation](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tensor-memory-allocation)

            *   [9.7.16.2. Matrix and Data Movement Shape](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-matrix-data-movement-shape)
                *   [9.7.16.2.1. Matrix Shape](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-matrix-shape)
                    *   [9.7.16.2.1.1. Target ISA Note](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-matrix-shape-target-isa-note)

                *   [9.7.16.2.2. Specifying Matrix Shape](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-specify-matrix-shape)
                *   [9.7.16.2.3. Data Movement Shape](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-data-movement-shape)
                    *   [9.7.16.2.3.1. Memory Layout](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-memory-layout)
                        *   [9.7.16.2.3.1.1. Matrix fragments for shape .32x32b](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-matrix-fragments-shape-3232b)
                        *   [9.7.16.2.3.1.2. Matrix fragments for shape .16x64b](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-matrix-fragments-shape-6464b)
                        *   [9.7.16.2.3.1.3. Matrix fragments for shape .16x128b](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-matrix-fragments-shape-16128b)
                        *   [9.7.16.2.3.1.4. Matrix fragments for shape .16x256b](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-matrix-fragments-shape-16256b)
                        *   [9.7.16.2.3.1.5. Matrix fragments for shape .16x32bx2](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-matrix-fragments-shape-1632b2)

            *   [9.7.16.3. Major-ness supported by Strides](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-majorness-supported-by-strides)
                *   [9.7.16.3.1. Leading Dimension Stride: relative offset or absolute address](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-leading-dimension-byte-offset)
                    *   [9.7.16.3.1.1. Relative offset mode](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-leading-dimension-byte-offset-relative-offset)
                    *   [9.7.16.3.1.2. Absolute address mode for K dimension being 48B](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-leading-dimension-byte-offset-absolute-address)
                        *   [9.7.16.3.1.2.1. Restrictions on the Leading Dimension Absolute Address Stride](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-leading-dimension-byte-offset-absolute-address-restriction)

                *   [9.7.16.3.2. Stride Dimension Byte Offset](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-stride-dimension-byte-offset)
                *   [9.7.16.3.3. Canonical Layouts](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-canonical-layouts)

            *   [9.7.16.4. Matrix Descriptors](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-matrix-descriptors)
                *   [9.7.16.4.1. Shared memory descriptor](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-shared-memory-descriptor)
                    *   [9.7.16.4.1.1. Target ISA Note](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-shared-memory-descriptor-target-isa-note)

                *   [9.7.16.4.2. Instruction descriptor](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-instruction-descriptor)
                *   [9.7.16.4.3. Zero-Column Mask Descriptor](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-zero-column-mask-descriptor)

            *   [9.7.16.5. Issue Granularity](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-issue-granularity)
                *   [9.7.16.5.1. CTA Pair](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-cta-pair)
                *   [9.7.16.5.2. Peer CTA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-peer-cta)

            *   [9.7.16.6. Memory Consistency Model for 5th generation of TensorCore operations](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-memory-consistency-model)
                *   [9.7.16.6.1. Asynchronous Operations](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-memory-consistency-model-async-operations)
                *   [9.7.16.6.2. Pipelined tcgen05 Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-memory-consistency-model-pipelined-instructions)
                    *   [9.7.16.6.2.1. Implicitly pipelined tcgen05 Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-memory-consistency-model-pipelined-instructions-implicit)
                        *   [9.7.16.6.2.1.1. mbarrier based completion mechanism](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-memory-consistency-model-mbarrier-completion)
                        *   [9.7.16.6.2.1.2. `tcgen05.wait` instruction based completion mechanism](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-memory-consistency-model-wait-completion)

                *   [9.7.16.6.3. Specialized Inter-thread Synchronization for tcgen05 instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-memory-consistency-model-inter-thread-sync)
                *   [9.7.16.6.4. Canonical synchronization patterns](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-memory-consistency-model-canonical-sync-patterns)
                    *   [9.7.16.6.4.1. Pipelined instructions, same thread](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-memory-consistency-model-canonical-sync-patterns-pipelined-same-thread)
                    *   [9.7.16.6.4.2. Non-pipelined instructions, same thread](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-memory-consistency-model-canonical-sync-patterns-non-pipelined-same-thread)
                    *   [9.7.16.6.4.3. Pipelined instructions, different thread](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-memory-consistency-model-canonical-sync-patterns-pipelined-diff-thread)
                    *   [9.7.16.6.4.4. Non-pipelined instructions, different thread](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-memory-consistency-model-canonical-sync-patterns-non-pipelined-diff-thread)
                    *   [9.7.16.6.4.5. Register dependencies, same thread](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-memory-consistency-model-canonical-sync-patterns-reg-dependency-same-thread)

                *   [9.7.16.6.5. Shared Memory Accesses](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-memory-consistency-model-smem-access)

            *   [9.7.16.7. Tensor Memory Allocation and Management Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-memory-alloc-manage-instructions)
                *   [9.7.16.7.1. Tensorcore 5th Generation Instructions: `tcgen05.alloc`, `tcgen05.dealloc`, `tcgen05.relinquish_alloc_permit`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-instructions-tcgen05-alloc-dealloc-relinquish-alloc-permit)

            *   [9.7.16.8. Tensor Memory and Register Load/Store Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-tensor-memory-ld-st)
                *   [9.7.16.8.1. Access restrictions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-tensor-memory-ld-st-access-restrictions)
                *   [9.7.16.8.2. Packing and Unpacking](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-tensor-memory-ld-st-packing-unpacking)
                *   [9.7.16.8.3. Tensorcore 5th Generation Instructions: `tcgen05.ld`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-instructions-tcgen05-ld)
                *   [9.7.16.8.4. Tensorcore 5th Generation Instructions: `tcgen05.st`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-instructions-tcgen05-st)
                *   [9.7.16.8.5. Tensorcore 5th Generation Instructions: `tcgen05.wait`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-instructions-tcgen05-wait)

            *   [9.7.16.9. Tensor Memory Data Movement Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-data-movement-instructions)
                *   [9.7.16.9.1. Optional Decompression](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-optional-decompression)
                    *   [9.7.16.9.1.1. Decompression of 4-bit floating point to 8-bit type](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-optional-decompression-4bit-8bit)
                    *   [9.7.16.9.1.2. Decompression of 6-bit floating point to 8-bit type](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-optional-decompression-6bit-8bit)

                *   [9.7.16.9.2. Tensorcore 5th Generation Instructions: `tcgen05.cp`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-instructions-tcgen05-cp)
                *   [9.7.16.9.3. Tensorcore 5th Generation Instructions: `tcgen05.shift`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-instructions-tcgen05-shift)

            *   [9.7.16.10. TensorCore 5th Generation Matrix Multiply and accumulate Operations](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-mma)
                *   [9.7.16.10.1. Transpose and Negate operations](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-transpose-and-negate-operations)
                *   [9.7.16.10.2. Matrix Layout Organization](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-matrix-layout-organization)
                *   [9.7.16.10.3. Valid Combinations of Type-Size, Major-ness and Swizzling](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-matrix-layout-organization-valid-comb-type-size-majorness-swizzle)
                *   [9.7.16.10.4. Packing formats of elements in Tensor and Shared memory](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-packing-formats)
                    *   [9.7.16.10.4.1. Packing format for matrix D in Tensor Memory](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-packing-formats-mat-d)
                    *   [9.7.16.10.4.2. Packing format for matrix A and B](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-packing-formats-mat-a-b)
                    *   [9.7.16.10.4.3. Packing format used for matrix A by `.kind::mxf8f6f4` in Tensor Memory](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-packing-formats-mxf8f6f4-tmem)
                    *   [9.7.16.10.4.4. Packing format used for matrix A and B by `.kind::mxf8f6f4` in Shared Memory](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-packing-formats-mxf8f6f4-smem)
                    *   [9.7.16.10.4.5. Packing format used for matrix A by `.kind::mxf4` and `.kind::mxf4nvf4` in Tensor Memory](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-packing-formats-mxf4-tmem)
                    *   [9.7.16.10.4.6. Packing format used for matrix A and B by `.kind::mxf4` and `.kind::mxf4nvf4` in Shared Memory](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-packing-formats-mxf4-smem)

                *   [9.7.16.10.5. Data Path Layout Organization](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-data-path-layout-organization)
                    *   [9.7.16.10.5.1. Layout A (M = 256)](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-data-path-layout-a)
                    *   [9.7.16.10.5.2. Layout B (M = 128 + cta-group::2 + Dense A matrix)](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-data-path-layout-b)
                    *   [9.7.16.10.5.3. Layout C (M = 128 + cta-group::2 + Sparse A matrix)](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-data-path-layout-c)
                    *   [9.7.16.10.5.4. Layout D (M = 128 + cta-group::1)](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-data-path-layout-d)
                    *   [9.7.16.10.5.5. Layout E (M = 64 + .ws mode)](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-data-path-layout-e)
                    *   [9.7.16.10.5.6. Layout F (M = 64 + non .ws mode)](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-data-path-layout-f)
                    *   [9.7.16.10.5.7. Layout G (M = 32)](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-data-path-layout-g)

                *   [9.7.16.10.6. Shared Memory Layout and Swizzling](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-shared-memory-layout-swizzling)
                *   [9.7.16.10.7. Block Scaling for `tcgen05.mma.sync`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-block-scaling)
                    *   [9.7.16.10.7.1. Valid combinations of scale_vectorsize with types and MMA-Kind](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-mma-scale-valid-vec-size)
                    *   [9.7.16.10.7.2. Scale Factor A ID](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-mma-scale-factor-a)
                        *   [9.7.16.10.7.2.1. Layout of the Scale Factor A Matrix for scale_vec::1X/block32 with K=32/K=64](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-mma-scale-factor-a-layout-1x)
                        *   [9.7.16.10.7.2.2. Layout of the Scale Factor A Matrix for scale_vec::2X/block32 with K=64/K=128](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-mma-scale-factor-a-layout-2x)
                        *   [9.7.16.10.7.2.3. Layout of the Scale Factor A Matrix for scale_vec::4X/block16 with K=64/K=128](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-mma-scale-factor-a-layout-4x)
                        *   [9.7.16.10.7.2.4. Layout of the Scale Factor A Matrix for block32 with K=96 (Semantically equivalent to scale_vec::3X)](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-mma-scale-factor-a-layout-block32-k96)
                        *   [9.7.16.10.7.2.5. Layout of the Scale Factor A Matrix for block16 with K=96 (Semantically equivalent to scale_vec::6X)](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-mma-scale-factor-a-layout-block16-k96)

                    *   [9.7.16.10.7.3. Scale Factor B ID](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-mma-scale-factor-b)
                        *   [9.7.16.10.7.3.1. Layout of the Scale Factor B Matrix for scale_vec::1X/block32 with K=32/K=64](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-mma-scale-factor-b-layout-1x)
                        *   [9.7.16.10.7.3.2. Layout of the Scale Factor B Matrix for scale_vec::2X/block32 with K=64/K=128](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-mma-scale-factor-b-layout-2x)
                        *   [9.7.16.10.7.3.3. Layout of the Scale Factor B Matrix for scale_vec::4X/block16 with K=64/K=128](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-mma-scale-factor-b-layout-4x)
                        *   [9.7.16.10.7.3.4. Layout of the Scale Factor B Matrix for block32 with K=96 (Semantically equivalent to scale_vec::3X)](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-mma-scale-factor-b-layout-block32-k96)
                        *   [9.7.16.10.7.3.5. Layout of the Scale Factor B Matrix for block16 with K=96 (Semantically equivalent to scale_vec::6X)](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-mma-scale-factor-b-layout-block16-k96)

                *   [9.7.16.10.8. Sparse Matrices](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-sparse-matrices)
                    *   [9.7.16.10.8.1. Sparse `tcgen05.mma.sp` with `.kind::tf32`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-sparse-matrices-kind-tf32)
                    *   [9.7.16.10.8.2. Sparse `tcgen05.mma.sp` with `.kind::f16`, `.kind::f8f6f4`, `.kind::mxf8f6f4`, `.kind::i8`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-sparse-matrices-kind-f16-f8f8f4-mxf8f6f4)
                    *   [9.7.16.10.8.3. Sparse `tcgen05.mma.sp` with `.kind::mxf4` and `.kind::mxf4nvf4`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-sparse-matrices-kind-mxf4)
                    *   [9.7.16.10.8.4. Sparsity selector](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-sparse-matrices-sparsity-selector)
                        *   [9.7.16.10.8.4.1. Layout of the Sparsity Metadata Matrix for M = 64 for `.kind::f16`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-sparse-matrices-sparsity-selector-kind-f16-m64)
                        *   [9.7.16.10.8.4.2. Layout of the Sparsity Metadata Matrix for M = 128 / M = 256 for `.kind::f16`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-sparse-matrices-sparsity-selector-kind-f16-m128-256)
                        *   [9.7.16.10.8.4.3. Layout of the Sparsity Metadata Matrix for M = 64 for `.kind::tf32`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-sparse-matrices-sparsity-selector-kind-tf32-m64)
                        *   [9.7.16.10.8.4.4. Layout of the Sparsity Metadata Matrix for M = 128 / M = 256 for `.kind::tf32`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-sparse-matrices-sparsity-selector-kind-tf32-m128-256)
                        *   [9.7.16.10.8.4.5. Layout of the Sparsity Metadata Matrix for M = 64 for `.kind::f8f6f4`, `.kind::mxf8f6f4`, `.kind::i8`, `.kind::mxf4`, `.kind::mxf4nvf4`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-sparse-matrices-sparsity-selector-kind-f8f6f4-mxf8f6f4-m64)
                        *   [9.7.16.10.8.4.6. Layout of the Sparsity Metadata Matrix for M = 128 / M = 256 for `.kind::f8f6f4`, `.kind::mxf8f6f4`, `.kind::i8`, `.kind::mxf4`, `.kind::mxf4nvf4`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-sparse-matrices-sparsity-selector-kind-f8f6f4-mxf8f6f4-m128-256)

                    *   [9.7.16.10.8.5. Alignment restriction](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-sparse-matrices-alignment-restriction)

                *   [9.7.16.10.9. TensorCore 5th Generation of MMA Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-mma-instructions)
                    *   [9.7.16.10.9.1. TensorCore 5th Generation Instructions: `tcgen05.mma`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-mma-instructions-mma)
                    *   [9.7.16.10.9.2. TensorCore 5th Generation Instructions: `tcgen05.mma.sp`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-mma-instructions-mma-sp)
                    *   [9.7.16.10.9.3. TensorCore 5th Generation Instructions: `tcgen05.mma.ws`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-mma-instructions-mma-ws)
                    *   [9.7.16.10.9.4. TensorCore 5th Generation Instructions: `tcgen05.mma.ws.sp`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-mma-instructions-mma-ws-sp)

            *   [9.7.16.11. TensorCore 5th Generation Specialized Synchronization Operations](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-special-sync-operations)
                *   [9.7.16.11.1. TensorCore 5th Generation Instructions: `tcgen05.fence`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-special-sync-operations-fence)

            *   [9.7.16.12. TensorCore 5th Generation Async Synchronization Operations](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen-async-sync-operations)
                *   [9.7.16.12.1. TensorCore 5th Generation Instructions: `tcgen05.commit`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen-async-sync-operations-commit)

        *   [9.7.17. Stack Manipulation Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#stack-manipulation-instructions)
            *   [9.7.17.1. Stack Manipulation Instructions: `stacksave`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#stack-manipulation-instructions-stacksave)
            *   [9.7.17.2. Stack Manipulation Instructions: `stackrestore`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#stack-manipulation-instructions-stackrestore)
            *   [9.7.17.3. Stack Manipulation Instructions: `alloca`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#stack-manipulation-instructions-alloca)

        *   [9.7.18. Video Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#video-instructions)
            *   [9.7.18.1. Scalar Video Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#scalar-video-instructions)
                *   [9.7.18.1.1. Scalar Video Instructions: `vadd`, `vsub`, `vabsdiff`, `vmin`, `vmax`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#scalar-video-instructions-vadd-vsub-vabsdiff-vmin-vmax)
                *   [9.7.18.1.2. Scalar Video Instructions: `vshl`, `vshr`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#scalar-video-instructions-vshl-vshr)
                *   [9.7.18.1.3. Scalar Video Instructions: `vmad`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#scalar-video-instructions-vmad)
                *   [9.7.18.1.4. Scalar Video Instructions: `vset`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#scalar-video-instructions-vset)

            *   [9.7.18.2. SIMD Video Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#simd-video-instructions)
                *   [9.7.18.2.1. SIMD Video Instructions: `vadd2`, `vsub2`, `vavrg2`, `vabsdiff2`, `vmin2`, `vmax2`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#simd-video-instructions-vadd2-vsub2-vavrg2-vabsdiff2-vmin2-vmax2)
                *   [9.7.18.2.2. SIMD Video Instructions: `vset2`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#simd-video-instructions-vset2)
                *   [9.7.18.2.3. SIMD Video Instructions: `vadd4`, `vsub4`, `vavrg4`, `vabsdiff4`, `vmin4`, `vmax4`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#simd-video-instructions-vadd4-vsub4-vavrg4-vabsdiff4-vmin4-vmax4)
                *   [9.7.18.2.4. SIMD Video Instructions: `vset4`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#simd-video-instructions-vset4)

        *   [9.7.19. Miscellaneous Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#miscellaneous-instructions)
            *   [9.7.19.1. Miscellaneous Instructions: `brkpt`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#miscellaneous-instructions-brkpt)
            *   [9.7.19.2. Miscellaneous Instructions: `nanosleep`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#miscellaneous-instructions-nanosleep)
            *   [9.7.19.3. Miscellaneous Instructions: `pmevent`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#miscellaneous-instructions-pmevent)
            *   [9.7.19.4. Miscellaneous Instructions: `trap`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#miscellaneous-instructions-trap)
            *   [9.7.19.5. Miscellaneous Instructions: `setmaxnreg`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#miscellaneous-instructions-setmaxnreg)

*   [10. Special Registers](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#special-registers)
    *   [10.1. Special Registers: `%tid`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers-tid)
    *   [10.2. Special Registers: `%ntid`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers-ntid)
    *   [10.3. Special Registers: `%laneid`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers-laneid)
    *   [10.4. Special Registers: `%warpid`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers-warpid)
    *   [10.5. Special Registers: `%nwarpid`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers-nwarpid)
    *   [10.6. Special Registers: `%ctaid`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers-ctaid)
    *   [10.7. Special Registers: `%nctaid`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers-nctaid)
    *   [10.8. Special Registers: `%smid`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers-smid)
    *   [10.9. Special Registers: `%nsmid`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers-nsmid)
    *   [10.10. Special Registers: `%gridid`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers-gridid)
    *   [10.11. Special Registers: `%is_explicit_cluster`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers-is-explicit-cluster)
    *   [10.12. Special Registers: `%clusterid`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers-clusterid)
    *   [10.13. Special Registers: `%nclusterid`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers-nclusterid)
    *   [10.14. Special Registers: `%cluster_ctaid`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers-cluster-ctaid)
    *   [10.15. Special Registers: `%cluster_nctaid`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers-cluster-nctaid)
    *   [10.16. Special Registers: `%cluster_ctarank`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers-cluster-ctarank)
    *   [10.17. Special Registers: `%cluster_nctarank`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers-cluster-nctarank)
    *   [10.18. Special Registers: `%lanemask_eq`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers-lanemask-eq)
    *   [10.19. Special Registers: `%lanemask_le`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers-lanemask-le)
    *   [10.20. Special Registers: `%lanemask_lt`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers-lanemask-lt)
    *   [10.21. Special Registers: `%lanemask_ge`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers-lanemask-ge)
    *   [10.22. Special Registers: `%lanemask_gt`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers-lanemask-gt)
    *   [10.23. Special Registers: `%clock`, `%clock_hi`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers-clock)
    *   [10.24. Special Registers: `%clock64`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers-clock64)
    *   [10.25. Special Registers: `%pm0` … `%pm7`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers-pm0-pm7)
    *   [10.26. Special Registers: `%pm0_64` … `%pm7_64`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers-pm0-64-pm7-64)
    *   [10.27. Special Registers: `%envreg<32>`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers-envreg-32)
    *   [10.28. Special Registers: `%globaltimer`, `%globaltimer_lo`, `%globaltimer_hi`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers-globaltimer)
    *   [10.29. Special Registers: `%reserved_smem_offset_begin`, `%reserved_smem_offset_end`, `%reserved_smem_offset_cap`, `%reserved_smem_offset_<2>`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers-reserved-smem)
    *   [10.30. Special Registers: `%total_smem_size`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers-total-smem-size)
    *   [10.31. Special Registers: `%aggr_smem_size`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers-aggr-smem-size)
    *   [10.32. Special Registers: `%dynamic_smem_size`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers-dynamic-smem-size)
    *   [10.33. Special Registers: `%current_graph_exec`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers-current-graph-exec)

*   [11. Directives](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#directives)
    *   [11.1. PTX Module Directives](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#ptx-module-directives)
        *   [11.1.1. PTX Module Directives: `.version`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#ptx-module-directives-version)
        *   [11.1.2. PTX Module Directives: `.target`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#ptx-module-directives-target)
        *   [11.1.3. PTX Module Directives: `.address_size`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#ptx-module-directives-address-size)

    *   [11.2. Specifying Kernel Entry Points and Functions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#specifying-kernel-entry-points-and-functions)
        *   [11.2.1. Kernel and Function Directives: `.entry`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#kernel-and-function-directives-entry)
        *   [11.2.2. Kernel and Function Directives: `.func`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#kernel-and-function-directives-func)
        *   [11.2.3. Kernel and Function Directives: `.alias`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#kernel-and-function-directives-alias)

    *   [11.3. Control Flow Directives](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#control-flow-directives)
        *   [11.3.1. Control Flow Directives: `.branchtargets`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#control-flow-directives-branchtargets)
        *   [11.3.2. Control Flow Directives: `.calltargets`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#control-flow-directives-calltargets)
        *   [11.3.3. Control Flow Directives: `.callprototype`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#control-flow-directives-callprototype)

    *   [11.4. Performance-Tuning Directives](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#performance-tuning-directives)
        *   [11.4.1. Performance-Tuning Directives: `.maxnreg`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#performance-tuning-directives-maxnreg)
        *   [11.4.2. Performance-Tuning Directives: `.maxntid`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#performance-tuning-directives-maxntid)
        *   [11.4.3. Performance-Tuning Directives: `.reqntid`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#performance-tuning-directives-reqntid)
        *   [11.4.4. Performance-Tuning Directives: `.minnctapersm`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#performance-tuning-directives-minnctapersm)
        *   [11.4.5. Performance-Tuning Directives: `.maxnctapersm` (deprecated)](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#performance-tuning-directives-maxnctapersm)
        *   [11.4.6. Performance-Tuning Directives: `.noreturn`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#performance-tuning-directives-noreturn)
        *   [11.4.7. Performance-Tuning Directives: `.pragma`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#performance-tuning-directives-pragma)
        *   [11.4.8. Performance-Tuning Directives: `.abi_preserve`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#performance-tuning-directives-abi-preserve)
        *   [11.4.9. Performance-Tuning Directives: `.abi_preserve_control`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#performance-tuning-directives-abi-preserve-control)

    *   [11.5. Debugging Directives](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#debugging-directives)
        *   [11.5.1. Debugging Directives: `@@dwarf`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#debugging-directives-atatdwarf)
        *   [11.5.2. Debugging Directives: `.section`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#debugging-directives-section)
        *   [11.5.3. Debugging Directives: `.file`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#debugging-directives-file)
        *   [11.5.4. Debugging Directives: `.loc`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#debugging-directives-loc)

    *   [11.6. Linking Directives](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#linking-directives)
        *   [11.6.1. Linking Directives: `.extern`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#linking-directives-extern)
        *   [11.6.2. Linking Directives: `.visible`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#linking-directives-visible)
        *   [11.6.3. Linking Directives: `.weak`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#linking-directives-weak)
        *   [11.6.4. Linking Directives: `.common`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#linking-directives-common)

    *   [11.7. Cluster Dimension Directives](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#cluster-dimension-directives)
        *   [11.7.1. Cluster Dimension Directives: `.reqnctapercluster`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#cluster-dimension-directives-reqnctapercluster)
        *   [11.7.2. Cluster Dimension Directives: `.explicitcluster`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#cluster-dimension-directives-explicitcluster)
        *   [11.7.3. Cluster Dimension Directives: `.maxclusterrank`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#cluster-dimension-directives-maxclusterrank)

    *   [11.8. Miscellaneous Directives](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#miscellaneous-directives)
        *   [11.8.1. Miscellaneous Directives: `.blocksareclusters`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#miscellaneous-directives-blocksareclusters)

*   [12. Descriptions of `.pragma` Strings](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#descriptions-pragma-strings)
    *   [12.1. Pragma Strings: `"nounroll"`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#pragma-strings-nounroll)
    *   [12.2. Pragma Strings: `"used_bytes_mask"`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#pragma-strings-used-bytes-mask)
    *   [12.3. Pragma Strings: `"enable_smem_spilling"`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#pragma-strings-enable-smem-spilling)
    *   [12.4. Pragma Strings: `"frequency"`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#pragma-strings-frequency)

*   [13. Release Notes](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#release-notes)
    *   [13.1. Changes in PTX ISA Version 9.1](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-9-1)
    *   [13.2. Changes in PTX ISA Version 9.0](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-9-0)
    *   [13.3. Changes in PTX ISA Version 8.8](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-8-8)
    *   [13.4. Changes in PTX ISA Version 8.7](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-8-7)
    *   [13.5. Changes in PTX ISA Version 8.6](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-8-6)
    *   [13.6. Changes in PTX ISA Version 8.5](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-8-5)
    *   [13.7. Changes in PTX ISA Version 8.4](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-8-4)
    *   [13.8. Changes in PTX ISA Version 8.3](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-8-3)
    *   [13.9. Changes in PTX ISA Version 8.2](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-8-2)
    *   [13.10. Changes in PTX ISA Version 8.1](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-8-1)
    *   [13.11. Changes in PTX ISA Version 8.0](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-8-0)
    *   [13.12. Changes in PTX ISA Version 7.8](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-7-8)
    *   [13.13. Changes in PTX ISA Version 7.7](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-7-7)
    *   [13.14. Changes in PTX ISA Version 7.6](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-7-6)
    *   [13.15. Changes in PTX ISA Version 7.5](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-7-5)
    *   [13.16. Changes in PTX ISA Version 7.4](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-7-4)
    *   [13.17. Changes in PTX ISA Version 7.3](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-7-3)
    *   [13.18. Changes in PTX ISA Version 7.2](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-7-2)
    *   [13.19. Changes in PTX ISA Version 7.1](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-7-1)
    *   [13.20. Changes in PTX ISA Version 7.0](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-7-0)
    *   [13.21. Changes in PTX ISA Version 6.5](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-6-5)
    *   [13.22. Changes in PTX ISA Version 6.4](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-6-4)
    *   [13.23. Changes in PTX ISA Version 6.3](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-6-3)
    *   [13.24. Changes in PTX ISA Version 6.2](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-6-2)
    *   [13.25. Changes in PTX ISA Version 6.1](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-6-1)
    *   [13.26. Changes in PTX ISA Version 6.0](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-6-0)
    *   [13.27. Changes in PTX ISA Version 5.0](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-5-0)
    *   [13.28. Changes in PTX ISA Version 4.3](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-4-3)
    *   [13.29. Changes in PTX ISA Version 4.2](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-4-2)
    *   [13.30. Changes in PTX ISA Version 4.1](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-4-1)
    *   [13.31. Changes in PTX ISA Version 4.0](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-4-0)
    *   [13.32. Changes in PTX ISA Version 3.2](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-3-2)
    *   [13.33. Changes in PTX ISA Version 3.1](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-3-1)
    *   [13.34. Changes in PTX ISA Version 3.0](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-3-0)
    *   [13.35. Changes in PTX ISA Version 2.3](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-2-3)
    *   [13.36. Changes in PTX ISA Version 2.2](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-2-2)
    *   [13.37. Changes in PTX ISA Version 2.1](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-2-1)
    *   [13.38. Changes in PTX ISA Version 2.0](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-2-0)

*   [14. Notices](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#notices)
    *   [14.1. Notice](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#notice)
    *   [14.2. OpenCL](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#opencl)
    *   [14.3. Trademarks](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#trademarks)

* * *

![Image 3](https://docs.nvidia.com/cuda/_static/NVIDIA-LogoBlack.svg)![Image 4](https://docs.nvidia.com/cuda/_static/NVIDIA-LogoWhite.svg)
[Privacy Policy](https://www.nvidia.com/en-us/about-nvidia/privacy-policy.md/) | [Manage My Privacy](https://www.nvidia.com/en-us/about-nvidia/privacy-center.md/) | [Do Not Sell or Share My Data](https://www.nvidia.com/en-us/preferences/start.md/) | [Terms of Service](https://www.nvidia.com/en-us/about-nvidia/terms-of-service.md/) | [Accessibility](https://www.nvidia.com/en-us/about-nvidia/accessibility.md/) | [Corporate Policies](https://www.nvidia.com/en-us/about-nvidia/company-policies.md/) | [Product Security](https://www.nvidia.com/en-us/product-security.md/) | [Contact](https://www.nvidia.com/en-us/contact/)

Copyright © 2007-2025, NVIDIA Corporation & affiliates. All rights reserved.

Last updated on Dec 02, 2025.

Links/Buttons:
- [](https://docs.nvidia.com/cuda/index.html)
- [PDF](https://docs.nvidia.com/cuda/pdf/ptx_isa_9.1.pdf)
- [Archive](https://developer.nvidia.com/cuda-toolkit-archive.md)
- [](https://docs.nvidia.com/cuda/parallel-thread-execution/contents.html.md#contents)
- [1. Introduction](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html)
- [1.1. Scalable Data-Parallel Computing using GPUs](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#scalable-data-parallel-computing-using-gpus)
- [1.2. Goals of PTX](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#goals-of-ptx)
- [1.3. PTX ISA Version 9.1](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#ptx-isa-version-9-1)
- [1.4. Document Structure](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#document-structure)
- [2. Programming Model](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#programming-model)
- [2.1. A Highly Multithreaded Coprocessor](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#highly-multithreaded-coprocessor)
- [2.2. Thread Hierarchy](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#thread-hierarchy)
- [2.2.1. Cooperative Thread Arrays](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#cooperative-thread-arrays)
- [2.2.2. Cluster of Cooperative Thread Arrays](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#cluster-of-cooperative-thread-arrays)
- [2.2.3. Grid of Clusters](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#grid-of-clusters)
- [2.3. Memory Hierarchy](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#memory-hierarchy)
- [3. PTX Machine Model](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#ptx-machine-model)
- [3.1. A Set of SIMT Multiprocessors](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#set-of-simt-multiprocessors)
- [3.2. Independent Thread Scheduling](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#independent-thread-scheduling)
- [3.3. On-chip Shared Memory](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#on-chip-shared-memory)
- [4. Syntax](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#syntax)
- [4.1. Source Format](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#source-format)
- [4.2. Comments](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#comments)
- [4.3. Statements](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#statements)
- [4.3.1. Directive Statements](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#directive-statements)
- [4.3.2. Instruction Statements](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#instruction-statements)
- [4.4. Identifiers](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#identifiers)
- [4.5. Constants](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#constants)
- [4.5.1. Integer Constants](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-constants)
- [4.5.2. Floating-Point Constants](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#floating-point-constants)
- [4.5.3. Predicate Constants](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#predicate-constants)
- [4.5.4. Constant Expressions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#constant-expressions)
- [4.5.5. Integer Constant Expression Evaluation](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-constant-expression-evaluation)
- [4.5.6. Summary of Constant Expression Evaluation Rules](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#summary-of-constant-expression-evaluation-rules)
- [5. State Spaces, Types, and Variables](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#state-spaces-types-and-variables)
- [5.1. State Spaces](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#state-spaces)
- [5.1.1. Register State Space](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#register-state-space)
- [5.1.2. Special Register State Space](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-register-state-space)
- [5.1.3. Constant State Space](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#constant-state-space)
- [5.1.3.1. Banked Constant State Space (deprecated)](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#banked-constant-state-space-deprecated)
- [5.1.4. Global State Space](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#global-state-space)
- [5.1.5. Local State Space](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#local-state-space)
- [5.1.6. Parameter State Space](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parameter-state-space)
- [5.1.6.1. Kernel Function Parameters](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#kernel-function-parameters)
- [5.1.6.2. Kernel Function Parameter Attributes](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#kernel-function-parameter-attributes)
- [5.1.6.3. Kernel Parameter Attribute: .ptr](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#kernel-parameter-attribute-ptr)
- [5.1.6.4. Device Function Parameters](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#device-function-parameters)
- [5.1.7. Shared State Space](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#shared-state-space)
- [5.1.8. Texture State Space (deprecated)](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#texture-state-space-deprecated)
- [5.2. Types](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#types)
- [5.2.1. Fundamental Types](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#fundamental-types)
- [5.2.2. Restricted Use of Sub-Word Sizes](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#restricted-use-of-sub-word-sizes)
- [5.2.3. Alternate Floating-Point Data Formats](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#alternate-floating-point-data-formats)
- [5.2.4. Fixed-point Data format](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#fixed-point-data-formats)
- [5.2.5. Packed Data Types](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#packed-data-types)
- [5.2.5.1. Packed Floating Point Data Types](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#packed-floating-point-data-types)
- [5.2.5.2. Packed Integer Data Types](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#packed-integer-data-types)
- [5.2.5.3. Packed Fixed-Point Data Types](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#packed-fixed-point-data-types)
- [5.3. Texture Sampler and Surface Types](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#texture-sampler-and-surface-types)
- [5.3.1. Texture and Surface Properties](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#texture-surface-properties)
- [5.3.2. Sampler Properties](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#sampler-properties)
- [5.3.3. Channel Data Type and Channel Order Fields](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#channel-data-type-and-channel-order-fields)
- [5.4. Variables](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#variables)
- [5.4.1. Variable Declarations](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#variable-declarations)
- [5.4.2. Vectors](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#vectors)
- [5.4.3. Array Declarations](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#array-declarations)
- [5.4.4. Initializers](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#initializers)
- [5.4.5. Alignment](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#alignment)
- [5.4.6. Parameterized Variable Names](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parameterized-variable-names)
- [5.4.7. Variable Attributes](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#variable-attributes)
- [5.4.8. Variable and Function Attribute Directive: .attribute](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#variable-and-function-attribute-directive-attribute)
- [5.5. Tensors](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tensors)
- [5.5.1. Tensor Dimension, size and format](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tensor-dimension-size-format)
- [5.5.1.1. Sub-byte Types](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tensor-dimension-size-format-sub-bytes)
- [5.5.1.1.1. Padding and alignment of the sub-byte types](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tensor-dimension-size-format-sub-bytes-padding-align)
- [5.5.2. Tensor Access Modes](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tensor-access-modes)
- [5.5.3. Tiled Mode](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tensor-tiled-mode)
- [5.5.3.1. Bounding Box](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tensor-tiled-mode-bounding-box)
- [5.5.3.2. Traversal-Stride](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tensor-tiled-mode-traversal-stride)
- [5.5.3.3. Out of Boundary Access](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tensor-tiled-mode-oob-access)
- [5.5.3.4. .tile::scatter4 and .tile::gather4 modes](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tensor-tiled-scatter4-gather4-modes)
- [5.5.3.4.1. Bounding Box](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tensor-tiled-scatter4-gather4-modes-bounding-box)
- [5.5.4. im2col mode](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tensor-im2col-mode)
- [5.5.4.1. Bounding Box](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tensor-im2col-mode-bounding-box)
- [5.5.4.2. Traversal Stride](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tensor-im2col-mode-traversal-stride)
- [5.5.4.3. Out of Boundary Access](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tensor-im2col-mode-oob-access)
- [5.5.5. im2col::w and im2col::w::128 modes](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tensor-im2col-w-w128-modes)
- [5.5.5.1. Bounding Box](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tensor-im2col-w-w128-modes-bounding-box)
- [5.5.5.2. Traversal Stride](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tensor-im2col-w-w128-modes-traversal-stride)
- [5.5.5.3. wHalo](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tensor-im2col-w-w128-modes-whalo)
- [5.5.5.4. wOffset](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tensor-im2col-w-w128-modes-woffset)
- [5.5.6. Interleave layout](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tensor-interleaved-layout)
- [5.5.7. Swizzling Modes](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tensor-swizzling-modes)
- [5.5.8. Tensor-map](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tensor-tensormap)
- [6. Instruction Operands](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#instruction-operands)
- [6.1. Operand Type Information](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#operand-type-information)
- [6.2. Source Operands](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#source-operands)
- [6.3. Destination Operands](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#destination-operands)
- [6.4. Using Addresses, Arrays, and Vectors](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#using-addresses-arrays-and-vectors)
- [6.4.1. Addresses as Operands](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#addresses-as-operands)
- [6.4.1.1. Generic Addressing](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#generic-addressing)
- [6.4.2. Arrays as Operands](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#arrays-as-operands)
- [6.4.3. Vectors as Operands](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#vectors-as-operands)
- [6.4.4. Labels and Function Names as Operands](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#labels-and-function-names-as-operands)
- [6.5. Type Conversion](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#type-conversion)
- [6.5.1. Scalar Conversions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#scalar-conversions)
- [6.5.2. Rounding Modifiers](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#rounding-modifiers)
- [6.6. Operand Costs](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#operand-costs)
- [7. Abstracting the ABI](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#abstracting-abi)
- [7.1. Function Declarations and Definitions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#function-declarations-and-definitions)
- [7.1.1. Changes from PTX ISA Version 1.x](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#changes-from-ptx-isa-version-1-x)
- [7.2. Variadic Functions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#variadic-functions)
- [7.3. Alloca](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#alloca)
- [8. Memory Consistency Model](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#memory-consistency-model)
- [8.1. Scope and applicability of the model](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#scope-and-applicability)
- [8.1.1. Limitations on atomicity at system scope](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#limitations-system-scope-atomicity)
- [8.2. Memory operations](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#memory-operations)
- [8.2.1. Overlap](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#overlap)
- [8.2.2. Aliases](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#aliases)
- [8.2.3. Multimem Addresses](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#multimem-addresses)
- [8.2.4. Memory Operations on Vector Data Types](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#memory-operations-on-vector-data-types)
- [8.2.5. Memory Operations on Packed Data Types](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#memory-operations-on-packed-data-types)
- [8.2.6. Initialization](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#initialization)
- [8.3. State spaces](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#memory-consistency-state-spaces)
- [8.4. Operation types](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#operation-types)
- [8.4.1. mmio Operation](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#mmio-operation)
- [8.4.2. volatile Operation](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#volatile-operation)
- [8.5. Scope](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#scope)
- [8.6. Proxies](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#proxies)
- [8.7. Morally strong operations](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#morally-strong-operations)
- [8.7.1. Conflict and Data-races](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#conflict-and-data-races)
- [8.7.2. Limitations on Mixed-size Data-races](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#mixed-size-limitations)
- [8.8. Release and Acquire Patterns](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#release-acquire-patterns)
- [8.9. Ordering of memory operations](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#ordering-memory-operations)
- [8.9.1. Program Order](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#program-order)
- [8.9.1.1. Asynchronous Operations](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#program-order-async-operations)
- [8.9.2. Observation Order](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#observation-order)
- [8.9.3. Fence-SC Order](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#fence-sc-order)
- [8.9.4. Memory synchronization](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#memory-synchronization)
- [8.9.5. Causality Order](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#causality-order)
- [8.9.6. Coherence Order](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#coherence-order)
- [8.9.7. Communication Order](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#communication-order)
- [8.10. Axioms](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#axioms)
- [8.10.1. Coherence](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#coherence-axiom)
- [8.10.2. Fence-SC](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#fence-sc-axiom)
- [8.10.3. Atomicity](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#atomicity-axiom)
- [8.10.4. No Thin Air](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#no-thin-air-axiom)
- [8.10.5. Sequential Consistency Per Location](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#sc-per-loc-axiom)
- [8.10.6. Causality](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#causality-axiom)
- [8.11. Special Cases](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-cases)
- [8.11.1. Reductions do not form Acquire Patterns](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#red-read)
- [9. Instruction Set](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#instruction-set)
- [9.1. Format and Semantics of Instruction Descriptions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#format-and-semantics-of-instruction-descriptions)
- [9.2. PTX Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#ptx-instructions)
- [9.3. Predicated Execution](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#predicated-execution)
- [9.3.1. Comparisons](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#comparisons)
- [9.3.1.1. Integer and Bit-Size Comparisons](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-and-bit-size-comparisons)
- [9.3.1.2. Floating Point Comparisons](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#floating-point-comparisons)
- [9.3.2. Manipulating Predicates](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#manipulating-predicates)
- [9.4. Type Information for Instructions and Operands](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#type-information-for-instructions-and-operands)
- [9.4.1. Operand Size Exceeding Instruction-Type Size](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#operand-size-exceeding-instruction-type-size)
- [9.5. Divergence of Threads in Control Constructs](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#divergence-of-threads-in-control-constructs)
- [9.6. Semantics](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#semantics)
- [9.6.1. Machine-Specific Semantics of 16-bit Code](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#machine-specific-semantics-of-16-bit-code)
- [9.7. Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#instructions)
- [9.7.1. Integer Arithmetic Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions)
- [9.7.1.1. Integer Arithmetic Instructions: add](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#integer-arithmetic-instructions-add)
- [9.7.1.2. Integer Arithmetic Instructions: sub](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#integer-arithmetic-instructions-sub)
- [9.7.1.3. Integer Arithmetic Instructions: mul](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#integer-arithmetic-instructions-mul)
- [9.7.1.4. Integer Arithmetic Instructions: mad](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#integer-arithmetic-instructions-mad)
- [9.7.1.5. Integer Arithmetic Instructions: mul24](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#integer-arithmetic-instructions-mul24)
- [9.7.1.6. Integer Arithmetic Instructions: mad24](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#integer-arithmetic-instructions-mad24)
- [9.7.1.7. Integer Arithmetic Instructions: sad](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#integer-arithmetic-instructions-sad)
- [9.7.1.8. Integer Arithmetic Instructions: div](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#integer-arithmetic-instructions-div)
- [9.7.1.9. Integer Arithmetic Instructions: rem](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#integer-arithmetic-instructions-rem)
- [9.7.1.10. Integer Arithmetic Instructions: abs](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#integer-arithmetic-instructions-abs)
- [9.7.1.11. Integer Arithmetic Instructions: neg](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#integer-arithmetic-instructions-neg)
- [9.7.1.12. Integer Arithmetic Instructions: min](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#integer-arithmetic-instructions-min)
- [9.7.1.13. Integer Arithmetic Instructions: max](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#integer-arithmetic-instructions-max)
- [9.7.1.14. Integer Arithmetic Instructions: popc](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#integer-arithmetic-instructions-popc)
- [9.7.1.15. Integer Arithmetic Instructions: clz](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#integer-arithmetic-instructions-clz)
- [9.7.1.16. Integer Arithmetic Instructions: bfind](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#integer-arithmetic-instructions-bfind)
- [9.7.1.17. Integer Arithmetic Instructions: fns](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#integer-arithmetic-instructions-fns)
- [9.7.1.18. Integer Arithmetic Instructions: brev](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#integer-arithmetic-instructions-brev)
- [9.7.1.19. Integer Arithmetic Instructions: bfe](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#integer-arithmetic-instructions-bfe)
- [9.7.1.20. Integer Arithmetic Instructions: bfi](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#integer-arithmetic-instructions-bfi)
- [9.7.1.21. Integer Arithmetic Instructions: szext](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#integer-arithmetic-instructions-szext)
- [9.7.1.22. Integer Arithmetic Instructions: bmsk](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#integer-arithmetic-instructions-bmsk)
- [9.7.1.23. Integer Arithmetic Instructions: dp4a](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#integer-arithmetic-instructions-dp4a)
- [9.7.1.24. Integer Arithmetic Instructions: dp2a](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#integer-arithmetic-instructions-dp2a)
- [9.7.2. Extended-Precision Integer Arithmetic Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#extended-precision-integer-arithmetic-instructions)
- [9.7.2.1. Extended-Precision Arithmetic Instructions: add.cc](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#extended-precision-arithmetic-instructions-add-cc)
- [9.7.2.2. Extended-Precision Arithmetic Instructions: addc](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#extended-precision-arithmetic-instructions-addc)
- [9.7.2.3. Extended-Precision Arithmetic Instructions: sub.cc](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#extended-precision-arithmetic-instructions-sub-cc)
- [9.7.2.4. Extended-Precision Arithmetic Instructions: subc](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#extended-precision-arithmetic-instructions-subc)
- [9.7.2.5. Extended-Precision Arithmetic Instructions: mad.cc](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#extended-precision-arithmetic-instructions-mad-cc)
- [9.7.2.6. Extended-Precision Arithmetic Instructions: madc](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#extended-precision-arithmetic-instructions-madc)
- [9.7.3. Floating-Point Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#floating-point-instructions)
- [9.7.3.1. Floating Point Instructions: testp](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#floating-point-instructions-testp)
- [9.7.3.2. Floating Point Instructions: copysign](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#floating-point-instructions-copysign)
- [9.7.3.3. Floating Point Instructions: add](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#floating-point-instructions-add)
- [9.7.3.4. Floating Point Instructions: sub](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#floating-point-instructions-sub)
- [9.7.3.5. Floating Point Instructions: mul](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#floating-point-instructions-mul)
- [9.7.3.6. Floating Point Instructions: fma](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#floating-point-instructions-fma)
- [9.7.3.7. Floating Point Instructions: mad](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#floating-point-instructions-mad)
- [9.7.3.8. Floating Point Instructions: div](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#floating-point-instructions-div)
- [9.7.3.9. Floating Point Instructions: abs](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#floating-point-instructions-abs)
- [9.7.3.10. Floating Point Instructions: neg](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#floating-point-instructions-neg)
- [9.7.3.11. Floating Point Instructions: min](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#floating-point-instructions-min)
- [9.7.3.12. Floating Point Instructions: max](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#floating-point-instructions-max)
- [9.7.3.13. Floating Point Instructions: rcp](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#floating-point-instructions-rcp)
- [9.7.3.14. Floating Point Instructions: rcp.approx.ftz.f64](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#floating-point-instructions-rcp-approx-ftz-f64)
- [9.7.3.15. Floating Point Instructions: sqrt](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#floating-point-instructions-sqrt)
- [9.7.3.16. Floating Point Instructions: rsqrt](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#floating-point-instructions-rsqrt)
- [9.7.3.17. Floating Point Instructions: rsqrt.approx.ftz.f64](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#floating-point-instructions-rsqrt-approx-ftz-f64)
- [9.7.3.18. Floating Point Instructions: sin](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#floating-point-instructions-sin)
- [9.7.3.19. Floating Point Instructions: cos](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#floating-point-instructions-cos)
- [9.7.3.20. Floating Point Instructions: lg2](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#floating-point-instructions-lg2)
- [9.7.3.21. Floating Point Instructions: ex2](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#floating-point-instructions-ex2)
- [9.7.3.22. Floating Point Instructions: tanh](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#floating-point-instructions-tanh)
- [9.7.4. Half Precision Floating-Point Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#half-precision-floating-point-instructions)
- [9.7.4.1. Half Precision Floating Point Instructions: add](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#half-precision-floating-point-instructions-add)
- [9.7.4.2. Half Precision Floating Point Instructions: sub](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#half-precision-floating-point-instructions-sub)
- [9.7.4.3. Half Precision Floating Point Instructions: mul](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#half-precision-floating-point-instructions-mul)
- [9.7.4.4. Half Precision Floating Point Instructions: fma](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#half-precision-floating-point-instructions-fma)
- [9.7.4.5. Half Precision Floating Point Instructions: neg](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#half-precision-floating-point-instructions-neg)
- [9.7.4.6. Half Precision Floating Point Instructions: abs](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#half-precision-floating-point-instructions-abs)
- [9.7.4.7. Half Precision Floating Point Instructions: min](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#half-precision-floating-point-instructions-min)
- [9.7.4.8. Half Precision Floating Point Instructions: max](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#half-precision-floating-point-instructions-max)
- [9.7.4.9. Half Precision Floating Point Instructions: tanh](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#half-precision-floating-point-instructions-tanh)
- [9.7.4.10. Half Precision Floating Point Instructions: ex2](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#half-precision-floating-point-instructions-ex2)
- [9.7.5. Mixed Precision Floating-Point Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#mixed-precision-floating-point-instructions)
- [9.7.5.1. Mixed Precision Floating Point Instructions: add](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#mixed-precision-floating-point-instructions-add)
- [9.7.5.2. Mixed Precision Floating Point Instructions: sub](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#mixed-precision-floating-point-instructions-sub)
- [9.7.5.3. Mixed Precision Floating Point Instructions: fma](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#mixed-precision-floating-point-instructions-fma)
- [9.7.6. Comparison and Selection Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#comparison-and-selection-instructions)
- [9.7.6.1. Comparison and Selection Instructions: set](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#comparison-and-selection-instructions-set)
- [9.7.6.2. Comparison and Selection Instructions: setp](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#comparison-and-selection-instructions-setp)
- [9.7.6.3. Comparison and Selection Instructions: selp](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#comparison-and-selection-instructions-selp)
- [9.7.6.4. Comparison and Selection Instructions: slct](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#comparison-and-selection-instructions-slct)
- [9.7.7. Half Precision Comparison Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#half-precision-comparison-instructions)
- [9.7.7.1. Half Precision Comparison Instructions: set](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#half-precision-comparison-instructions-set)
- [9.7.7.2. Half Precision Comparison Instructions: setp](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#half-precision-comparison-instructions-setp)
- [9.7.8. Logic and Shift Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#logic-and-shift-instructions)
- [9.7.8.1. Logic and Shift Instructions: and](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#logic-and-shift-instructions-and)
- [9.7.8.2. Logic and Shift Instructions: or](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#logic-and-shift-instructions-or)
- [9.7.8.3. Logic and Shift Instructions: xor](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#logic-and-shift-instructions-xor)
- [9.7.8.4. Logic and Shift Instructions: not](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#logic-and-shift-instructions-not)
- [9.7.8.5. Logic and Shift Instructions: cnot](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#logic-and-shift-instructions-cnot)
- [9.7.8.6. Logic and Shift Instructions: lop3](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#logic-and-shift-instructions-lop3)
- [9.7.8.7. Logic and Shift Instructions: shf](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#logic-and-shift-instructions-shf)
- [9.7.8.8. Logic and Shift Instructions: shl](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#logic-and-shift-instructions-shl)
- [9.7.8.9. Logic and Shift Instructions: shr](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#logic-and-shift-instructions-shr)
- [9.7.9. Data Movement and Conversion Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions)
- [9.7.9.1. Cache Operators](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#cache-operators)
- [9.7.9.2. Cache Eviction Priority Hints](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#cache-eviction-priority-hints)
- [9.7.9.3. Data Movement and Conversion Instructions: mov](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-mov)
- [9.7.9.4. Data Movement and Conversion Instructions: mov](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-mov-2)
- [9.7.9.5. Data Movement and Conversion Instructions: shfl (deprecated)](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-shfl)
- [9.7.9.6. Data Movement and Conversion Instructions: shfl.sync](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-shfl-sync)
- [9.7.9.7. Data Movement and Conversion Instructions: prmt](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-prmt)
- [9.7.9.8. Data Movement and Conversion Instructions: ld](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-ld)
- [9.7.9.9. Data Movement and Conversion Instructions: ld.global.nc](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-ld-global-nc)
- [9.7.9.10. Data Movement and Conversion Instructions: ldu](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-ldu)
- [9.7.9.11. Data Movement and Conversion Instructions: st](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-st)
- [9.7.9.12. Data Movement and Conversion Instructions: st.async](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-st-async)
- [9.7.9.13. Data Movement and Conversion Instructions: st.bulk](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-st-bulk)
- [9.7.9.14. Data Movement and Conversion Instructions: multimem.ld_reduce, multimem.st, multimem.red](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-multimem)
- [9.7.9.15. Data Movement and Conversion Instructions: prefetch, prefetchu](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-prefetch-prefetchu)
- [9.7.9.16. Data Movement and Conversion Instructions: applypriority](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-applypriority)
- [9.7.9.17. Data Movement and Conversion Instructions: discard](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-discard)
- [9.7.9.18. Data Movement and Conversion Instructions: createpolicy](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-createpolicy)
- [9.7.9.19. Data Movement and Conversion Instructions: isspacep](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-isspacep)
- [9.7.9.20. Data Movement and Conversion Instructions: cvta](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-cvta)
- [9.7.9.21. Data Movement and Conversion Instructions: cvt](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-cvt)
- [9.7.9.22. Data Movement and Conversion Instructions: cvt.pack](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-cvt-pack)
- [9.7.9.23. Data Movement and Conversion Instructions: mapa](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-mapa)
- [9.7.9.24. Data Movement and Conversion Instructions: getctarank](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-getctarank)
- [9.7.9.25. Data Movement and Conversion Instructions: Asynchronous copy](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-asynchronous-copy)
- [9.7.9.25.1. Completion Mechanisms for Asynchronous Copy Operations](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-asynchronous-copy-completion-mechanisms)
- [9.7.9.25.1.1. Async-group mechanism](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-asynchronous-copy-completion-mechanisms-async-group)
- [9.7.9.25.1.2. Mbarrier-based mechanism](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-asynchronous-copy-completion-mechanisms-mbarrier)
- [9.7.9.25.2. Async Proxy](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#async-proxy)
- [9.7.9.25.3. Data Movement and Conversion Instructions: Non-bulk copy](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-non-bulk-copy)
- [9.7.9.25.3.1. Data Movement and Conversion Instructions: cp.async](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-cp-async)
- [9.7.9.25.3.2. Data Movement and Conversion Instructions: cp.async.commit_group](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-cp-async-commit-group)
- [9.7.9.25.3.3. Data Movement and Conversion Instructions: cp.async.wait_group / cp.async.wait_all](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-cp-async-wait-group)
- [9.7.9.25.4. Data Movement and Conversion Instructions: Bulk copy](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-bulk-copy)
- [9.7.9.25.4.1. Data Movement and Conversion Instructions: cp.async.bulk](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-cp-async-bulk)
- [9.7.9.25.4.2. Data Movement and Conversion Instructions: cp.reduce.async.bulk](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-cp-reduce-async-bulk)
- [9.7.9.25.4.3. Data Movement and Conversion Instructions: cp.async.bulk.prefetch](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-cp-async-bulk-prefetch)
- [9.7.9.26. Data Movement and Conversion Instructions: multimem.cp.async.bulk](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-multimem-cp-async-bulk)
- [9.7.9.27. Data Movement and Conversion Instructions: multimem.cp.reduce.async.bulk](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-multimem-cp-reduce-async-bulk)
- [9.7.9.27.1. Data Movement and Conversion Instructions: Tensor copy](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-tensor-copy)
- [9.7.9.27.1.1. Restriction on Tensor Copy instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-tensor-copy-restrictions)
- [9.7.9.27.1.2. Data Movement and Conversion Instructions: cp.async.bulk.tensor](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-cp-async-bulk-tensor)
- [9.7.9.27.1.3. Data Movement and Conversion Instructions: cp.reduce.async.bulk.tensor](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-cp-reduce-async-bulk-tensor)
- [9.7.9.27.1.4. Data Movement and Conversion Instructions: cp.async.bulk.prefetch.tensor](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-cp-async-bulk-prefetch-tensor)
- [9.7.9.27.2. Data Movement and Conversion Instructions: Bulk and Tensor copy completion instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-bulk-tensor-copy-completion)
- [9.7.9.27.2.1. Data Movement and Conversion Instructions: cp.async.bulk.commit_group](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-cp-async-bulk-commit-group)
- [9.7.9.27.2.2. Data Movement and Conversion Instructions: cp.async.bulk.wait_group](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-cp-async-bulk-wait-group)
- [9.7.9.28. Data Movement and Conversion Instructions: tensormap.replace](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#data-movement-and-conversion-instructions-tensormap-replace)
- [9.7.10. Texture Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#texture-instructions)
- [9.7.10.1. Texturing Modes](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#texturing-modes)
- [9.7.10.2. Mipmaps](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#mipmaps)
- [9.7.10.3. Texture Instructions: tex](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#texture-instructions-tex)
- [9.7.10.4. Texture Instructions: tld4](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#texture-instructions-tld4)
- [9.7.10.5. Texture Instructions: txq](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#texture-instructions-txq)
- [9.7.10.6. Texture Instructions: istypep](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#texture-instructions-istypep)
- [9.7.11. Surface Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#surface-instructions)
- [9.7.11.1. Surface Instructions: suld](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#surface-instructions-suld)
- [9.7.11.2. Surface Instructions: sust](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#surface-instructions-sust)
- [9.7.11.3. Surface Instructions: sured](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#surface-instructions-sured)
- [9.7.11.4. Surface Instructions: suq](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#surface-instructions-suq)
- [9.7.12. Control Flow Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#control-flow-instructions)
- [9.7.12.1. Control Flow Instructions: {}](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#control-flow-instructions-curly-braces)
- [9.7.12.2. Control Flow Instructions: @](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#control-flow-instructions-at)
- [9.7.12.3. Control Flow Instructions: bra](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#control-flow-instructions-bra)
- [9.7.12.4. Control Flow Instructions: brx.idx](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#control-flow-instructions-brx-idx)
- [9.7.12.5. Control Flow Instructions: call](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#control-flow-instructions-call)
- [9.7.12.6. Control Flow Instructions: ret](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#control-flow-instructions-ret)
- [9.7.12.7. Control Flow Instructions: exit](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#control-flow-instructions-exit)
- [9.7.13. Parallel Synchronization and Communication Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions)
- [9.7.13.1. Parallel Synchronization and Communication Instructions: bar, barrier](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-bar)
- [9.7.13.2. Parallel Synchronization and Communication Instructions: bar.warp.sync](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-bar-warp-sync)
- [9.7.13.3. Parallel Synchronization and Communication Instructions: barrier.cluster](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-barrier-cluster)
- [9.7.13.4. Parallel Synchronization and Communication Instructions: membar / fence](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-membar)
- [9.7.13.5. Parallel Synchronization and Communication Instructions: atom](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-atom)
- [9.7.13.6. Parallel Synchronization and Communication Instructions: red](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-red)
- [9.7.13.7. Parallel Synchronization and Communication Instructions: red.async](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-red-async)
- [9.7.13.8. Parallel Synchronization and Communication Instructions: vote (deprecated)](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-vote)
- [9.7.13.9. Parallel Synchronization and Communication Instructions: vote.sync](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-vote-sync)
- [9.7.13.10. Parallel Synchronization and Communication Instructions: match.sync](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-match-sync)
- [9.7.13.11. Parallel Synchronization and Communication Instructions: activemask](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-activemask)
- [9.7.13.12. Parallel Synchronization and Communication Instructions: redux.sync](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-redux-sync)
- [9.7.13.13. Parallel Synchronization and Communication Instructions: griddepcontrol](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-griddepcontrol)
- [9.7.13.14. Parallel Synchronization and Communication Instructions: elect.sync](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-elect-sync)
- [9.7.13.15. Parallel Synchronization and Communication Instructions: mbarrier](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-mbarrier)
- [9.7.13.15.1. Size and alignment of mbarrier object](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-mbarrier-size-alignment)
- [9.7.13.15.2. Contents of the mbarrier object](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-mbarrier-contents)
- [9.7.13.15.3. Lifecycle of the mbarrier object](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-mbarrier-lifecycle)
- [9.7.13.15.4. Phase of the mbarrier object](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-mbarrier-phase)
- [9.7.13.15.5. Tracking asynchronous operations by the mbarrier object](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-mbarrier-tracking-async-operations)
- [9.7.13.15.5.1. expect-tx operation](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-mbarrier-expect-tx-operation)
- [9.7.13.15.5.2. complete-tx operation](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-mbarrier-complete-tx-operation)
- [9.7.13.15.6. Phase Completion of the mbarrier object](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-mbarrier-phase-completion)
- [9.7.13.15.7. Arrive-on operation on mbarrier object](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-mbarrier-arrive-on)
- [9.7.13.15.8. mbarrier support with shared memory](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-mbarrier-smem)
- [9.7.13.15.9. Parallel Synchronization and Communication Instructions: mbarrier.init](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-mbarrier-init)
- [9.7.13.15.10. Parallel Synchronization and Communication Instructions: mbarrier.inval](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-mbarrier-inval)
- [9.7.13.15.11. Parallel Synchronization and Communication Instructions: mbarrier.expect_tx](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-mbarrier-expect-tx)
- [9.7.13.15.12. Parallel Synchronization and Communication Instructions: mbarrier.complete_tx](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-mbarrier-complete-tx)
- [9.7.13.15.13. Parallel Synchronization and Communication Instructions: mbarrier.arrive](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-mbarrier-arrive)
- [9.7.13.15.14. Parallel Synchronization and Communication Instructions: mbarrier.arrive_drop](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-mbarrier-arrive-drop)
- [9.7.13.15.15. Parallel Synchronization and Communication Instructions: cp.async.mbarrier.arrive](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-cp-async-mbarrier-arrive)
- [9.7.13.15.16. Parallel Synchronization and Communication Instructions: mbarrier.test_wait / mbarrier.try_wait](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-mbarrier-test-wait-try-wait)
- [9.7.13.15.17. Parallel Synchronization and Communication Instructions: mbarrier.pending_count](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-mbarrier-pending-count)
- [9.7.13.16. Parallel Synchronization and Communication Instructions: tensormap.cp_fenceproxy](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-tensormap-cp-fenceproxy)
- [9.7.13.17. Parallel Synchronization and Communication Instructions: clusterlaunchcontrol.try_cancel](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-clusterlaunchcontrol-try-cancel)
- [9.7.13.18. Parallel Synchronization and Communication Instructions: clusterlaunchcontrol.query_cancel](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#parallel-synchronization-and-communication-instructions-clusterlaunchcontrol-query-cancel)
- [9.7.14. Warp Level Matrix Multiply-Accumulate Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-instructions)
- [9.7.14.1. Matrix Shape](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-shape)
- [9.7.14.2. Matrix Data-types](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-data-types)
- [9.7.14.3. Block Scaling for mma.sync](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-block-scaling)
- [9.7.14.4. Matrix multiply-accumulate operation using wmma instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-instructions-wmma)
- [9.7.14.4.1. Matrix Fragments for WMMA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-fragment)
- [9.7.14.4.2. Matrix Storage for WMMA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-storage)
- [9.7.14.4.3. Warp-level Matrix Load Instruction: wmma.load](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-instructions-wmma-ld)
- [9.7.14.4.4. Warp-level Matrix Store Instruction: wmma.store](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-instructions-wmma-st)
- [9.7.14.4.5. Warp-level Matrix Multiply-and-Accumulate Instruction: wmma.mma](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-instructions-wmma-mma)
- [9.7.14.5. Matrix multiply-accumulate operation using mma instruction](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-instructions-for-mma)
- [9.7.14.5.1. Matrix Fragments for mma.m8n8k4 with .f16 floating point type](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-fragment-mma-884-f16)
- [9.7.14.5.2. Matrix Fragments for mma.m8n8k4 with .f64 floating point type](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-fragment-mma-884-f64)
- [9.7.14.5.3. Matrix Fragments for mma.m8n8k16](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-fragment-mma-8816)
- [9.7.14.5.4. Matrix Fragments for mma.m8n8k32](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-fragment-mma-8832)
- [9.7.14.5.5. Matrix Fragments for mma.m8n8k128](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-fragment-mma-88128)
- [9.7.14.5.6. Matrix Fragments for mma.m16n8k4](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-fragment-mma-1684)
- [9.7.14.5.7. Matrix Fragments for mma.m16n8k8](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-fragment-mma-1688)
- [9.7.14.5.8. Matrix Fragments for mma.m16n8k16 with floating point type](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-fragment-mma-16816-float)
- [9.7.14.5.9. Matrix Fragments for mma.m16n8k16 with integer type](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-fragment-mma-16816-i8-f8)
- [9.7.14.5.10. Matrix Fragments for mma.m16n8k32](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-fragment-mma-16832)
- [9.7.14.5.11. Matrix Fragments for mma.m16n8k64](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-fragment-mma-16864)
- [9.7.14.5.12. Matrix Fragments for mma.m16n8k128](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-fragment-mma-168128)
- [9.7.14.5.13. Matrix Fragments for mma.m16n8k256](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-fragment-mma-168256)
- [9.7.14.5.14. Multiply-and-Accumulate Instruction: mma](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-instructions-mma)
- [9.7.14.5.15. Warp-level matrix load instruction: ldmatrix](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-instructions-ldmatrix)
- [9.7.14.5.16. Warp-level matrix store instruction: stmatrix](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-instructions-stmatrix)
- [9.7.14.5.17. Warp-level matrix transpose instruction: movmatrix](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-instructions-movmatrix)
- [9.7.14.6. Matrix multiply-accumulate operation using mma.sp instruction with sparse matrix A](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-instructions-for-sparse-mma)
- [9.7.14.6.1. Sparse matrix storage](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-sparse-matrix-storage)
- [9.7.14.6.2. Matrix fragments for multiply-accumulate operation with sparse matrix A](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-fragments-for-sparse-mma)
- [9.7.14.6.2.1. Matrix Fragments for sparse mma.m16n8k16 with .f16 and .bf16 types](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-fragment-sparse-mma-16816-f16bf16)
- [9.7.14.6.2.2. Matrix Fragments for sparse mma.m16n8k32 with .f16 and .bf16 types](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-fragment-sparse-mma-16832-f16bf16)
- [9.7.14.6.2.3. Matrix Fragments for sparse mma.m16n8k16 with .tf32 floating point type](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-fragment-sparse-mma-16816-tf32)
- [9.7.14.6.2.4. Matrix Fragments for sparse mma.m16n8k8 with .tf32 floating point type](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-fragment-sparse-mma-1688-tf32)
- [9.7.14.6.2.5. Matrix Fragments for sparse mma.m16n8k32 with .u8 / .s8 integer type](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-fragment-sparse-mma-16832-u8s8)
- [9.7.14.6.2.6. Matrix Fragments for sparse mma.m16n8k64 with .u8 / .s8 / .e4m3 / .e5m2 type](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-fragment-sparse-mma-16864-u8s8-fp8)
- [9.7.14.6.2.7. Matrix Fragments for sparse mma.m16n8k64 with .u4 / .s4 integer type](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-fragment-sparse-mma-16864-u4s4)
- [9.7.14.6.2.8. Matrix Fragments for sparse mma.m16n8k128 with .u4 / .s4 integer type](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-fragment-sparse-mma-168128-u4s4)
- [9.7.14.6.3. Multiply-and-Accumulate Instruction: mma.sp / mma.sp::ordered_metadata](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#warp-level-matrix-instructions-sparse-mma)
- [9.7.15. Asynchronous Warpgroup Level Matrix Multiply-Accumulate Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-instructions)
- [9.7.15.1. Warpgroup](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-instructions-warpgroup)
- [9.7.15.2. Matrix Shape](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-shape)
- [9.7.15.3. Matrix Data-types](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-data-types)
- [9.7.15.4. Async Proxy](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-async-proxy)
- [9.7.15.5. Asynchronous Warpgroup Level Matrix Multiply-Accumulate Operation using wgmma.mma_async instruction](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#asynchronous-warpgroup-level-matrix-operation-wgmma-mma-async)
- [9.7.15.5.1. Register Fragments and Shared Memory Matrix Layouts](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-fragment)
- [9.7.15.5.1.1. Register Fragments](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-register-fragment)
- [9.7.15.5.1.1.1. Matrix Fragments for wgmma.mma_async.m64nNk16](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#asynchronous-warpgroup-level-matrix-register-fragment-wgmma-64n16)
- [9.7.15.5.1.1.2. Matrix Fragments for wgmma.mma_async.m64nNk8](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#asynchronous-warpgroup-level-matrix-register-fragment-wgmma-64n8)
- [9.7.15.5.1.1.3. Matrix Fragments for wgmma.mma_async.m64nNk32](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#asynchronous-warpgroup-level-matrix-register-fragment-wgmma-64n32)
- [9.7.15.5.1.1.4. Matrix Fragments for wgmma.mma_async.m64nNk256](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#asynchronous-warpgroup-level-matrix-register-fragment-wgmma-64n256)
- [9.7.15.5.1.2. Shared Memory Matrix Layout](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-shared-memory-layout)
- [9.7.15.5.1.2.1. Major-ness supported by Strides](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-majorness-supported-by-strides)
- [9.7.15.5.1.2.1.1. Leading Dimension Byte Offset](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-leading-dimension-byte-offset)
- [9.7.15.5.1.2.1.2. Stride Dimension Byte Offset](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-stride-dimension-byte-offset)
- [9.7.15.5.1.2.1.3. Canonical Layouts](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-canonical-layouts)
- [9.7.15.5.1.2.2. Matrix Descriptor Format](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-shared-memory-layout-matrix-descriptor)
- [9.7.15.5.2. Asynchronous Multiply-and-Accumulate Instruction: wgmma.mma_async](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#asynchronous-warpgroup-level-matrix-instructions-wgmma-mma)
- [9.7.15.6. Asynchronous Warpgroup Level Multiply-and-Accumulate Operation using wgmma.mma_async.sp instruction](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#asynchronous-warpgroup-level-matrix-instructions-for-sparse-wgmma)
- [9.7.15.6.1. Sparse matrix storage](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-sparse-matrix-storage)
- [9.7.15.6.2. Matrix fragments for warpgroup-level multiply-accumulate operation with sparse matrix A](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-fragments-for-sparse-wgmma)
- [9.7.15.6.2.1. Matrix Fragments for sparse wgmma.mma_async.m64nNk32](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#asynchronous-warpgroup-level-matrix-fragment-sparse-wgmma-64n32)
- [9.7.15.6.2.2. Matrix Fragments for sparse wgmma.mma_async.m64nNk16](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-fragment-sparse-wgmma-64n16)
- [9.7.15.6.2.3. Matrix Fragments for sparse wgmma.mma_async.m64nNk64](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#asynchronous-warpgroup-level-matrix-fragment-sparse-wgmma-64n64)
- [9.7.15.6.3. Asynchronous Multiply-and-Accumulate Instruction: wgmma.mma_async.sp](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#asynchronous-warpgroup-level-matrix-instructions-wgmma-mma-sp)
- [9.7.15.7. Asynchronous wgmma Proxy Operations](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#asynchronous-wgmma-proxy-operations)
- [9.7.15.7.1. Asynchronous Multiply-and-Accumulate Instruction: wgmma.fence](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#asynchronous-warpgroup-level-matrix-instructions-wgmma-fence)
- [9.7.15.7.2. Asynchronous Multiply-and-Accumulate Instruction: wgmma.commit_group](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#asynchronous-warpgroup-level-matrix-instructions-wgmma-commit-group)
- [9.7.15.7.3. Asynchronous Multiply-and-Accumulate Instruction: wgmma.wait_group](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#asynchronous-warpgroup-level-matrix-instructions-wgmma-wait-group)
- [9.7.16. TensorCore 5th Generation Family Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tensorcore-5th-generation-instructions)
- [9.7.16.1. Tensor Memory](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tensor-memory)
- [9.7.16.1.1. Tensor Memory Addressing](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tensor-memory-addressing)
- [9.7.16.1.2. Tensor Memory Allocation](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tensor-memory-allocation)
- [9.7.16.2. Matrix and Data Movement Shape](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-matrix-data-movement-shape)
- [9.7.16.2.1. Matrix Shape](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-matrix-shape)
- [9.7.16.2.1.1. Target ISA Note](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-matrix-shape-target-isa-note)
- [9.7.16.2.2. Specifying Matrix Shape](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-specify-matrix-shape)
- [9.7.16.2.3. Data Movement Shape](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-data-movement-shape)
- [9.7.16.2.3.1. Memory Layout](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-memory-layout)
- [9.7.16.2.3.1.1. Matrix fragments for shape .32x32b](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-matrix-fragments-shape-3232b)
- [9.7.16.2.3.1.2. Matrix fragments for shape .16x64b](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-matrix-fragments-shape-6464b)
- [9.7.16.2.3.1.3. Matrix fragments for shape .16x128b](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-matrix-fragments-shape-16128b)
- [9.7.16.2.3.1.4. Matrix fragments for shape .16x256b](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-matrix-fragments-shape-16256b)
- [9.7.16.2.3.1.5. Matrix fragments for shape .16x32bx2](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-matrix-fragments-shape-1632b2)
- [9.7.16.3. Major-ness supported by Strides](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-majorness-supported-by-strides)
- [9.7.16.3.1. Leading Dimension Stride: relative offset or absolute address](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-leading-dimension-byte-offset)
- [9.7.16.3.1.1. Relative offset mode](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-leading-dimension-byte-offset-relative-offset)
- [9.7.16.3.1.2. Absolute address mode for K dimension being 48B](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-leading-dimension-byte-offset-absolute-address)
- [9.7.16.3.1.2.1. Restrictions on the Leading Dimension Absolute Address Stride](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-leading-dimension-byte-offset-absolute-address-restriction)
- [9.7.16.3.2. Stride Dimension Byte Offset](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-stride-dimension-byte-offset)
- [9.7.16.3.3. Canonical Layouts](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-canonical-layouts)
- [9.7.16.4. Matrix Descriptors](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-matrix-descriptors)
- [9.7.16.4.1. Shared memory descriptor](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-shared-memory-descriptor)
- [9.7.16.4.1.1. Target ISA Note](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-shared-memory-descriptor-target-isa-note)
- [9.7.16.4.2. Instruction descriptor](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-instruction-descriptor)
- [9.7.16.4.3. Zero-Column Mask Descriptor](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-zero-column-mask-descriptor)
- [9.7.16.5. Issue Granularity](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-issue-granularity)
- [9.7.16.5.1. CTA Pair](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-cta-pair)
- [9.7.16.5.2. Peer CTA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-peer-cta)
- [9.7.16.6. Memory Consistency Model for 5th generation of TensorCore operations](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-memory-consistency-model)
- [9.7.16.6.1. Asynchronous Operations](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-memory-consistency-model-async-operations)
- [9.7.16.6.2. Pipelined tcgen05 Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-memory-consistency-model-pipelined-instructions)
- [9.7.16.6.2.1. Implicitly pipelined tcgen05 Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-memory-consistency-model-pipelined-instructions-implicit)
- [9.7.16.6.2.1.1. mbarrier based completion mechanism](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-memory-consistency-model-mbarrier-completion)
- [9.7.16.6.2.1.2. tcgen05.wait instruction based completion mechanism](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-memory-consistency-model-wait-completion)
- [9.7.16.6.3. Specialized Inter-thread Synchronization for tcgen05 instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-memory-consistency-model-inter-thread-sync)
- [9.7.16.6.4. Canonical synchronization patterns](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-memory-consistency-model-canonical-sync-patterns)
- [9.7.16.6.4.1. Pipelined instructions, same thread](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-memory-consistency-model-canonical-sync-patterns-pipelined-same-thread)
- [9.7.16.6.4.2. Non-pipelined instructions, same thread](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-memory-consistency-model-canonical-sync-patterns-non-pipelined-same-thread)
- [9.7.16.6.4.3. Pipelined instructions, different thread](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-memory-consistency-model-canonical-sync-patterns-pipelined-diff-thread)
- [9.7.16.6.4.4. Non-pipelined instructions, different thread](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-memory-consistency-model-canonical-sync-patterns-non-pipelined-diff-thread)
- [9.7.16.6.4.5. Register dependencies, same thread](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-memory-consistency-model-canonical-sync-patterns-reg-dependency-same-thread)
- [9.7.16.6.5. Shared Memory Accesses](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-memory-consistency-model-smem-access)
- [9.7.16.7. Tensor Memory Allocation and Management Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-memory-alloc-manage-instructions)
- [9.7.16.7.1. Tensorcore 5th Generation Instructions: tcgen05.alloc, tcgen05.dealloc, tcgen05.relinquish_alloc_permit](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-instructions-tcgen05-alloc-dealloc-relinquish-alloc-permit)
- [9.7.16.8. Tensor Memory and Register Load/Store Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-tensor-memory-ld-st)
- [9.7.16.8.1. Access restrictions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-tensor-memory-ld-st-access-restrictions)
- [9.7.16.8.2. Packing and Unpacking](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-tensor-memory-ld-st-packing-unpacking)
- [9.7.16.8.3. Tensorcore 5th Generation Instructions: tcgen05.ld](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-instructions-tcgen05-ld)
- [9.7.16.8.4. Tensorcore 5th Generation Instructions: tcgen05.st](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-instructions-tcgen05-st)
- [9.7.16.8.5. Tensorcore 5th Generation Instructions: tcgen05.wait](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-instructions-tcgen05-wait)
- [9.7.16.9. Tensor Memory Data Movement Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-data-movement-instructions)
- [9.7.16.9.1. Optional Decompression](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-optional-decompression)
- [9.7.16.9.1.1. Decompression of 4-bit floating point to 8-bit type](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-optional-decompression-4bit-8bit)
- [9.7.16.9.1.2. Decompression of 6-bit floating point to 8-bit type](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-optional-decompression-6bit-8bit)
- [9.7.16.9.2. Tensorcore 5th Generation Instructions: tcgen05.cp](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-instructions-tcgen05-cp)
- [9.7.16.9.3. Tensorcore 5th Generation Instructions: tcgen05.shift](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-instructions-tcgen05-shift)
- [9.7.16.10. TensorCore 5th Generation Matrix Multiply and accumulate Operations](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-mma)
- [9.7.16.10.1. Transpose and Negate operations](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-transpose-and-negate-operations)
- [9.7.16.10.2. Matrix Layout Organization](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-matrix-layout-organization)
- [9.7.16.10.3. Valid Combinations of Type-Size, Major-ness and Swizzling](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-matrix-layout-organization-valid-comb-type-size-majorness-swizzle)
- [9.7.16.10.4. Packing formats of elements in Tensor and Shared memory](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-packing-formats)
- [9.7.16.10.4.1. Packing format for matrix D in Tensor Memory](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-packing-formats-mat-d)
- [9.7.16.10.4.2. Packing format for matrix A and B](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-packing-formats-mat-a-b)
- [9.7.16.10.4.3. Packing format used for matrix A by .kind::mxf8f6f4 in Tensor Memory](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-packing-formats-mxf8f6f4-tmem)
- [9.7.16.10.4.4. Packing format used for matrix A and B by .kind::mxf8f6f4 in Shared Memory](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-packing-formats-mxf8f6f4-smem)
- [9.7.16.10.4.5. Packing format used for matrix A by .kind::mxf4 and .kind::mxf4nvf4 in Tensor Memory](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-packing-formats-mxf4-tmem)
- [9.7.16.10.4.6. Packing format used for matrix A and B by .kind::mxf4 and .kind::mxf4nvf4 in Shared Memory](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-packing-formats-mxf4-smem)
- [9.7.16.10.5. Data Path Layout Organization](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-data-path-layout-organization)
- [9.7.16.10.5.1. Layout A (M = 256)](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-data-path-layout-a)
- [9.7.16.10.5.2. Layout B (M = 128 + cta-group::2 + Dense A matrix)](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-data-path-layout-b)
- [9.7.16.10.5.3. Layout C (M = 128 + cta-group::2 + Sparse A matrix)](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-data-path-layout-c)
- [9.7.16.10.5.4. Layout D (M = 128 + cta-group::1)](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-data-path-layout-d)
- [9.7.16.10.5.5. Layout E (M = 64 + .ws mode)](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-data-path-layout-e)
- [9.7.16.10.5.6. Layout F (M = 64 + non .ws mode)](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-data-path-layout-f)
- [9.7.16.10.5.7. Layout G (M = 32)](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-data-path-layout-g)
- [9.7.16.10.6. Shared Memory Layout and Swizzling](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-shared-memory-layout-swizzling)
- [9.7.16.10.7. Block Scaling for tcgen05.mma.sync](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-block-scaling)
- [9.7.16.10.7.1. Valid combinations of scale_vectorsize with types and MMA-Kind](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-mma-scale-valid-vec-size)
- [9.7.16.10.7.2. Scale Factor A ID](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-mma-scale-factor-a)
- [9.7.16.10.7.2.1. Layout of the Scale Factor A Matrix for scale_vec::1X/block32 with K=32/K=64](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-mma-scale-factor-a-layout-1x)
- [9.7.16.10.7.2.2. Layout of the Scale Factor A Matrix for scale_vec::2X/block32 with K=64/K=128](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-mma-scale-factor-a-layout-2x)
- [9.7.16.10.7.2.3. Layout of the Scale Factor A Matrix for scale_vec::4X/block16 with K=64/K=128](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-mma-scale-factor-a-layout-4x)
- [9.7.16.10.7.2.4. Layout of the Scale Factor A Matrix for block32 with K=96 (Semantically equivalent to scale_vec::3X)](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-mma-scale-factor-a-layout-block32-k96)
- [9.7.16.10.7.2.5. Layout of the Scale Factor A Matrix for block16 with K=96 (Semantically equivalent to scale_vec::6X)](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-mma-scale-factor-a-layout-block16-k96)
- [9.7.16.10.7.3. Scale Factor B ID](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-mma-scale-factor-b)
- [9.7.16.10.7.3.1. Layout of the Scale Factor B Matrix for scale_vec::1X/block32 with K=32/K=64](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-mma-scale-factor-b-layout-1x)
- [9.7.16.10.7.3.2. Layout of the Scale Factor B Matrix for scale_vec::2X/block32 with K=64/K=128](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-mma-scale-factor-b-layout-2x)
- [9.7.16.10.7.3.3. Layout of the Scale Factor B Matrix for scale_vec::4X/block16 with K=64/K=128](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-mma-scale-factor-b-layout-4x)
- [9.7.16.10.7.3.4. Layout of the Scale Factor B Matrix for block32 with K=96 (Semantically equivalent to scale_vec::3X)](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-mma-scale-factor-b-layout-block32-k96)
- [9.7.16.10.7.3.5. Layout of the Scale Factor B Matrix for block16 with K=96 (Semantically equivalent to scale_vec::6X)](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-mma-scale-factor-b-layout-block16-k96)
- [9.7.16.10.8. Sparse Matrices](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-sparse-matrices)
- [9.7.16.10.8.1. Sparse tcgen05.mma.sp with .kind::tf32](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-sparse-matrices-kind-tf32)
- [9.7.16.10.8.2. Sparse tcgen05.mma.sp with .kind::f16, .kind::f8f6f4, .kind::mxf8f6f4, .kind::i8](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-sparse-matrices-kind-f16-f8f8f4-mxf8f6f4)
- [9.7.16.10.8.3. Sparse tcgen05.mma.sp with .kind::mxf4 and .kind::mxf4nvf4](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-sparse-matrices-kind-mxf4)
- [9.7.16.10.8.4. Sparsity selector](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-sparse-matrices-sparsity-selector)
- [9.7.16.10.8.4.1. Layout of the Sparsity Metadata Matrix for M = 64 for .kind::f16](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-sparse-matrices-sparsity-selector-kind-f16-m64)
- [9.7.16.10.8.4.2. Layout of the Sparsity Metadata Matrix for M = 128 / M = 256 for .kind::f16](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-sparse-matrices-sparsity-selector-kind-f16-m128-256)
- [9.7.16.10.8.4.3. Layout of the Sparsity Metadata Matrix for M = 64 for .kind::tf32](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-sparse-matrices-sparsity-selector-kind-tf32-m64)
- [9.7.16.10.8.4.4. Layout of the Sparsity Metadata Matrix for M = 128 / M = 256 for .kind::tf32](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-sparse-matrices-sparsity-selector-kind-tf32-m128-256)
- [9.7.16.10.8.4.5. Layout of the Sparsity Metadata Matrix for M = 64 for .kind::f8f6f4, .kind::mxf8f6f4, .kind::i8, .kind::mxf4, .kind::mxf4nvf4](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-sparse-matrices-sparsity-selector-kind-f8f6f4-mxf8f6f4-m64)
- [9.7.16.10.8.4.6. Layout of the Sparsity Metadata Matrix for M = 128 / M = 256 for .kind::f8f6f4, .kind::mxf8f6f4, .kind::i8, .kind::mxf4, .kind::mxf4nvf4](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-sparse-matrices-sparsity-selector-kind-f8f6f4-mxf8f6f4-m128-256)
- [9.7.16.10.8.5. Alignment restriction](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-sparse-matrices-alignment-restriction)
- [9.7.16.10.9. TensorCore 5th Generation of MMA Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-mma-instructions)
- [9.7.16.10.9.1. TensorCore 5th Generation Instructions: tcgen05.mma](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-mma-instructions-mma)
- [9.7.16.10.9.2. TensorCore 5th Generation Instructions: tcgen05.mma.sp](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-mma-instructions-mma-sp)
- [9.7.16.10.9.3. TensorCore 5th Generation Instructions: tcgen05.mma.ws](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-mma-instructions-mma-ws)
- [9.7.16.10.9.4. TensorCore 5th Generation Instructions: tcgen05.mma.ws.sp](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-mma-instructions-mma-ws-sp)
- [9.7.16.11. TensorCore 5th Generation Specialized Synchronization Operations](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen05-special-sync-operations)
- [9.7.16.11.1. TensorCore 5th Generation Instructions: tcgen05.fence](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen05-special-sync-operations-fence)
- [9.7.16.12. TensorCore 5th Generation Async Synchronization Operations](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tcgen-async-sync-operations)
- [9.7.16.12.1. TensorCore 5th Generation Instructions: tcgen05.commit](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#tcgen-async-sync-operations-commit)
- [9.7.17. Stack Manipulation Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#stack-manipulation-instructions)
- [9.7.17.1. Stack Manipulation Instructions: stacksave](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#stack-manipulation-instructions-stacksave)
- [9.7.17.2. Stack Manipulation Instructions: stackrestore](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#stack-manipulation-instructions-stackrestore)
- [9.7.17.3. Stack Manipulation Instructions: alloca](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#stack-manipulation-instructions-alloca)
- [9.7.18. Video Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#video-instructions)
- [9.7.18.1. Scalar Video Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#scalar-video-instructions)
- [9.7.18.1.1. Scalar Video Instructions: vadd, vsub, vabsdiff, vmin, vmax](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#scalar-video-instructions-vadd-vsub-vabsdiff-vmin-vmax)
- [9.7.18.1.2. Scalar Video Instructions: vshl, vshr](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#scalar-video-instructions-vshl-vshr)
- [9.7.18.1.3. Scalar Video Instructions: vmad](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#scalar-video-instructions-vmad)
- [9.7.18.1.4. Scalar Video Instructions: vset](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#scalar-video-instructions-vset)
- [9.7.18.2. SIMD Video Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#simd-video-instructions)
- [9.7.18.2.1. SIMD Video Instructions: vadd2, vsub2, vavrg2, vabsdiff2, vmin2, vmax2](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#simd-video-instructions-vadd2-vsub2-vavrg2-vabsdiff2-vmin2-vmax2)
- [9.7.18.2.2. SIMD Video Instructions: vset2](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#simd-video-instructions-vset2)
- [9.7.18.2.3. SIMD Video Instructions: vadd4, vsub4, vavrg4, vabsdiff4, vmin4, vmax4](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#simd-video-instructions-vadd4-vsub4-vavrg4-vabsdiff4-vmin4-vmax4)
- [9.7.18.2.4. SIMD Video Instructions: vset4](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#simd-video-instructions-vset4)
- [9.7.19. Miscellaneous Instructions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#miscellaneous-instructions)
- [9.7.19.1. Miscellaneous Instructions: brkpt](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#miscellaneous-instructions-brkpt)
- [9.7.19.2. Miscellaneous Instructions: nanosleep](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#miscellaneous-instructions-nanosleep)
- [9.7.19.3. Miscellaneous Instructions: pmevent](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#miscellaneous-instructions-pmevent)
- [9.7.19.4. Miscellaneous Instructions: trap](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#miscellaneous-instructions-trap)
- [9.7.19.5. Miscellaneous Instructions: setmaxnreg](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#miscellaneous-instructions-setmaxnreg)
- [10. Special Registers](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers)
- [10.1. Special Registers: %tid](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#special-registers-tid)
- [10.2. Special Registers: %ntid](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#special-registers-ntid)
- [10.3. Special Registers: %laneid](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#special-registers-laneid)
- [10.4. Special Registers: %warpid](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#special-registers-warpid)
- [10.5. Special Registers: %nwarpid](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#special-registers-nwarpid)
- [10.6. Special Registers: %ctaid](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#special-registers-ctaid)
- [10.7. Special Registers: %nctaid](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#special-registers-nctaid)
- [10.8. Special Registers: %smid](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#special-registers-smid)
- [10.9. Special Registers: %nsmid](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#special-registers-nsmid)
- [10.10. Special Registers: %gridid](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#special-registers-gridid)
- [10.11. Special Registers: %is_explicit_cluster](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#special-registers-is-explicit-cluster)
- [10.12. Special Registers: %clusterid](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#special-registers-clusterid)
- [10.13. Special Registers: %nclusterid](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#special-registers-nclusterid)
- [10.14. Special Registers: %cluster_ctaid](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#special-registers-cluster-ctaid)
- [10.15. Special Registers: %cluster_nctaid](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#special-registers-cluster-nctaid)
- [10.16. Special Registers: %cluster_ctarank](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#special-registers-cluster-ctarank)
- [10.17. Special Registers: %cluster_nctarank](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#special-registers-cluster-nctarank)
- [10.18. Special Registers: %lanemask_eq](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#special-registers-lanemask-eq)
- [10.19. Special Registers: %lanemask_le](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#special-registers-lanemask-le)
- [10.20. Special Registers: %lanemask_lt](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#special-registers-lanemask-lt)
- [10.21. Special Registers: %lanemask_ge](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#special-registers-lanemask-ge)
- [10.22. Special Registers: %lanemask_gt](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#special-registers-lanemask-gt)
- [10.23. Special Registers: %clock, %clock_hi](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#special-registers-clock)
- [10.24. Special Registers: %clock64](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#special-registers-clock64)
- [10.25. Special Registers: %pm0 … %pm7](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#special-registers-pm0-pm7)
- [10.26. Special Registers: %pm0_64 … %pm7_64](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#special-registers-pm0-64-pm7-64)
- [10.27. Special Registers: %envreg<32>](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#special-registers-envreg-32)
- [10.28. Special Registers: %globaltimer, %globaltimer_lo, %globaltimer_hi](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#special-registers-globaltimer)
- [10.29. Special Registers: %reserved_smem_offset_begin, %reserved_smem_offset_end, %reserved_smem_offset_cap, %reserved_smem_offset_<2>](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#special-registers-reserved-smem)
- [10.30. Special Registers: %total_smem_size](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#special-registers-total-smem-size)
- [10.31. Special Registers: %aggr_smem_size](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#special-registers-aggr-smem-size)
- [10.32. Special Registers: %dynamic_smem_size](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#special-registers-dynamic-smem-size)
- [10.33. Special Registers: %current_graph_exec](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#special-registers-current-graph-exec)
- [11. Directives](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#directives)
- [11.1. PTX Module Directives](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#ptx-module-directives)
- [11.1.1. PTX Module Directives: .version](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#ptx-module-directives-version)
- [11.1.2. PTX Module Directives: .target](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#ptx-module-directives-target)
- [11.1.3. PTX Module Directives: .address_size](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#ptx-module-directives-address-size)
- [11.2. Specifying Kernel Entry Points and Functions](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#specifying-kernel-entry-points-and-functions)
- [11.2.1. Kernel and Function Directives: .entry](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#kernel-and-function-directives-entry)
- [11.2.2. Kernel and Function Directives: .func](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#kernel-and-function-directives-func)
- [11.2.3. Kernel and Function Directives: .alias](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#kernel-and-function-directives-alias)
- [11.3. Control Flow Directives](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#control-flow-directives)
- [11.3.1. Control Flow Directives: .branchtargets](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#control-flow-directives-branchtargets)
- [11.3.2. Control Flow Directives: .calltargets](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#control-flow-directives-calltargets)
- [11.3.3. Control Flow Directives: .callprototype](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#control-flow-directives-callprototype)
- [11.4. Performance-Tuning Directives](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#performance-tuning-directives)
- [11.4.1. Performance-Tuning Directives: .maxnreg](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#performance-tuning-directives-maxnreg)
- [11.4.2. Performance-Tuning Directives: .maxntid](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#performance-tuning-directives-maxntid)
- [11.4.3. Performance-Tuning Directives: .reqntid](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#performance-tuning-directives-reqntid)
- [11.4.4. Performance-Tuning Directives: .minnctapersm](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#performance-tuning-directives-minnctapersm)
- [11.4.5. Performance-Tuning Directives: .maxnctapersm (deprecated)](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#performance-tuning-directives-maxnctapersm)
- [11.4.6. Performance-Tuning Directives: .noreturn](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#performance-tuning-directives-noreturn)
- [11.4.7. Performance-Tuning Directives: .pragma](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#performance-tuning-directives-pragma)
- [11.4.8. Performance-Tuning Directives: .abi_preserve](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#performance-tuning-directives-abi-preserve)
- [11.4.9. Performance-Tuning Directives: .abi_preserve_control](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#performance-tuning-directives-abi-preserve-control)
- [11.5. Debugging Directives](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#debugging-directives)
- [11.5.1. Debugging Directives: @@dwarf](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#debugging-directives-atatdwarf)
- [11.5.2. Debugging Directives: .section](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#debugging-directives-section)
- [11.5.3. Debugging Directives: .file](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#debugging-directives-file)
- [11.5.4. Debugging Directives: .loc](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#debugging-directives-loc)
- [11.6. Linking Directives](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#linking-directives)
- [11.6.1. Linking Directives: .extern](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#linking-directives-extern)
- [11.6.2. Linking Directives: .visible](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#linking-directives-visible)
- [11.6.3. Linking Directives: .weak](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#linking-directives-weak)
- [11.6.4. Linking Directives: .common](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#linking-directives-common)
- [11.7. Cluster Dimension Directives](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#cluster-dimension-directives)
- [11.7.1. Cluster Dimension Directives: .reqnctapercluster](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#cluster-dimension-directives-reqnctapercluster)
- [11.7.2. Cluster Dimension Directives: .explicitcluster](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#cluster-dimension-directives-explicitcluster)
- [11.7.3. Cluster Dimension Directives: .maxclusterrank](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#cluster-dimension-directives-maxclusterrank)
- [11.8. Miscellaneous Directives](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#miscellaneous-directives)
- [11.8.1. Miscellaneous Directives: .blocksareclusters](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#miscellaneous-directives-blocksareclusters)
- [12. Descriptions of .pragma Strings](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#descriptions-pragma-strings)
- [12.1. Pragma Strings: "nounroll"](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#pragma-strings-nounroll)
- [12.2. Pragma Strings: "used_bytes_mask"](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#pragma-strings-used-bytes-mask)
- [12.3. Pragma Strings: "enable_smem_spilling"](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#pragma-strings-enable-smem-spilling)
- [12.4. Pragma Strings: "frequency"](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#pragma-strings-frequency)
- [13. Release Notes](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#release-notes)
- [13.1. Changes in PTX ISA Version 9.1](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-9-1)
- [13.2. Changes in PTX ISA Version 9.0](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-9-0)
- [13.3. Changes in PTX ISA Version 8.8](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-8-8)
- [13.4. Changes in PTX ISA Version 8.7](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-8-7)
- [13.5. Changes in PTX ISA Version 8.6](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-8-6)
- [13.6. Changes in PTX ISA Version 8.5](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-8-5)
- [13.7. Changes in PTX ISA Version 8.4](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-8-4)
- [13.8. Changes in PTX ISA Version 8.3](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-8-3)
- [13.9. Changes in PTX ISA Version 8.2](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-8-2)
- [13.10. Changes in PTX ISA Version 8.1](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-8-1)
- [13.11. Changes in PTX ISA Version 8.0](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-8-0)
- [13.12. Changes in PTX ISA Version 7.8](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-7-8)
- [13.13. Changes in PTX ISA Version 7.7](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-7-7)
- [13.14. Changes in PTX ISA Version 7.6](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-7-6)
- [13.15. Changes in PTX ISA Version 7.5](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-7-5)
- [13.16. Changes in PTX ISA Version 7.4](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-7-4)
- [13.17. Changes in PTX ISA Version 7.3](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-7-3)
- [13.18. Changes in PTX ISA Version 7.2](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-7-2)
- [13.19. Changes in PTX ISA Version 7.1](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-7-1)
- [13.20. Changes in PTX ISA Version 7.0](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-7-0)
- [13.21. Changes in PTX ISA Version 6.5](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-6-5)
- [13.22. Changes in PTX ISA Version 6.4](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-6-4)
- [13.23. Changes in PTX ISA Version 6.3](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-6-3)
- [13.24. Changes in PTX ISA Version 6.2](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-6-2)
- [13.25. Changes in PTX ISA Version 6.1](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-6-1)
- [13.26. Changes in PTX ISA Version 6.0](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-6-0)
- [13.27. Changes in PTX ISA Version 5.0](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-5-0)
- [13.28. Changes in PTX ISA Version 4.3](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-4-3)
- [13.29. Changes in PTX ISA Version 4.2](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-4-2)
- [13.30. Changes in PTX ISA Version 4.1](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-4-1)
- [13.31. Changes in PTX ISA Version 4.0](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-4-0)
- [13.32. Changes in PTX ISA Version 3.2](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-3-2)
- [13.33. Changes in PTX ISA Version 3.1](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-3-1)
- [13.34. Changes in PTX ISA Version 3.0](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-3-0)
- [13.35. Changes in PTX ISA Version 2.3](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-2-3)
- [13.36. Changes in PTX ISA Version 2.2](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-2-2)
- [13.37. Changes in PTX ISA Version 2.1](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-2-1)
- [13.38. Changes in PTX ISA Version 2.0](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#changes-in-ptx-isa-version-2-0)
- [14. Notices](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#notices)
- [14.1. Notice](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#notice)
- [14.2. OpenCL](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#opencl)
- [14.3. Trademarks](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html.md#trademarks)
- [Privacy Policy](https://www.nvidia.com/en-us/about-nvidia/privacy-policy.md/)
- [Manage My Privacy](https://www.nvidia.com/en-us/about-nvidia/privacy-center.md/)
- [Do Not Sell or Share My Data](https://www.nvidia.com/en-us/preferences/start.md/)
- [Terms of Service](https://www.nvidia.com/en-us/about-nvidia/terms-of-service.md/)
- [Accessibility](https://www.nvidia.com/en-us/about-nvidia/accessibility.md/)
- [Corporate Policies](https://www.nvidia.com/en-us/about-nvidia/company-policies.md/)
- [Product Security](https://www.nvidia.com/en-us/product-security.md/)
- [Contact](https://www.nvidia.com/en-us/contact/)