Contents

• 1. Introduction
  • 1.1. The Benefits of Using GPUs
  • 1.2. CUDA®: A General-Purpose Parallel Computing Platform and Programming Model
  • 1.3. A Scalable Programming Model
  • 1.4. Document Structure
• 2. Programming Model
  • 2.1. Kernels
  • 2.2. Thread Hierarchy
    • 2.2.1. Thread Block Clusters
  • 2.3. Memory Hierarchy
  • 2.4. Heterogeneous Programming
  • 2.5. Asynchronous SIMT Programming Model
    • 2.5.1. Asynchronous Operations
  • 2.6. Compute Capability
• 3. Programming Interface
  • 3.1. Compilation with NVCC
    • 3.1.1. Compilation Workflow
      • 3.1.1.1. Offline Compilation
      • 3.1.1.2. Just-in-Time Compilation
    • 3.1.2. Binary Compatibility
    • 3.1.3. PTX Compatibility
    • 3.1.4. Application Compatibility
    • 3.1.5. C++ Compatibility
    • 3.1.6. 64-Bit Compatibility
  • 3.2. CUDA Runtime
    • 3.2.1. Initialization
    • 3.2.2. Device Memory
    • 3.2.3. Device Memory L2 Access Management
      • 3.2.3.1. L2 cache Set-Aside for Persisting Accesses
      • 3.2.3.2. L2 Policy for Persisting Accesses
      • 3.2.3.3. L2 Access Properties
      • 3.2.3.4. L2 Persistence Example
      • 3.2.3.5. Reset L2 Access to Normal
      • 3.2.3.6. Manage Utilization of L2 set-aside cache
      • 3.2.3.7. Query L2 cache Properties
      • 3.2.3.8. Control L2 Cache Set-Aside Size for Persisting Memory Access
    • 3.2.4. Shared Memory
    • 3.2.5. Distributed Shared Memory
    • 3.2.6. Page-Locked Host Memory
      • 3.2.6.1. Portable Memory
      • 3.2.6.2. Write-Combining Memory
      • 3.2.6.3. Mapped Memory
    • 3.2.7. Memory Synchronization Domains
      • 3.2.7.1. Memory Fence Interference
      • 3.2.7.2. Isolating Traffic with Domains
      • 3.2.7.3. Using Domains in CUDA
    • 3.2.8. Asynchronous Concurrent Execution
      • 3.2.8.1. Concurrent Execution between Host and Device
      • 3.2.8.2. Concurrent Kernel Execution
      • 3.2.8.3. Overlap of Data Transfer and Kernel Execution
      • 3.2.8.4. Concurrent Data Transfers
      • 3.2.8.5. Streams
        • 3.2.8.5.1. Creation and Destruction of Streams
        • 3.2.8.5.2. Default Stream
        • 3.2.8.5.3. Explicit Synchronization
        • 3.2.8.5.4. Implicit Synchronization
        • 3.2.8.5.5. Overlapping Behavior
        • 3.2.8.5.6. Host Functions (Callbacks)
        • 3.2.8.5.7. Stream Priorities
      • 3.2.8.6. Programmatic Dependent Launch and Synchronization
        • 3.2.8.6.1. Background
        • 3.2.8.6.2. API Description
        • 3.2.8.6.3. Use in CUDA Graphs
      • 3.2.8.7. CUDA Graphs
        • 3.2.8.7.1. Graph Structure
          • 3.2.8.7.1.1. Node Types
          • 3.2.8.7.1.2. Edge Data
        • 3.2.8.7.2. Creating a Graph Using Graph APIs
        • 3.2.8.7.3. Creating a Graph Using Stream Capture
          • 3.2.8.7.3.1. Cross-stream Dependencies and Events
          • 3.2.8.7.3.2. Prohibited and Unhandled Operations
          • 3.2.8.7.3.3. Invalidation
        • 3.2.8.7.4. CUDA User Objects
        • 3.2.8.7.5. Updating Instantiated Graphs
          • 3.2.8.7.5.1. Graph Update Limitations
          • 3.2.8.7.5.2. Whole Graph Update
          • 3.2.8.7.5.3. Individual node update
          • 3.2.8.7.5.4. Individual node enable
        • 3.2.8.7.6. Using Graph APIs
        • 3.2.8.7.7. Device Graph Launch
          • 3.2.8.7.7.1. Device Graph Creation
            • 3.2.8.7.7.1.1. Device Graph Requirements
            • 3.2.8.7.7.1.2. Device Graph Upload
            • 3.2.8.7.7.1.3. Device Graph Update
          • 3.2.8.7.7.2. Device Launch
            • 3.2.8.7.7.2.1. Device Launch Modes
              • 3.2.8.7.7.2.1.1. Fire and Forget Launch
              • 3.2.8.7.7.2.1.2. Graph Execution Environments
              • 3.2.8.7.7.2.1.3. Tail Launch
                • 3.2.8.7.7.2.1.3.1. Tail Self-launch
              • 3.2.8.7.7.2.1.4. Sibling Launch
        • 3.2.8.7.8. Conditional Graph Nodes
          • 3.2.8.7.8.1. Conditional Handles
          • 3.2.8.7.8.2. Condtional Node Body Graph Requirements
          • 3.2.8.7.8.3. Conditional IF Nodes
          • 3.2.8.7.8.4. Conditional WHILE Nodes
      • 3.2.8.8. Events
        • 3.2.8.8.1. Creation and Destruction of Events
        • 3.2.8.8.2. Elapsed Time
      • 3.2.8.9. Synchronous Calls
    • 3.2.9. Multi-Device System
      • 3.2.9.1. Device Enumeration
      • 3.2.9.2. Device Selection
      • 3.2.9.3. Stream and Event Behavior
      • 3.2.9.4. Peer-to-Peer Memory Access
        • 3.2.9.4.1. IOMMU on Linux
      • 3.2.9.5. Peer-to-Peer Memory Copy
    • 3.2.10. Unified Virtual Address Space
    • 3.2.11. Interprocess Communication
    • 3.2.12. Error Checking
    • 3.2.13. Call Stack
    • 3.2.14. Texture and Surface Memory
      • 3.2.14.1. Texture Memory
        • 3.2.14.1.1. Texture Object API
        • 3.2.14.1.2. 16-Bit Floating-Point Textures
        • 3.2.14.1.3. Layered Textures
        • 3.2.14.1.4. Cubemap Textures
        • 3.2.14.1.5. Cubemap Layered Textures
        • 3.2.14.1.6. Texture Gather
      • 3.2.14.2. Surface Memory
        • 3.2.14.2.1. Surface Object API
        • 3.2.14.2.2. Cubemap Surfaces
        • 3.2.14.2.3. Cubemap Layered Surfaces
      • 3.2.14.3. CUDA Arrays
      • 3.2.14.4. Read/Write Coherency
    • 3.2.15. Graphics Interoperability
      • 3.2.15.1. OpenGL Interoperability
      • 3.2.15.2. Direct3D Interoperability
        • 3.2.15.2.1. Direct3D 9 Version
        • 3.2.15.2.2. Direct3D 10 Version
        • 3.2.15.2.3. Direct3D 11 Version
      • 3.2.15.3. SLI Interoperability
    • 3.2.16. External Resource Interoperability
      • 3.2.16.1. Vulkan Interoperability
        • 3.2.16.1.1. Matching device UUIDs
        • 3.2.16.1.2. Importing Memory Objects
        • 3.2.16.1.3. Mapping Buffers onto Imported Memory Objects
        • 3.2.16.1.4. Mapping Mipmapped Arrays onto Imported Memory Objects
        • 3.2.16.1.5. Importing Synchronization Objects
        • 3.2.16.1.6. Signaling/Waiting on Imported Synchronization Objects
      • 3.2.16.2. OpenGL Interoperability
      • 3.2.16.3. Direct3D 12 Interoperability
        • 3.2.16.3.1. Matching Device LUIDs
        • 3.2.16.3.2. Importing Memory Objects
        • 3.2.16.3.3. Mapping Buffers onto Imported Memory Objects
        • 3.2.16.3.4. Mapping Mipmapped Arrays onto Imported Memory Objects
        • 3.2.16.3.5. Importing Synchronization Objects
        • 3.2.16.3.6. Signaling/Waiting on Imported Synchronization Objects
      • 3.2.16.4. Direct3D 11 Interoperability
        • 3.2.16.4.1. Matching Device LUIDs
        • 3.2.16.4.2. Importing Memory Objects
        • 3.2.16.4.3. Mapping Buffers onto Imported Memory Objects
        • 3.2.16.4.4. Mapping Mipmapped Arrays onto Imported Memory Objects
        • 3.2.16.4.5. Importing Synchronization Objects
        • 3.2.16.4.6. Signaling/Waiting on Imported Synchronization Objects
      • 3.2.16.5. NVIDIA Software Communication Interface Interoperability (NVSCI)
        • 3.2.16.5.1. Importing Memory Objects
        • 3.2.16.5.2. Mapping Buffers onto Imported Memory Objects
        • 3.2.16.5.3. Mapping Mipmapped Arrays onto Imported Memory Objects
        • 3.2.16.5.4. Importing Synchronization Objects
        • 3.2.16.5.5. Signaling/Waiting on Imported Synchronization Objects
  • 3.3. Versioning and Compatibility
  • 3.4. Compute Modes
  • 3.5. Mode Switches
  • 3.6. Tesla Compute Cluster Mode for Windows
• 4. Hardware Implementation
  • 4.1. SIMT Architecture
  • 4.2. Hardware Multithreading
• 5. Performance Guidelines
  • 5.1. Overall Performance Optimization Strategies
  • 5.2. Maximize Utilization
    • 5.2.1. Application Level
    • 5.2.2. Device Level
    • 5.2.3. Multiprocessor Level
      • 5.2.3.1. Occupancy Calculator
  • 5.3. Maximize Memory Throughput
    • 5.3.1. Data Transfer between Host and Device
    • 5.3.2. Device Memory Accesses
  • 5.4. Maximize Instruction Throughput
    • 5.4.1. Arithmetic Instructions
    • 5.4.2. Control Flow Instructions
    • 5.4.3. Synchronization Instruction
  • 5.5. Minimize Memory Thrashing
• 6. CUDA-Enabled GPUs
• 7. C++ Language Extensions
  • 7.1. Function Execution Space Specifiers
    • 7.1.1. __global__
    • 7.1.2. __device__
    • 7.1.3. __host__
    • 7.1.4. Undefined behavior
    • 7.1.5. __noinline__ and __forceinline__
    • 7.1.6. __inline_hint__
  • 7.2. Variable Memory Space Specifiers
    • 7.2.1. __device__
    • 7.2.2. __constant__
    • 7.2.3. __shared__
    • 7.2.4. __grid_constant__
    • 7.2.5. __managed__
    • 7.2.6. __restrict__
  • 7.3. Built-in Vector Types
    • 7.3.1. char, short, int, long, longlong, float, double
    • 7.3.2. dim3
  • 7.4. Built-in Variables
    • 7.4.1. gridDim
    • 7.4.2. blockIdx
    • 7.4.3. blockDim
    • 7.4.4. threadIdx
    • 7.4.5. warpSize
  • 7.5. Memory Fence Functions
  • 7.6. Synchronization Functions
  • 7.7. Mathematical Functions
  • 7.8. Texture Functions
    • 7.8.1. Texture Object API
      • 7.8.1.1. tex1Dfetch()
      • 7.8.1.2. tex1D()
      • 7.8.1.3. tex1DLod()
      • 7.8.1.4. tex1DGrad()
      • 7.8.1.5. tex2D()
      • 7.8.1.6. tex2D() for sparse CUDA arrays
      • 7.8.1.7. tex2Dgather()
      • 7.8.1.8. tex2Dgather() for sparse CUDA arrays
      • 7.8.1.9. tex2DGrad()
      • 7.8.1.10. tex2DGrad() for sparse CUDA arrays
      • 7.8.1.11. tex2DLod()
      • 7.8.1.12. tex2DLod() for sparse CUDA arrays
      • 7.8.1.13. tex3D()
      • 7.8.1.14. tex3D() for sparse CUDA arrays
      • 7.8.1.15. tex3DLod()
      • 7.8.1.16. tex3DLod() for sparse CUDA arrays
      • 7.8.1.17. tex3DGrad()
      • 7.8.1.18. tex3DGrad() for sparse CUDA arrays
      • 7.8.1.19. tex1DLayered()
      • 7.8.1.20. tex1DLayeredLod()
      • 7.8.1.21. tex1DLayeredGrad()
      • 7.8.1.22. tex2DLayered()
      • 7.8.1.23. tex2DLayered() for sparse CUDA arrays
      • 7.8.1.24. tex2DLayeredLod()
      • 7.8.1.25. tex2DLayeredLod() for sparse CUDA arrays
      • 7.8.1.26. tex2DLayeredGrad()
      • 7.8.1.27. tex2DLayeredGrad() for sparse CUDA arrays
      • 7.8.1.28. texCubemap()
      • 7.8.1.29. texCubemapGrad()
      • 7.8.1.30. texCubemapLod()
      • 7.8.1.31. texCubemapLayered()
      • 7.8.1.32. texCubemapLayeredGrad()
      • 7.8.1.33. texCubemapLayeredLod()
  • 7.9. Surface Functions
    • 7.9.1. Surface Object API
      • 7.9.1.1. surf1Dread()
      • 7.9.1.2. surf1Dwrite
      • 7.9.1.3. surf2Dread()
      • 7.9.1.4. surf2Dwrite()
      • 7.9.1.5. surf3Dread()
      • 7.9.1.6. surf3Dwrite()
      • 7.9.1.7. surf1DLayeredread()
      • 7.9.1.8. surf1DLayeredwrite()
      • 7.9.1.9. surf2DLayeredread()
      • 7.9.1.10. surf2DLayeredwrite()
      • 7.9.1.11. surfCubemapread()
      • 7.9.1.12. surfCubemapwrite()
      • 7.9.1.13. surfCubemapLayeredread()
      • 7.9.1.14. surfCubemapLayeredwrite()
  • 7.10. Read-Only Data Cache Load Function
  • 7.11. Load Functions Using Cache Hints
  • 7.12. Store Functions Using Cache Hints
  • 7.13. Time Function
  • 7.14. Atomic Functions
    • 7.14.1. Arithmetic Functions
      • 7.14.1.1. atomicAdd()
      • 7.14.1.2. atomicSub()
      • 7.14.1.3. atomicExch()
      • 7.14.1.4. atomicMin()
      • 7.14.1.5. atomicMax()
      • 7.14.1.6. atomicInc()
      • 7.14.1.7. atomicDec()
      • 7.14.1.8. atomicCAS()
    • 7.14.2. Bitwise Functions
      • 7.14.2.1. atomicAnd()
      • 7.14.2.2. atomicOr()
      • 7.14.2.3. atomicXor()
  • 7.15. Address Space Predicate Functions
    • 7.15.1. __isGlobal()
    • 7.15.2. __isShared()
    • 7.15.3. __isConstant()
    • 7.15.4. __isGridConstant()
    • 7.15.5. __isLocal()
  • 7.16. Address Space Conversion Functions
    • 7.16.1. __cvta_generic_to_global()
    • 7.16.2. __cvta_generic_to_shared()
    • 7.16.3. __cvta_generic_to_constant()
    • 7.16.4. __cvta_generic_to_local()
    • 7.16.5. __cvta_global_to_generic()
    • 7.16.6. __cvta_shared_to_generic()
    • 7.16.7. __cvta_constant_to_generic()
    • 7.16.8. __cvta_local_to_generic()
  • 7.17. Alloca Function
    • 7.17.1. Synopsis
    • 7.17.2. Description
    • 7.17.3. Example
  • 7.18. Compiler Optimization Hint Functions
    • 7.18.1. __builtin_assume_aligned()
    • 7.18.2. __builtin_assume()
    • 7.18.3. __assume()
    • 7.18.4. __builtin_expect()
    • 7.18.5. __builtin_unreachable()
    • 7.18.6. Restrictions
  • 7.19. Warp Vote Functions
  • 7.20. Warp Match Functions
    • 7.20.1. Synopsis
    • 7.20.2. Description
  • 7.21. Warp Reduce Functions
    • 7.21.1. Synopsis
    • 7.21.2. Description
  • 7.22. Warp Shuffle Functions
    • 7.22.1. Synopsis
    • 7.22.2. Description
    • 7.22.3. Examples
      • 7.22.3.1. Broadcast of a single value across a warp
      • 7.22.3.2. Inclusive plus-scan across sub-partitions of 8 threads
      • 7.22.3.3. Reduction across a warp
  • 7.23. Nanosleep Function
    • 7.23.1. Synopsis
    • 7.23.2. Description
    • 7.23.3. Example
  • 7.24. Warp Matrix Functions
    • 7.24.1. Description
    • 7.24.2. Alternate Floating Point
    • 7.24.3. Double Precision
    • 7.24.4. Sub-byte Operations
    • 7.24.5. Restrictions
    • 7.24.6. Element Types and Matrix Sizes
    • 7.24.7. Example
  • 7.25. DPX
    • 7.25.1. Examples
  • 7.26. Asynchronous Barrier
    • 7.26.1. Simple Synchronization Pattern
    • 7.26.2. Temporal Splitting and Five Stages of Synchronization
    • 7.26.3. Bootstrap Initialization, Expected Arrival Count, and Participation
    • 7.26.4. A Barrier’s Phase: Arrival, Countdown, Completion, and Reset
    • 7.26.5. Spatial Partitioning (also known as Warp Specialization)
    • 7.26.6. Early Exit (Dropping out of Participation)
    • 7.26.7. Completion Function
    • 7.26.8. Memory Barrier Primitives Interface
      • 7.26.8.1. Data Types
      • 7.26.8.2. Memory Barrier Primitives API
  • 7.27. Asynchronous Data Copies
    • 7.27.1. memcpy_async API
    • 7.27.2. Copy and Compute Pattern - Staging Data Through Shared Memory
    • 7.27.3. Without memcpy_async
    • 7.27.4. With memcpy_async
    • 7.27.5. Asynchronous Data Copies using cuda::barrier
    • 7.27.6. Performance Guidance for memcpy_async
      • 7.27.6.1. Alignment
      • 7.27.6.2. Trivially copyable
      • 7.27.6.3. Warp Entanglement - Commit
      • 7.27.6.4. Warp Entanglement - Wait
      • 7.27.6.5. Warp Entanglement - Arrive-On
      • 7.27.6.6. Keep Commit and Arrive-On Operations Converged
  • 7.28. Asynchronous Data Copies using cuda::pipeline
    • 7.28.1. Single-Stage Asynchronous Data Copies using cuda::pipeline
    • 7.28.2. Multi-Stage Asynchronous Data Copies using cuda::pipeline
    • 7.28.3. Pipeline Interface
    • 7.28.4. Pipeline Primitives Interface
      • 7.28.4.1. memcpy_async Primitive
      • 7.28.4.2. Commit Primitive
      • 7.28.4.3. Wait Primitive
      • 7.28.4.4. Arrive On Barrier Primitive
  • 7.29. Asynchronous Data Copies using the Tensor Memory Accelerator (TMA)
    • 7.29.1. Using TMA to transfer one-dimensional arrays
    • 7.29.2. Using TMA to transfer multi-dimensional arrays
      • 7.29.2.1. Multi-dimensional TMA PTX wrappers
  • 7.30. Encoding a Tensor Map on Device
    • 7.30.1. Device-side Encoding and Modification of a Tensor Map
    • 7.30.2. Usage of a Modified Tensor Map
    • 7.30.3. Creating a Template Tensor Map Value Using the Driver API
  • 7.31. Profiler Counter Function
  • 7.32. Assertion
  • 7.33. Trap function
  • 7.34. Breakpoint Function
  • 7.35. Formatted Output
    • 7.35.1. Format Specifiers
    • 7.35.2. Limitations
    • 7.35.3. Associated Host-Side API
    • 7.35.4. Examples
  • 7.36. Dynamic Global Memory Allocation and Operations
    • 7.36.1. Heap Memory Allocation
    • 7.36.2. Interoperability with Host Memory API
    • 7.36.3. Examples
      • 7.36.3.1. Per Thread Allocation
      • 7.36.3.2. Per Thread Block Allocation
      • 7.36.3.3. Allocation Persisting Between Kernel Launches
  • 7.37. Execution Configuration
  • 7.38. Launch Bounds
  • 7.39. Maximum Number of Registers per Thread
  • 7.40. #pragma unroll
  • 7.41. SIMD Video Instructions
  • 7.42. Diagnostic Pragmas
• 8. Cooperative Groups
  • 8.1. Introduction
  • 8.2. What’s New in Cooperative Groups
    • 8.2.1. CUDA 12.2
    • 8.2.2. CUDA 12.1
    • 8.2.3. CUDA 12.0
  • 8.3. Programming Model Concept
    • 8.3.1. Composition Example
  • 8.4. Group Types
    • 8.4.1. Implicit Groups
      • 8.4.1.1. Thread Block Group
      • 8.4.1.2. Cluster Group
      • 8.4.1.3. Grid Group
      • 8.4.1.4. Multi Grid Group
    • 8.4.2. Explicit Groups
      • 8.4.2.1. Thread Block Tile
        • 8.4.2.1.1. Warp-Synchronous Code Pattern
        • 8.4.2.1.2. Single thread group
      • 8.4.2.2. Coalesced Groups
        • 8.4.2.2.1. Discovery Pattern
  • 8.5. Group Partitioning
    • 8.5.1. tiled_partition
    • 8.5.2. labeled_partition
    • 8.5.3. binary_partition
  • 8.6. Group Collectives
    • 8.6.1. Synchronization
      • 8.6.1.1. barrier_arrive and barrier_wait
      • 8.6.1.2. sync
    • 8.6.2. Data Transfer
      • 8.6.2.1. memcpy_async
      • 8.6.2.2. wait and wait_prior
    • 8.6.3. Data Manipulation
      • 8.6.3.1. reduce
      • 8.6.3.2. Reduce Operators
      • 8.6.3.3. inclusive_scan and exclusive_scan
    • 8.6.4. Execution control
      • 8.6.4.1. invoke_one and invoke_one_broadcast
  • 8.7. Grid Synchronization
  • 8.8. Multi-Device Synchronization
• 9. CUDA Dynamic Parallelism
  • 9.1. Introduction
    • 9.1.1. Overview
    • 9.1.2. Glossary
  • 9.2. Execution Environment and Memory Model
    • 9.2.1. Execution Environment
      • 9.2.1.1. Parent and Child Grids
      • 9.2.1.2. Scope of CUDA Primitives
      • 9.2.1.3. Synchronization
      • 9.2.1.4. Streams and Events
      • 9.2.1.5. Ordering and Concurrency
      • 9.2.1.6. Device Management
    • 9.2.2. Memory Model
      • 9.2.2.1. Coherence and Consistency
        • 9.2.2.1.1. Global Memory
        • 9.2.2.1.2. Zero Copy Memory
        • 9.2.2.1.3. Constant Memory
        • 9.2.2.1.4. Shared and Local Memory
        • 9.2.2.1.5. Local Memory
        • 9.2.2.1.6. Texture Memory
  • 9.3. Programming Interface
    • 9.3.1. CUDA C++ Reference
      • 9.3.1.1. Device-Side Kernel Launch
        • 9.3.1.1.1. Launches are Asynchronous
        • 9.3.1.1.2. Launch Environment Configuration
      • 9.3.1.2. Streams
        • 9.3.1.2.1. The Implicit (NULL) Stream
        • 9.3.1.2.2. The Fire-and-Forget Stream
        • 9.3.1.2.3. The Tail Launch Stream
      • 9.3.1.3. Events
      • 9.3.1.4. Synchronization
      • 9.3.1.5. Device Management
      • 9.3.1.6. Memory Declarations
        • 9.3.1.6.1. Device and Constant Memory
        • 9.3.1.6.2. Textures and Surfaces
        • 9.3.1.6.3. Shared Memory Variable Declarations
        • 9.3.1.6.4. Symbol Addresses
      • 9.3.1.7. API Errors and Launch Failures
        • 9.3.1.7.1. Launch Setup APIs
      • 9.3.1.8. API Reference
    • 9.3.2. Device-side Launch from PTX
      • 9.3.2.1. Kernel Launch APIs
        • 9.3.2.1.1. cudaLaunchDevice
        • 9.3.2.1.2. cudaGetParameterBuffer
      • 9.3.2.2. Parameter Buffer Layout
    • 9.3.3. Toolkit Support for Dynamic Parallelism
      • 9.3.3.1. Including Device Runtime API in CUDA Code
      • 9.3.3.2. Compiling and Linking
  • 9.4. Programming Guidelines
    • 9.4.1. Basics
    • 9.4.2. Performance
      • 9.4.2.1. Dynamic-parallelism-enabled Kernel Overhead
    • 9.4.3. Implementation Restrictions and Limitations
      • 9.4.3.1. Runtime
        • 9.4.3.1.1. Memory Footprint
        • 9.4.3.1.2. Pending Kernel Launches
        • 9.4.3.1.3. Configuration Options
        • 9.4.3.1.4. Memory Allocation and Lifetime
        • 9.4.3.1.5. SM Id and Warp Id
        • 9.4.3.1.6. ECC Errors
  • 9.5. CDP2 vs CDP1
    • 9.5.1. Differences Between CDP1 and CDP2
    • 9.5.2. Compatibility and Interoperability
  • 9.6. Legacy CUDA Dynamic Parallelism (CDP1)
    • 9.6.1. Execution Environment and Memory Model (CDP1)
      • 9.6.1.1. Execution Environment (CDP1)
        • 9.6.1.1.1. Parent and Child Grids (CDP1)
        • 9.6.1.1.2. Scope of CUDA Primitives (CDP1)
        • 9.6.1.1.3. Synchronization (CDP1)
        • 9.6.1.1.4. Streams and Events (CDP1)
        • 9.6.1.1.5. Ordering and Concurrency (CDP1)
        • 9.6.1.1.6. Device Management (CDP1)
      • 9.6.1.2. Memory Model (CDP1)
        • 9.6.1.2.1. Coherence and Consistency (CDP1)
          • 9.6.1.2.1.1. Global Memory (CDP1)
          • 9.6.1.2.1.2. Zero Copy Memory (CDP1)
          • 9.6.1.2.1.3. Constant Memory (CDP1)
          • 9.6.1.2.1.4. Shared and Local Memory (CDP1)
          • 9.6.1.2.1.5. Local Memory (CDP1)
          • 9.6.1.2.1.6. Texture Memory (CDP1)
    • 9.6.2. Programming Interface (CDP1)
      • 9.6.2.1. CUDA C++ Reference (CDP1)
        • 9.6.2.1.1. Device-Side Kernel Launch (CDP1)
          • 9.6.2.1.1.1. Launches are Asynchronous (CDP1)
          • 9.6.2.1.1.2. Launch Environment Configuration (CDP1)
        • 9.6.2.1.2. Streams (CDP1)
          • 9.6.2.1.2.1. The Implicit (NULL) Stream (CDP1)
        • 9.6.2.1.3. Events (CDP1)
        • 9.6.2.1.4. Synchronization (CDP1)
          • 9.6.2.1.4.1. Block Wide Synchronization (CDP1)
        • 9.6.2.1.5. Device Management (CDP1)
        • 9.6.2.1.6. Memory Declarations (CDP1)
          • 9.6.2.1.6.1. Device and Constant Memory (CDP1)
          • 9.6.2.1.6.2. Textures and Surfaces (CDP1)
          • 9.6.2.1.6.3. Shared Memory Variable Declarations (CDP1)
          • 9.6.2.1.6.4. Symbol Addresses (CDP1)
        • 9.6.2.1.7. API Errors and Launch Failures (CDP1)
          • 9.6.2.1.7.1. Launch Setup APIs (CDP1)
        • 9.6.2.1.8. API Reference (CDP1)
      • 9.6.2.2. Device-side Launch from PTX (CDP1)
        • 9.6.2.2.1. Kernel Launch APIs (CDP1)
          • 9.6.2.2.1.1. cudaLaunchDevice (CDP1)
          • 9.6.2.2.1.2. cudaGetParameterBuffer (CDP1)
        • 9.6.2.2.2. Parameter Buffer Layout (CDP1)
      • 9.6.2.3. Toolkit Support for Dynamic Parallelism (CDP1)
        • 9.6.2.3.1. Including Device Runtime API in CUDA Code (CDP1)
        • 9.6.2.3.2. Compiling and Linking (CDP1)
    • 9.6.3. Programming Guidelines (CDP1)
      • 9.6.3.1. Basics (CDP1)
      • 9.6.3.2. Performance (CDP1)
        • 9.6.3.2.1. Synchronization (CDP1)
        • 9.6.3.2.2. Dynamic-parallelism-enabled Kernel Overhead (CDP1)
      • 9.6.3.3. Implementation Restrictions and Limitations (CDP1)
        • 9.6.3.3.1. Runtime (CDP1)
          • 9.6.3.3.1.1. Memory Footprint (CDP1)
          • 9.6.3.3.1.2. Nesting and Synchronization Depth (CDP1)
          • 9.6.3.3.1.3. Pending Kernel Launches (CDP1)
          • 9.6.3.3.1.4. Configuration Options (CDP1)
          • 9.6.3.3.1.5. Memory Allocation and Lifetime (CDP1)
          • 9.6.3.3.1.6. SM Id and Warp Id (CDP1)
          • 9.6.3.3.1.7. ECC Errors (CDP1)
• 10. Virtual Memory Management
  • 10.1. Introduction
  • 10.2. Query for Support
  • 10.3. Allocating Physical Memory
    • 10.3.1. Shareable Memory Allocations
    • 10.3.2. Memory Type
      • 10.3.2.1. Compressible Memory
  • 10.4. Reserving a Virtual Address Range
  • 10.5. Virtual Aliasing Support
  • 10.6. Mapping Memory
  • 10.7. Controlling Access Rights
• 11. Stream Ordered Memory Allocator
  • 11.1. Introduction
  • 11.2. Query for Support
  • 11.3. API Fundamentals (cudaMallocAsync and cudaFreeAsync)
  • 11.4. Memory Pools and the cudaMemPool_t
  • 11.5. Default/Implicit Pools
  • 11.6. Explicit Pools
  • 11.7. Physical Page Caching Behavior
  • 11.8. Resource Usage Statistics
  • 11.9. Memory Reuse Policies
    • 11.9.1. cudaMemPoolReuseFollowEventDependencies
    • 11.9.2. cudaMemPoolReuseAllowOpportunistic
    • 11.9.3. cudaMemPoolReuseAllowInternalDependencies
    • 11.9.4. Disabling Reuse Policies
  • 11.10. Device Accessibility for Multi-GPU Support
  • 11.11. IPC Memory Pools
    • 11.11.1. Creating and Sharing IPC Memory Pools
    • 11.11.2. Set Access in the Importing Process
    • 11.11.3. Creating and Sharing Allocations from an Exported Pool
    • 11.11.4. IPC Export Pool Limitations
    • 11.11.5. IPC Import Pool Limitations
  • 11.12. Synchronization API Actions
  • 11.13. Addendums
    • 11.13.1. cudaMemcpyAsync Current Context/Device Sensitivity
    • 11.13.2. cuPointerGetAttribute Query
    • 11.13.3. cuGraphAddMemsetNode
    • 11.13.4. Pointer Attributes
• 12. Graph Memory Nodes
  • 12.1. Introduction
  • 12.2. Support and Compatibility
  • 12.3. API Fundamentals
    • 12.3.1. Graph Node APIs
    • 12.3.2. Stream Capture
    • 12.3.3. Accessing and Freeing Graph Memory Outside of the Allocating Graph
    • 12.3.4. cudaGraphInstantiateFlagAutoFreeOnLaunch
  • 12.4. Optimized Memory Reuse
    • 12.4.1. Address Reuse within a Graph
    • 12.4.2. Physical Memory Management and Sharing
  • 12.5. Performance Considerations
    • 12.5.1. First Launch / cudaGraphUpload
  • 12.6. Physical Memory Footprint
  • 12.7. Peer Access
    • 12.7.1. Peer Access with Graph Node APIs
    • 12.7.2. Peer Access with Stream Capture
• 13. Mathematical Functions
  • 13.1. Standard Functions
  • 13.2. Intrinsic Functions
• 14. C++ Language Support
  • 14.1. C++11 Language Features
  • 14.2. C++14 Language Features
  • 14.3. C++17 Language Features
  • 14.4. C++20 Language Features
  • 14.5. Restrictions
    • 14.5.1. Host Compiler Extensions
    • 14.5.2. Preprocessor Symbols
      • 14.5.2.1. __CUDA_ARCH__
    • 14.5.3. Qualifiers
      • 14.5.3.1. Device Memory Space Specifiers
      • 14.5.3.2. __managed__ Memory Space Specifier
      • 14.5.3.3. Volatile Qualifier
    • 14.5.4. Pointers
    • 14.5.5. Operators
      • 14.5.5.1. Assignment Operator
      • 14.5.5.2. Address Operator
    • 14.5.6. Run Time Type Information (RTTI)
    • 14.5.7. Exception Handling
    • 14.5.8. Standard Library
    • 14.5.9. Namespace Reservations
    • 14.5.10. Functions
      • 14.5.10.1. External Linkage
      • 14.5.10.2. Implicitly-declared and explicitly-defaulted functions
      • 14.5.10.3. Function Parameters
        • 14.5.10.3.1. __global__ Function Argument Processing
        • 14.5.10.3.2. Toolkit and Driver Compatibility
        • 14.5.10.3.3. Link Compatibility across Toolkit Revisions
      • 14.5.10.4. Static Variables within Function
      • 14.5.10.5. Function Pointers
      • 14.5.10.6. Function Recursion
      • 14.5.10.7. Friend Functions
      • 14.5.10.8. Operator Function
      • 14.5.10.9. Allocation and Deallocation Functions
    • 14.5.11. Classes
      • 14.5.11.1. Data Members
      • 14.5.11.2. Function Members
      • 14.5.11.3. Virtual Functions
      • 14.5.11.4. Virtual Base Classes
      • 14.5.11.5. Anonymous Unions
      • 14.5.11.6. Windows-Specific
    • 14.5.12. Templates
    • 14.5.13. Trigraphs and Digraphs
    • 14.5.14. Const-qualified variables
    • 14.5.15. Long Double
    • 14.5.16. Deprecation Annotation
    • 14.5.17. Noreturn Annotation
    • 14.5.18. [[likely]] / [[unlikely]] Standard Attributes
    • 14.5.19. const and pure GNU Attributes
    • 14.5.20. __nv_pure__ Attribute
    • 14.5.21. Intel Host Compiler Specific
    • 14.5.22. C++11 Features
      • 14.5.22.1. Lambda Expressions
      • 14.5.22.2. std::initializer_list
      • 14.5.22.3. Rvalue references
      • 14.5.22.4. Constexpr functions and function templates
      • 14.5.22.5. Constexpr variables
      • 14.5.22.6. Inline namespaces
        • 14.5.22.6.1. Inline unnamed namespaces
      • 14.5.22.7. thread_local
      • 14.5.22.8. __global__ functions and function templates
      • 14.5.22.9. __managed__ and __shared__ variables
      • 14.5.22.10. Defaulted functions
    • 14.5.23. C++14 Features
      • 14.5.23.1. Functions with deduced return type
      • 14.5.23.2. Variable templates
    • 14.5.24. C++17 Features
      • 14.5.24.1. Inline Variable
      • 14.5.24.2. Structured Binding
    • 14.5.25. C++20 Features
      • 14.5.25.1. Module support
      • 14.5.25.2. Coroutine support
      • 14.5.25.3. Three-way comparison operator
      • 14.5.25.4. Consteval functions
  • 14.6. Polymorphic Function Wrappers
  • 14.7. Extended Lambdas
    • 14.7.1. Extended Lambda Type Traits
    • 14.7.2. Extended Lambda Restrictions
    • 14.7.3. Notes on __host__ __device__ lambdas
    • 14.7.4. *this Capture By Value
    • 14.7.5. Additional Notes
  • 14.8. Code Samples
    • 14.8.1. Data Aggregation Class
    • 14.8.2. Derived Class
    • 14.8.3. Class Template
    • 14.8.4. Function Template
    • 14.8.5. Functor Class
• 15. Texture Fetching
  • 15.1. Nearest-Point Sampling
  • 15.2. Linear Filtering
  • 15.3. Table Lookup
• 16. Compute Capabilities
  • 16.1. Feature Availability
  • 16.2. Features and Technical Specifications
  • 16.3. Floating-Point Standard
  • 16.4. Compute Capability 5.x
    • 16.4.1. Architecture
    • 16.4.2. Global Memory
    • 16.4.3. Shared Memory
  • 16.5. Compute Capability 6.x
    • 16.5.1. Architecture
    • 16.5.2. Global Memory
    • 16.5.3. Shared Memory
  • 16.6. Compute Capability 7.x
    • 16.6.1. Architecture
    • 16.6.2. Independent Thread Scheduling
    • 16.6.3. Global Memory
    • 16.6.4. Shared Memory
  • 16.7. Compute Capability 8.x
    • 16.7.1. Architecture
    • 16.7.2. Global Memory
    • 16.7.3. Shared Memory
  • 16.8. Compute Capability 9.0
    • 16.8.1. Architecture
    • 16.8.2. Global Memory
    • 16.8.3. Shared Memory
    • 16.8.4. Features Accelerating Specialized Computations
• 17. Driver API
  • 17.1. Context
  • 17.2. Module
  • 17.3. Kernel Execution
  • 17.4. Interoperability between Runtime and Driver APIs
  • 17.5. Driver Entry Point Access
    • 17.5.1. Introduction
    • 17.5.2. Driver Function Typedefs
    • 17.5.3. Driver Function Retrieval
      • 17.5.3.1. Using the driver API
      • 17.5.3.2. Using the runtime API
      • 17.5.3.3. Retrieve per-thread default stream versions
      • 17.5.3.4. Access new CUDA features
    • 17.5.4. Potential Implications with cuGetProcAddress
      • 17.5.4.1. Implications with cuGetProcAddress vs Implicit Linking
      • 17.5.4.2. Compile Time vs Runtime Version Usage in cuGetProcAddress
      • 17.5.4.3. API Version Bumps with Explicit Version Checks
      • 17.5.4.4. Issues with Runtime API Usage
      • 17.5.4.5. Issues with Runtime API and Dynamic Versioning
      • 17.5.4.6. Issues with Runtime API allowing CUDA Version
      • 17.5.4.7. Implications to API/ABI
    • 17.5.5. Determining cuGetProcAddress Failure Reasons
• 18. CUDA Environment Variables
• 19. Unified Memory Programming
  • 19.1. Unified Memory Introduction
    • 19.1.1. System Requirements for Unified Memory
    • 19.1.2. Programming Model
      • 19.1.2.1. Allocation APIs for System-Allocated Memory
      • 19.1.2.2. Allocation API for CUDA Managed Memory: cudaMallocManaged()
      • 19.1.2.3. Global-Scope Managed Variables Using __managed__
      • 19.1.2.4. Difference between Unified Memory and Mapped Memory
      • 19.1.2.5. Pointer Attributes
      • 19.1.2.6. Runtime detection of Unified Memory Support Level
      • 19.1.2.7. GPU Memory Oversubscription
      • 19.1.2.8. Performance Hints
        • 19.1.2.8.1. Data Prefetching
        • 19.1.2.8.2. Data Usage Hints
        • 19.1.2.8.3. Querying Data Usage Attributes on Managed Memory
  • 19.2. Unified memory on devices with full CUDA Unified Memory support
    • 19.2.1. System-Allocated Memory: in-depth examples
      • 19.2.1.1. File-backed Unified Memory
      • 19.2.1.2. Inter-Process Communication (IPC) with Unified Memory
    • 19.2.2. Performance Tuning
      • 19.2.2.1. Memory Paging and Page Sizes
        • 19.2.2.1.1. Choosing the right page size
        • 19.2.2.1.2. CPU and GPU page tables: hardware coherency vs. software coherency
      • 19.2.2.2. Direct Unified Memory Access from host
      • 19.2.2.3. Host Native Atomics
      • 19.2.2.4. Atomic accesses & synchronization primitives
      • 19.2.2.5. Memcpy()/Memset() Behavior With Unified Memory
  • 19.3. Unified memory on devices without full CUDA Unified Memory support
    • 19.3.1. Unified memory on devices with only CUDA Managed Memory support
    • 19.3.2. Unified memory on Windows or devices with compute capability 5.x
      • 19.3.2.1. Data Migration and Coherency
      • 19.3.2.2. GPU Memory Oversubscription
      • 19.3.2.3. Multi-GPU
      • 19.3.2.4. Coherency and Concurrency
        • 19.3.2.4.1. GPU Exclusive Access To Managed Memory
        • 19.3.2.4.2. Explicit Synchronization and Logical GPU Activity
        • 19.3.2.4.3. Managing Data Visibility and Concurrent CPU + GPU Access with Streams
        • 19.3.2.4.4. Stream Association Examples
        • 19.3.2.4.5. Stream Attach With Multithreaded Host Programs
        • 19.3.2.4.6. Advanced Topic: Modular Programs and Data Access Constraints
        • 19.3.2.4.7. Memcpy()/Memset() Behavior With Stream-associated Unified Memory
• 20. Lazy Loading
  • 20.1. What is Lazy Loading?
  • 20.2. Lazy Loading version support
    • 20.2.1. Driver
    • 20.2.2. Toolkit
    • 20.2.3. Compiler
  • 20.3. Triggering loading of kernels in lazy mode
    • 20.3.1. CUDA Driver API
    • 20.3.2. CUDA Runtime API
  • 20.4. Querying whether Lazy Loading is turned on
  • 20.5. Possible issues when adopting lazy loading
    • 20.5.1. Concurrent execution
    • 20.5.2. Allocators
    • 20.5.3. Autotuning
• 21. Notices
  • 21.1. Notice
  • 21.2. OpenCL
  • 21.3. Trademarks