Contents

• 1. Introduction
  • 1.1. Scalable Data-Parallel Computing using GPUs
  • 1.2. Goals of PTX
  • 1.3. PTX ISA Version 8.5
  • 1.4. Document Structure
• 2. Programming Model
  • 2.1. A Highly Multithreaded Coprocessor
  • 2.2. Thread Hierarchy
    • 2.2.1. Cooperative Thread Arrays
    • 2.2.2. Cluster of Cooperative Thread Arrays
    • 2.2.3. Grid of Clusters
  • 2.3. Memory Hierarchy
• 3. PTX Machine Model
  • 3.1. A Set of SIMT Multiprocessors
  • 3.2. Independent Thread Scheduling
  • 3.3. On-chip Shared Memory
• 4. Syntax
  • 4.1. Source Format
  • 4.2. Comments
  • 4.3. Statements
    • 4.3.1. Directive Statements
    • 4.3.2. Instruction Statements
  • 4.4. Identifiers
  • 4.5. Constants
    • 4.5.1. Integer Constants
    • 4.5.2. Floating-Point Constants
    • 4.5.3. Predicate Constants
    • 4.5.4. Constant Expressions
    • 4.5.5. Integer Constant Expression Evaluation
    • 4.5.6. Summary of Constant Expression Evaluation Rules
• 5. State Spaces, Types, and Variables
  • 5.1. State Spaces
    • 5.1.1. Register State Space
    • 5.1.2. Special Register State Space
    • 5.1.3. Constant State Space
      • 5.1.3.1. Banked Constant State Space (deprecated)
    • 5.1.4. Global State Space
    • 5.1.5. Local State Space
    • 5.1.6. Parameter State Space
      • 5.1.6.1. Kernel Function Parameters
      • 5.1.6.2. Kernel Function Parameter Attributes
      • 5.1.6.3. Kernel Parameter Attribute: .ptr
      • 5.1.6.4. Device Function Parameters
    • 5.1.7. Shared State Space
    • 5.1.8. Texture State Space (deprecated)
  • 5.2. Types
    • 5.2.1. Fundamental Types
    • 5.2.2. Restricted Use of Sub-Word Sizes
    • 5.2.3. Alternate Floating-Point Data Formats
    • 5.2.4. Packed Data Types
      • 5.2.4.1. Packed Floating Point Data Types
      • 5.2.4.2. Packed Integer Data Types
  • 5.3. Texture Sampler and Surface Types
    • 5.3.1. Texture and Surface Properties
    • 5.3.2. Sampler Properties
    • 5.3.3. Channel Data Type and Channel Order Fields
  • 5.4. Variables
    • 5.4.1. Variable Declarations
    • 5.4.2. Vectors
    • 5.4.3. Array Declarations
    • 5.4.4. Initializers
    • 5.4.5. Alignment
    • 5.4.6. Parameterized Variable Names
    • 5.4.7. Variable Attributes
    • 5.4.8. Variable and Function Attribute Directive: .attribute
  • 5.5. Tensors
    • 5.5.1. Tensor Dimension, size and format
    • 5.5.2. Tensor Access Modes
    • 5.5.3. Tiled Mode
      • 5.5.3.1. Bounding Box
      • 5.5.3.2. Traversal-Stride
      • 5.5.3.3. Out of Boundary Access
    • 5.5.4. Im2col mode
      • 5.5.4.1. Bounding Box
      • 5.5.4.2. Traversal Stride
      • 5.5.4.3. Out of Boundary Access
    • 5.5.5. Interleave layout
    • 5.5.6. Swizzling Modes
    • 5.5.7. Tensor-map
• 6. Instruction Operands
  • 6.1. Operand Type Information
  • 6.2. Source Operands
  • 6.3. Destination Operands
  • 6.4. Using Addresses, Arrays, and Vectors
    • 6.4.1. Addresses as Operands
      • 6.4.1.1. Generic Addressing
    • 6.4.2. Arrays as Operands
    • 6.4.3. Vectors as Operands
    • 6.4.4. Labels and Function Names as Operands
  • 6.5. Type Conversion
    • 6.5.1. Scalar Conversions
    • 6.5.2. Rounding Modifiers
  • 6.6. Operand Costs
• 7. Abstracting the ABI
  • 7.1. Function Declarations and Definitions
    • 7.1.1. Changes from PTX ISA Version 1.x
  • 7.2. Variadic Functions
  • 7.3. Alloca
• 8. Memory Consistency Model
  • 8.1. Scope and applicability of the model
    • 8.1.1. Limitations on atomicity at system scope
  • 8.2. Memory operations
    • 8.2.1. Overlap
    • 8.2.2. Aliases
    • 8.2.3. Multimem Addresses
    • 8.2.4. Memory Operations on Vector Data Types
    • 8.2.5. Memory Operations on Packed Data Types
    • 8.2.6. Initialization
  • 8.3. State spaces
  • 8.4. Operation types
    • 8.4.1. mmio Operation
  • 8.5. Scope
  • 8.6. Proxies
  • 8.7. Morally strong operations
    • 8.7.1. Conflict and Data-races
    • 8.7.2. Limitations on Mixed-size Data-races
  • 8.8. Release and Acquire Patterns
  • 8.9. Ordering of memory operations
    • 8.9.1. Program Order
      • 8.9.1.1. Asynchronous Operations
    • 8.9.2. Observation Order
    • 8.9.3. Fence-SC Order
    • 8.9.4. Memory synchronization
    • 8.9.5. Causality Order
    • 8.9.6. Coherence Order
    • 8.9.7. Communication Order
  • 8.10. Axioms
    • 8.10.1. Coherence
    • 8.10.2. Fence-SC
    • 8.10.3. Atomicity
    • 8.10.4. No Thin Air
    • 8.10.5. Sequential Consistency Per Location
    • 8.10.6. Causality
• 9. Instruction Set
  • 9.1. Format and Semantics of Instruction Descriptions
  • 9.2. PTX Instructions
  • 9.3. Predicated Execution
    • 9.3.1. Comparisons
      • 9.3.1.1. Integer and Bit-Size Comparisons
      • 9.3.1.2. Floating Point Comparisons
    • 9.3.2. Manipulating Predicates
  • 9.4. Type Information for Instructions and Operands
    • 9.4.1. Operand Size Exceeding Instruction-Type Size
  • 9.5. Divergence of Threads in Control Constructs
  • 9.6. Semantics
    • 9.6.1. Machine-Specific Semantics of 16-bit Code
  • 9.7. Instructions
    • 9.7.1. Integer Arithmetic Instructions
      • 9.7.1.1. Integer Arithmetic Instructions: add
      • 9.7.1.2. Integer Arithmetic Instructions: sub
      • 9.7.1.3. Integer Arithmetic Instructions: mul
      • 9.7.1.4. Integer Arithmetic Instructions: mad
      • 9.7.1.5. Integer Arithmetic Instructions: mul24
      • 9.7.1.6. Integer Arithmetic Instructions: mad24
      • 9.7.1.7. Integer Arithmetic Instructions: sad
      • 9.7.1.8. Integer Arithmetic Instructions: div
      • 9.7.1.9. Integer Arithmetic Instructions: rem
      • 9.7.1.10. Integer Arithmetic Instructions: abs
      • 9.7.1.11. Integer Arithmetic Instructions: neg
      • 9.7.1.12. Integer Arithmetic Instructions: min
      • 9.7.1.13. Integer Arithmetic Instructions: max
      • 9.7.1.14. Integer Arithmetic Instructions: popc
      • 9.7.1.15. Integer Arithmetic Instructions: clz
      • 9.7.1.16. Integer Arithmetic Instructions: bfind
      • 9.7.1.17. Integer Arithmetic Instructions: fns
      • 9.7.1.18. Integer Arithmetic Instructions: brev
      • 9.7.1.19. Integer Arithmetic Instructions: bfe
      • 9.7.1.20. Integer Arithmetic Instructions: bfi
      • 9.7.1.21. Integer Arithmetic Instructions: szext
      • 9.7.1.22. Integer Arithmetic Instructions: bmsk
      • 9.7.1.23. Integer Arithmetic Instructions: dp4a
      • 9.7.1.24. Integer Arithmetic Instructions: dp2a
    • 9.7.2. Extended-Precision Integer Arithmetic Instructions
      • 9.7.2.1. Extended-Precision Arithmetic Instructions: add.cc
      • 9.7.2.2. Extended-Precision Arithmetic Instructions: addc
      • 9.7.2.3. Extended-Precision Arithmetic Instructions: sub.cc
      • 9.7.2.4. Extended-Precision Arithmetic Instructions: subc
      • 9.7.2.5. Extended-Precision Arithmetic Instructions: mad.cc
      • 9.7.2.6. Extended-Precision Arithmetic Instructions: madc
    • 9.7.3. Floating-Point Instructions
      • 9.7.3.1. Floating Point Instructions: testp
      • 9.7.3.2. Floating Point Instructions: copysign
      • 9.7.3.3. Floating Point Instructions: add
      • 9.7.3.4. Floating Point Instructions: sub
      • 9.7.3.5. Floating Point Instructions: mul
      • 9.7.3.6. Floating Point Instructions: fma
      • 9.7.3.7. Floating Point Instructions: mad
      • 9.7.3.8. Floating Point Instructions: div
      • 9.7.3.9. Floating Point Instructions: abs
      • 9.7.3.10. Floating Point Instructions: neg
      • 9.7.3.11. Floating Point Instructions: min
      • 9.7.3.12. Floating Point Instructions: max
      • 9.7.3.13. Floating Point Instructions: rcp
      • 9.7.3.14. Floating Point Instructions: rcp.approx.ftz.f64
      • 9.7.3.15. Floating Point Instructions: sqrt
      • 9.7.3.16. Floating Point Instructions: rsqrt
      • 9.7.3.17. Floating Point Instructions: rsqrt.approx.ftz.f64
      • 9.7.3.18. Floating Point Instructions: sin
      • 9.7.3.19. Floating Point Instructions: cos
      • 9.7.3.20. Floating Point Instructions: lg2
      • 9.7.3.21. Floating Point Instructions: ex2
      • 9.7.3.22. Floating Point Instructions: tanh
    • 9.7.4. Half Precision Floating-Point Instructions
      • 9.7.4.1. Half Precision Floating Point Instructions: add
      • 9.7.4.2. Half Precision Floating Point Instructions: sub
      • 9.7.4.3. Half Precision Floating Point Instructions: mul
      • 9.7.4.4. Half Precision Floating Point Instructions: fma
      • 9.7.4.5. Half Precision Floating Point Instructions: neg
      • 9.7.4.6. Half Precision Floating Point Instructions: abs
      • 9.7.4.7. Half Precision Floating Point Instructions: min
      • 9.7.4.8. Half Precision Floating Point Instructions: max
      • 9.7.4.9. Half Precision Floating Point Instructions: tanh
      • 9.7.4.10. Half Precision Floating Point Instructions: ex2
    • 9.7.7. Comparison and Selection Instructions
      • 9.7.7.1. Comparison and Selection Instructions: set
      • 9.7.7.2. Comparison and Selection Instructions: setp
      • 9.7.7.3. Comparison and Selection Instructions: selp
      • 9.7.7.4. Comparison and Selection Instructions: slct
    • 9.7.8. Half Precision Comparison Instructions
      • 9.7.8.1. Half Precision Comparison Instructions: set
      • 9.7.8.2. Half Precision Comparison Instructions: setp
    • 9.7.9. Logic and Shift Instructions
      • 9.7.9.1. Logic and Shift Instructions: and
      • 9.7.9.2. Logic and Shift Instructions: or
      • 9.7.9.3. Logic and Shift Instructions: xor
      • 9.7.9.4. Logic and Shift Instructions: not
      • 9.7.9.5. Logic and Shift Instructions: cnot
      • 9.7.9.6. Logic and Shift Instructions: lop3
      • 9.7.9.7. Logic and Shift Instructions: shf
      • 9.7.9.8. Logic and Shift Instructions: shl
      • 9.7.9.9. Logic and Shift Instructions: shr
    • 9.7.10. Data Movement and Conversion Instructions
      • 9.7.10.1. Cache Operators
      • 9.7.10.2. Cache Eviction Priority Hints
      • 9.7.10.3. Data Movement and Conversion Instructions: mov
      • 9.7.10.4. Data Movement and Conversion Instructions: mov
      • 9.7.10.5. Data Movement and Conversion Instructions: shfl (deprecated)
      • 9.7.10.6. Data Movement and Conversion Instructions: shfl.sync
      • 9.7.10.7. Data Movement and Conversion Instructions: prmt
      • 9.7.10.8. Data Movement and Conversion Instructions: ld
      • 9.7.10.9. Data Movement and Conversion Instructions: ld.global.nc
      • 9.7.10.10. Data Movement and Conversion Instructions: ldu
      • 9.7.10.11. Data Movement and Conversion Instructions: st
      • 9.7.10.12. Data Movement and Conversion Instructions: st.async
      • 9.7.10.13. Data Movement and Conversion Instructions: multimem.ld_reduce, multimem.st, multimem.red
      • 9.7.10.14. Data Movement and Conversion Instructions: prefetch, prefetchu
      • 9.7.10.15. Data Movement and Conversion Instructions: applypriority
      • 9.7.10.16. Data Movement and Conversion Instructions: discard
      • 9.7.10.17. Data Movement and Conversion Instructions: createpolicy
      • 9.7.10.18. Data Movement and Conversion Instructions: isspacep
      • 9.7.10.19. Data Movement and Conversion Instructions: cvta
      • 9.7.10.20. Data Movement and Conversion Instructions: cvt
      • 9.7.10.21. Data Movement and Conversion Instructions: cvt.pack
      • 9.7.10.22. Data Movement and Conversion Instructions: mapa
      • 9.7.10.23. Data Movement and Conversion Instructions: getctarank
      • 9.7.10.24. Data Movement and Conversion Instructions: Asynchronous copy
        • 9.7.10.24.1. Completion Mechanisms for Asynchronous Copy Operations
        • 9.7.10.24.2. Async Proxy
        • 9.7.10.24.3. Data Movement and Conversion Instructions: cp.async
        • 9.7.10.24.4. Data Movement and Conversion Instructions: cp.async.commit_group
        • 9.7.10.24.5. Data Movement and Conversion Instructions: cp.async.wait_group / cp.async.wait_all
        • 9.7.10.24.6. Data Movement and Conversion Instructions: cp.async.bulk
        • 9.7.10.24.7. Data Movement and Conversion Instructions: cp.reduce.async.bulk
        • 9.7.10.24.8. Data Movement and Conversion Instructions: cp.async.bulk.prefetch
        • 9.7.10.24.9. Data Movement and Conversion Instructions: cp.async.bulk.tensor
        • 9.7.10.24.10. Data Movement and Conversion Instructions: cp.reduce.async.bulk.tensor
        • 9.7.10.24.11. Data Movement and Conversion Instructions: cp.async.bulk.prefetch.tensor
        • 9.7.10.24.12. Data Movement and Conversion Instructions: cp.async.bulk.commit_group
        • 9.7.10.24.13. Data Movement and Conversion Instructions: cp.async.bulk.wait_group
      • 9.7.10.25. Data Movement and Conversion Instructions: tensormap.replace
    • 9.7.11. Texture Instructions
      • 9.7.11.1. Texturing Modes
      • 9.7.11.2. Mipmaps
      • 9.7.11.3. Texture Instructions: tex
      • 9.7.11.4. Texture Instructions: tld4
      • 9.7.11.5. Texture Instructions: txq
      • 9.7.11.6. Texture Instructions: istypep
    • 9.7.12. Surface Instructions
      • 9.7.12.1. Surface Instructions: suld
      • 9.7.12.2. Surface Instructions: sust
      • 9.7.12.3. Surface Instructions: sured
      • 9.7.12.4. Surface Instructions: suq
    • 9.7.13. Control Flow Instructions
      • 9.7.13.1. Control Flow Instructions: {}
      • 9.7.13.2. Control Flow Instructions: @
      • 9.7.13.3. Control Flow Instructions: bra
      • 9.7.13.4. Control Flow Instructions: brx.idx
      • 9.7.13.5. Control Flow Instructions: call
      • 9.7.13.6. Control Flow Instructions: ret
      • 9.7.13.7. Control Flow Instructions: exit
    • 9.7.14. Parallel Synchronization and Communication Instructions
      • 9.7.14.1. Parallel Synchronization and Communication Instructions: bar, barrier
      • 9.7.14.2. Parallel Synchronization and Communication Instructions: bar.warp.sync
      • 9.7.14.3. Parallel Synchronization and Communication Instructions: barrier.cluster
      • 9.7.14.4. Parallel Synchronization and Communication Instructions: membar/fence
      • 9.7.14.5. Parallel Synchronization and Communication Instructions: atom
      • 9.7.14.6. Parallel Synchronization and Communication Instructions: red
      • 9.7.14.7. Parallel Synchronization and Communication Instructions: red.async
      • 9.7.14.8. Parallel Synchronization and Communication Instructions: vote (deprecated)
      • 9.7.14.9. Parallel Synchronization and Communication Instructions: vote.sync
      • 9.7.14.10. Parallel Synchronization and Communication Instructions: match.sync
      • 9.7.14.11. Parallel Synchronization and Communication Instructions: activemask
      • 9.7.14.12. Parallel Synchronization and Communication Instructions: redux.sync
      • 9.7.14.13. Parallel Synchronization and Communication Instructions: griddepcontrol
      • 9.7.14.14. Parallel Synchronization and Communication Instructions: elect.sync
      • 9.7.14.15. Parallel Synchronization and Communication Instructions: mbarrier
        • 9.7.14.15.1. Size and alignment of mbarrier object
        • 9.7.14.15.2. Contents of the mbarrier object
        • 9.7.14.15.3. Lifecycle of the mbarrier object
        • 9.7.14.15.4. Phase of the mbarrier object
        • 9.7.14.15.5. Tracking asynchronous operations by the mbarrier object
          • 9.7.14.15.5.1. expect-tx operation
          • 9.7.14.15.5.2. complete-tx operation
        • 9.7.14.15.6. Phase Completion of the mbarrier object
        • 9.7.14.15.7. Arrive-on operation on mbarrier object
        • 9.7.14.15.8. mbarrier support with shared memory
        • 9.7.14.15.9. Parallel Synchronization and Communication Instructions: mbarrier.init
        • 9.7.14.15.10. Parallel Synchronization and Communication Instructions: mbarrier.inval
        • 9.7.14.15.11. Parallel Synchronization and Communication Instructions: mbarrier.expect_tx
        • 9.7.14.15.12. Parallel Synchronization and Communication Instructions: mbarrier.complete_tx
        • 9.7.14.15.13. Parallel Synchronization and Communication Instructions: mbarrier.arrive
        • 9.7.14.15.14. Parallel Synchronization and Communication Instructions: mbarrier.arrive_drop
        • 9.7.14.15.15. Parallel Synchronization and Communication Instructions: cp.async.mbarrier.arrive
        • 9.7.14.15.16. Parallel Synchronization and Communication Instructions: mbarrier.test_wait/mbarrier.try_wait
        • 9.7.14.15.17. Parallel Synchronization and Communication Instructions: mbarrier.pending_count
      • 9.7.14.16. Parallel Synchronization and Communication Instructions: tensormap.cp_fenceproxy
    • 9.7.15. Warp Level Matrix Multiply-Accumulate Instructions
      • 9.7.15.1. Matrix Shape
      • 9.7.15.2. Matrix Data-types
      • 9.7.15.3. Matrix multiply-accumulate operation using wmma instructions
        • 9.7.15.3.1. Matrix Fragments for WMMA
        • 9.7.15.3.2. Matrix Storage for WMMA
        • 9.7.15.3.3. Warp-level Matrix Load Instruction: wmma.load
        • 9.7.15.3.4. Warp-level Matrix Store Instruction: wmma.store
        • 9.7.15.3.5. Warp-level Matrix Multiply-and-Accumulate Instruction: wmma.mma
      • 9.7.15.4. Matrix multiply-accumulate operation using mma instruction
        • 9.7.15.4.1. Matrix Fragments for mma.m8n8k4 with .f16 floating point type
        • 9.7.15.4.2. Matrix Fragments for mma.m8n8k4 with .f64 floating point type
        • 9.7.15.4.3. Matrix Fragments for mma.m8n8k16
        • 9.7.15.4.4. Matrix Fragments for mma.m8n8k32
        • 9.7.15.4.5. Matrix Fragments for mma.m8n8k128
        • 9.7.15.4.6. Matrix Fragments for mma.m16n8k4
        • 9.7.15.4.7. Matrix Fragments for mma.m16n8k8
        • 9.7.15.4.8. Matrix Fragments for mma.m16n8k16 with floating point type
        • 9.7.15.4.9. Matrix Fragments for mma.m16n8k16 with integer type
        • 9.7.15.4.10. Matrix Fragments for mma.m16n8k32
        • 9.7.15.4.11. Matrix Fragments for mma.m16n8k64
        • 9.7.15.4.12. Matrix Fragments for mma.m16n8k128
        • 9.7.15.4.13. Matrix Fragments for mma.m16n8k256
        • 9.7.15.4.14. Multiply-and-Accumulate Instruction: mma
        • 9.7.15.4.15. Warp-level matrix load instruction: ldmatrix
        • 9.7.15.4.16. Warp-level matrix store instruction: stmatrix
        • 9.7.15.4.17. Warp-level matrix transpose instruction: movmatrix
      • 9.7.15.5. Matrix multiply-accumulate operation using mma.sp instruction with sparse matrix A
        • 9.7.15.5.1. Sparse matrix storage
        • 9.7.15.5.2. Matrix fragments for multiply-accumulate operation with sparse matrix A
          • 9.7.15.5.2.1. Matrix Fragments for sparse mma.m16n8k16 with .f16 and .bf16 types
          • 9.7.15.5.2.2. Matrix Fragments for sparse mma.m16n8k32 with .f16 and .bf16 types
          • 9.7.15.5.2.3. Matrix Fragments for sparse mma.m16n8k16 with .tf32 floating point type
          • 9.7.15.5.2.4. Matrix Fragments for sparse mma.m16n8k8 with .tf32 floating point type
          • 9.7.15.5.2.5. Matrix Fragments for sparse mma.m16n8k32 with .u8/.s8 integer type
          • 9.7.15.5.2.6. Matrix Fragments for sparse mma.m16n8k64 with .u8/.s8/.e4m3/.e5m2 type
          • 9.7.15.5.2.7. Matrix Fragments for sparse mma.m16n8k64 with .u4/.s4 integer type
          • 9.7.15.5.2.8. Matrix Fragments for sparse mma.m16n8k128 with .u4/.s4 integer type
        • 9.7.15.5.3. Multiply-and-Accumulate Instruction: mma.sp/mma.sp::ordered_metadata
    • 9.7.16. Asynchronous Warpgroup Level Matrix Multiply-Accumulate Instructions
      • 9.7.16.1. Warpgroup
      • 9.7.16.2. Matrix Shape
      • 9.7.16.3. Matrix Data-types
      • 9.7.16.4. Async Proxy
      • 9.7.16.5. Asynchronous Warpgroup Level Matrix Multiply-Accumulate Operation using wgmma.mma_async instruction
        • 9.7.16.5.1. Register Fragments and Shared Memory Matrix Layouts
          • 9.7.16.5.1.1. Register Fragments
            • 9.7.16.5.1.1.1. Matrix Fragments for wgmma.mma_async.m64nNk16
            • 9.7.16.5.1.1.2. Matrix Fragments for wgmma.mma_async.m64nNk8
            • 9.7.16.5.1.1.3. Matrix Fragments for wgmma.mma_async.m64nNk32
            • 9.7.16.5.1.1.4. Matrix Fragments for wgmma.mma_async.m64nNk256
          • 9.7.16.5.1.2. Shared Memory Matrix Layout
            • 9.7.16.5.1.2.1. Shared Memory Layout for wgmma.mma_async.m64nNk16
            • 9.7.16.5.1.2.2. Shared Memory Layout for wgmma.mma_async.m64nNk8
            • 9.7.16.5.1.2.3. Shared Memory Layout for wgmma.mma_async.m64nNk32
            • 9.7.16.5.1.2.4. Shared Memory Layout for wgmma.mma_async.m64nNk256
            • 9.7.16.5.1.2.5. Strides
            • 9.7.16.5.1.2.6. Swizzling Modes
            • 9.7.16.5.1.2.7. Matrix Descriptor Format
        • 9.7.16.5.2. Asynchronous Multiply-and-Accumulate Instruction: wgmma.mma_async
      • 9.7.16.6. Asynchronous Warpgroup Level Multiply-and-Accumulate Operation using wgmma.mma_async.sp instruction
        • 9.7.16.6.1. Sparse matrix storage
        • 9.7.16.6.2. Matrix fragments for warpgroup-level multiply-accumulate operation with sparse matrix A
          • 9.7.16.6.2.1. Matrix Fragments for sparse wgmma.mma_async.m64nNk32
          • 9.7.16.6.2.2. Matrix Fragments for sparse wgmma.mma_async.m64nNk16
          • 9.7.16.6.2.3. Matrix Fragments for sparse wgmma.mma_async.m64nNk64
        • 9.7.16.6.3. Shared Memory Matrix Layout
          • 9.7.16.6.3.1. Shared Memory Layout for wgmma.mma_async.sp.m64nNk32
          • 9.7.16.6.3.2. Shared Memory Layout for wgmma.mma_async.sp.m64nNk16
          • 9.7.16.6.3.3. Shared Memory Layout for wgmma.mma_async.sp.m64nNk64
        • 9.7.16.6.4. Asynchronous Multiply-and-Accumulate Instruction: wgmma.mma_async.sp
      • 9.7.16.7. Asynchronous wgmma Proxy Operations
        • 9.7.16.7.1. Asynchronous Multiply-and-Accumulate Instruction: wgmma.fence
        • 9.7.16.7.2. Asynchronous Multiply-and-Accumulate Instruction: wgmma.commit_group
        • 9.7.16.7.3. Asynchronous Multiply-and-Accumulate Instruction: wgmma.wait_group
    • 9.7.17. Stack Manipulation Instructions
      • 9.7.17.1. Stack Manipulation Instructions: stacksave
      • 9.7.17.2. Stack Manipulation Instructions: stackrestore
      • 9.7.17.3. Stack Manipulation Instructions: alloca
    • 9.7.18. Video Instructions
      • 9.7.18.1. Scalar Video Instructions
        • 9.7.18.1.1. Scalar Video Instructions: vadd, vsub, vabsdiff, vmin, vmax
        • 9.7.18.1.2. Scalar Video Instructions: vshl, vshr
        • 9.7.18.1.3. Scalar Video Instructions: vmad
        • 9.7.18.1.4. Scalar Video Instructions: vset
      • 9.7.18.2. SIMD Video Instructions
        • 9.7.18.2.1. SIMD Video Instructions: vadd2, vsub2, vavrg2, vabsdiff2, vmin2, vmax2
        • 9.7.18.2.2. SIMD Video Instructions: vset2
        • 9.7.18.2.3. SIMD Video Instructions: vadd4, vsub4, vavrg4, vabsdiff4, vmin4, vmax4
        • 9.7.18.2.4. SIMD Video Instructions: vset4
    • 9.7.19. Miscellaneous Instructions
      • 9.7.19.1. Miscellaneous Instructions: brkpt
      • 9.7.19.2. Miscellaneous Instructions: nanosleep
      • 9.7.19.3. Miscellaneous Instructions: pmevent
      • 9.7.19.4. Miscellaneous Instructions: trap
      • 9.7.19.5. Miscellaneous Instructions: setmaxnreg
• 10. Special Registers
  • 10.1. Special Registers: %tid
  • 10.2. Special Registers: %ntid
  • 10.3. Special Registers: %laneid
  • 10.4. Special Registers: %warpid
  • 10.5. Special Registers: %nwarpid
  • 10.6. Special Registers: %ctaid
  • 10.7. Special Registers: %nctaid
  • 10.8. Special Registers: %smid
  • 10.9. Special Registers: %nsmid
  • 10.10. Special Registers: %gridid
  • 10.11. Special Registers: %is_explicit_cluster
  • 10.12. Special Registers: %clusterid
  • 10.13. Special Registers: %nclusterid
  • 10.14. Special Registers: %cluster_ctaid
  • 10.15. Special Registers: %cluster_nctaid
  • 10.16. Special Registers: %cluster_ctarank
  • 10.17. Special Registers: %cluster_nctarank
  • 10.18. Special Registers: %lanemask_eq
  • 10.19. Special Registers: %lanemask_le
  • 10.20. Special Registers: %lanemask_lt
  • 10.21. Special Registers: %lanemask_ge
  • 10.22. Special Registers: %lanemask_gt
  • 10.23. Special Registers: %clock, %clock_hi
  • 10.24. Special Registers: %clock64
  • 10.25. Special Registers: %pm0..%pm7
  • 10.26. Special Registers: %pm0_64..%pm7_64
  • 10.27. Special Registers: %envreg<32>
  • 10.28. Special Registers: %globaltimer, %globaltimer_lo, %globaltimer_hi
  • 10.29. Special Registers: %reserved_smem_offset_begin, %reserved_smem_offset_end, %reserved_smem_offset_cap, %reserved_smem_offset_<2>
  • 10.30. Special Registers: %total_smem_size
  • 10.31. Special Registers: %aggr_smem_size
  • 10.32. Special Registers: %dynamic_smem_size
  • 10.33. Special Registers: %current_graph_exec
• 11. Directives
  • 11.1. PTX Module Directives
    • 11.1.1. PTX Module Directives: .version
    • 11.1.2. PTX Module Directives: .target
    • 11.1.3. PTX Module Directives: .address_size
  • 11.2. Specifying Kernel Entry Points and Functions
    • 11.2.1. Kernel and Function Directives: .entry
    • 11.2.2. Kernel and Function Directives: .func
    • 11.2.3. Kernel and Function Directives: .alias
  • 11.3. Control Flow Directives
    • 11.3.1. Control Flow Directives: .branchtargets
    • 11.3.2. Control Flow Directives: .calltargets
    • 11.3.3. Control Flow Directives: .callprototype
  • 11.4. Performance-Tuning Directives
    • 11.4.1. Performance-Tuning Directives: .maxnreg
    • 11.4.2. Performance-Tuning Directives: .maxntid
    • 11.4.3. Performance-Tuning Directives: .reqntid
    • 11.4.4. Performance-Tuning Directives: .minnctapersm
    • 11.4.5. Performance-Tuning Directives: .maxnctapersm (deprecated)
    • 11.4.6. Performance-Tuning Directives: .noreturn
    • 11.4.7. Performance-Tuning Directives: .pragma
  • 11.5. Debugging Directives
    • 11.5.1. Debugging Directives: @@dwarf
    • 11.5.2. Debugging Directives: .section
    • 11.5.3. Debugging Directives: .file
    • 11.5.4. Debugging Directives: .loc
  • 11.6. Linking Directives
    • 11.6.1. Linking Directives: .extern
    • 11.6.2. Linking Directives: .visible
    • 11.6.3. Linking Directives: .weak
    • 11.6.4. Linking Directives: .common
  • 11.7. Cluster Dimension Directives
    • 11.7.1. Cluster Dimension Directives: .reqnctapercluster
    • 11.7.2. Cluster Dimension Directives: .explicitcluster
    • 11.7.3. Cluster Dimension Directives: .maxclusterrank
• 12. Release Notes
  • 12.1. Changes in PTX ISA Version 8.5
  • 12.2. Changes in PTX ISA Version 8.4
  • 12.3. Changes in PTX ISA Version 8.3
  • 12.4. Changes in PTX ISA Version 8.2
  • 12.5. Changes in PTX ISA Version 8.1
  • 12.6. Changes in PTX ISA Version 8.0
  • 12.7. Changes in PTX ISA Version 7.8
  • 12.8. Changes in PTX ISA Version 7.7
  • 12.9. Changes in PTX ISA Version 7.6
  • 12.10. Changes in PTX ISA Version 7.5
  • 12.11. Changes in PTX ISA Version 7.4
  • 12.12. Changes in PTX ISA Version 7.3
  • 12.13. Changes in PTX ISA Version 7.2
  • 12.14. Changes in PTX ISA Version 7.1
  • 12.15. Changes in PTX ISA Version 7.0
  • 12.16. Changes in PTX ISA Version 6.5
  • 12.17. Changes in PTX ISA Version 6.4
  • 12.18. Changes in PTX ISA Version 6.3
  • 12.19. Changes in PTX ISA Version 6.2
  • 12.20. Changes in PTX ISA Version 6.1
  • 12.21. Changes in PTX ISA Version 6.0
  • 12.22. Changes in PTX ISA Version 5.0
  • 12.23. Changes in PTX ISA Version 4.3
  • 12.24. Changes in PTX ISA Version 4.2
  • 12.25. Changes in PTX ISA Version 4.1
  • 12.26. Changes in PTX ISA Version 4.0
  • 12.27. Changes in PTX ISA Version 3.2
  • 12.28. Changes in PTX ISA Version 3.1
  • 12.29. Changes in PTX ISA Version 3.0
  • 12.30. Changes in PTX ISA Version 2.3
  • 12.31. Changes in PTX ISA Version 2.2
  • 12.32. Changes in PTX ISA Version 2.1
  • 12.33. Changes in PTX ISA Version 2.0
• 14. Descriptions of .pragma Strings
  • 14.1. Pragma Strings: “nounroll”
  • 14.2. Pragma Strings: “used_bytes_mask”
• 15. Notices
  • 15.1. Notice
  • 15.2. OpenCL
  • 15.3. Trademarks