CUDA C++ Programming Guide (NVIDIA Corporation)（Z-Library）

Page 1

CUDA C++ Programming Guide Release 12.8 NVIDIA Corporation Feb 28, 2025

Page 2

(This page has no text content)

Page 3

Contents 1 The Benefits of Using GPUs 3 2 CUDA®: A General-Purpose Parallel Computing Platform and Programming Model 5 3 A Scalable Programming Model 7 4 Document Structure 9 5 Programming Model 11 5.1 Kernels . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 11 5.2 Thread Hierarchy . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 12 5.2.1 Thread Block Clusters . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 14 5.3 Memory Hierarchy . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 16 5.4 Heterogeneous Programming . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 17 5.5 Asynchronous SIMT Programming Model . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 17 5.5.1 Asynchronous Operations . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 17 5.6 Compute Capability . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 19 6 Programming Interface 21 6.1 Compilation with NVCC . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 21 6.1.1 Compilation Workflow . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 22 6.1.1.1 Offline Compilation . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 22 6.1.1.2 Just-in-Time Compilation . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 22 6.1.2 Binary Compatibility . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 22 6.1.3 PTX Compatibility . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 23 6.1.4 Application Compatibility . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 23 6.1.5 C++ Compatibility . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 24 6.1.6 64-Bit Compatibility . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 24 6.2 CUDA Runtime . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 24 6.2.1 Initialization . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 25 6.2.2 Device Memory . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 26 6.2.3 Device Memory L2 Access Management . . . . . . . . . . . . . . . . . . . . . . . . . . . . 29 6.2.3.1 L2 Cache Set-Aside for Persisting Accesses . . . . . . . . . . . . . . . . . . . . . . . 29 6.2.3.2 L2 Policy for Persisting Accesses . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 30 6.2.3.3 L2 Access Properties . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 31 6.2.3.4 L2 Persistence Example . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 31 6.2.3.5 Reset L2 Access to Normal . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 32 6.2.3.6 Manage Utilization of L2 set-aside cache . . . . . . . . . . . . . . . . . . . . . . . . 33 6.2.3.7 Query L2 cache Properties . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 33 6.2.3.8 Control L2 Cache Set-Aside Size for Persisting Memory Access . . . . . . . . . . 33 6.2.4 Shared Memory . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 34 6.2.5 Distributed Shared Memory . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 40 6.2.6 Page-Locked Host Memory . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 42 6.2.6.1 Portable Memory . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 43 i

Page 4

6.2.6.2 Write-Combining Memory . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 43 6.2.6.3 Mapped Memory . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 43 6.2.7 Memory Synchronization Domains . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 44 6.2.7.1 Memory Fence Interference . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 44 6.2.7.2 Isolating Traffic with Domains . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 45 6.2.7.3 Using Domains in CUDA . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 45 6.2.8 Asynchronous Concurrent Execution . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 46 6.2.8.1 Concurrent Execution between Host and Device . . . . . . . . . . . . . . . . . . . . 47 6.2.8.2 Concurrent Kernel Execution . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 47 6.2.8.3 Overlap of Data Transfer and Kernel Execution . . . . . . . . . . . . . . . . . . . . . 47 6.2.8.4 Concurrent Data Transfers . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 48 6.2.8.5 Streams . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 48 6.2.8.6 Programmatic Dependent Launch and Synchronization . . . . . . . . . . . . . . . 51 6.2.8.7 CUDA Graphs . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 54 6.2.8.8 Events . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 78 6.2.8.9 Synchronous Calls . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 78 6.2.9 Multi-Device System . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 79 6.2.9.1 Device Enumeration . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 79 6.2.9.2 Device Selection . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 79 6.2.9.3 Stream and Event Behavior . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 79 6.2.9.4 Peer-to-Peer Memory Access . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 80 6.2.9.5 Peer-to-Peer Memory Copy . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 81 6.2.10 Unified Virtual Address Space . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 81 6.2.11 Interprocess Communication . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 82 6.2.12 Error Checking . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 82 6.2.13 Call Stack . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 83 6.2.14 Texture and Surface Memory . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 83 6.2.14.1 Texture Memory . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 83 6.2.14.2 Surface Memory . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 89 6.2.14.3 CUDA Arrays . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 91 6.2.14.4 Read/Write Coherency . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 92 6.2.15 Graphics Interoperability . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 92 6.2.15.1 OpenGL Interoperability . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 92 6.2.15.2 Direct3D Interoperability . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 95 6.2.15.3 SLI Interoperability . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 101 6.2.16 External Resource Interoperability . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 102 6.2.16.1 Vulkan Interoperability . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 102 6.2.16.2 OpenGL Interoperability . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 110 6.2.16.3 Direct3D 12 Interoperability . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 111 6.2.16.4 Direct3D 11 Interoperability . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 117 6.2.16.5 NVIDIA Software Communication Interface Interoperability (NVSCI) . . . . . . . . 125 6.3 Versioning and Compatibility . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 131 6.4 Compute Modes . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 132 6.5 Mode Switches . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 133 6.6 Tesla Compute Cluster Mode for Windows . . . . . . . . . . . . . . . . . . . . . . . . . . . . 133 7 Hardware Implementation 135 7.1 SIMT Architecture . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 135 7.2 Hardware Multithreading . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 137 8 Performance Guidelines 139 8.1 Overall Performance Optimization Strategies . . . . . . . . . . . . . . . . . . . . . . . . . . . 139 8.2 Maximize Utilization . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 139 8.2.1 Application Level . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 140 ii

Page 5

8.2.2 Device Level . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 140 8.2.3 Multiprocessor Level . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 140 8.2.3.1 Occupancy Calculator . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 142 8.3 Maximize Memory Throughput . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 144 8.3.1 Data Transfer between Host and Device . . . . . . . . . . . . . . . . . . . . . . . . . . . . 145 8.3.2 Device Memory Accesses . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 146 8.4 Maximize Instruction Throughput . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 149 8.4.1 Arithmetic Instructions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 149 8.4.2 Control Flow Instructions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 156 8.4.3 Synchronization Instruction . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 156 8.5 Minimize Memory Thrashing . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 156 9 CUDA-Enabled GPUs 159 10 C++ Language Extensions 161 10.1 Function Execution Space Specifiers . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 161 10.1.1 __global__ . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 161 10.1.2 __device__ . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 161 10.1.3 __host__ . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 162 10.1.4 Undefined behavior . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 162 10.1.5 __noinline__ and __forceinline__ . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 162 10.1.6 __inline_hint__ . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 163 10.2 Variable Memory Space Specifiers . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 163 10.2.1 __device__ . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 163 10.2.2 __constant__ . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 163 10.2.3 __shared__ . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 164 10.2.4 __grid_constant__ . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 165 10.2.5 __managed__ . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 165 10.2.6 __restrict__ . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 166 10.3 Built-in Vector Types . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 167 10.3.1 char, short, int, long, longlong, float, double . . . . . . . . . . . . . . . . . . . . . . . . . . 167 10.3.2 dim3 . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 169 10.4 Built-in Variables . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 169 10.4.1 gridDim . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 169 10.4.2 blockIdx . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 169 10.4.3 blockDim . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 169 10.4.4 threadIdx . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 169 10.4.5 warpSize . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 169 10.5 Memory Fence Functions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 170 10.6 Synchronization Functions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 173 10.7 Mathematical Functions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 174 10.8 Texture Functions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 174 10.8.1 Texture Object API . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 174 10.8.1.1 tex1Dfetch() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 174 10.8.1.2 tex1D() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 174 10.8.1.3 tex1DLod() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 174 10.8.1.4 tex1DGrad() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 175 10.8.1.5 tex2D() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 175 10.8.1.6 tex2D() for sparse CUDA arrays . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 175 10.8.1.7 tex2Dgather() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 175 10.8.1.8 tex2Dgather() for sparse CUDA arrays . . . . . . . . . . . . . . . . . . . . . . . . . . 175 10.8.1.9 tex2DGrad() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 176 10.8.1.10 tex2DGrad() for sparse CUDA arrays . . . . . . . . . . . . . . . . . . . . . . . . . . . 176 10.8.1.11 tex2DLod() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 176 iii

Page 6

10.8.1.12 tex2DLod() for sparse CUDA arrays . . . . . . . . . . . . . . . . . . . . . . . . . . . . 176 10.8.1.13 tex3D() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 176 10.8.1.14 tex3D() for sparse CUDA arrays . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 177 10.8.1.15 tex3DLod() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 177 10.8.1.16 tex3DLod() for sparse CUDA arrays . . . . . . . . . . . . . . . . . . . . . . . . . . . . 177 10.8.1.17 tex3DGrad() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 177 10.8.1.18 tex3DGrad() for sparse CUDA arrays . . . . . . . . . . . . . . . . . . . . . . . . . . . 177 10.8.1.19 tex1DLayered() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 178 10.8.1.20 tex1DLayeredLod() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 178 10.8.1.21 tex1DLayeredGrad() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 178 10.8.1.22 tex2DLayered() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 178 10.8.1.23 tex2DLayered() for Sparse CUDA Arrays . . . . . . . . . . . . . . . . . . . . . . . . . 178 10.8.1.24 tex2DLayeredLod() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 179 10.8.1.25 tex2DLayeredLod() for sparse CUDA arrays . . . . . . . . . . . . . . . . . . . . . . . 179 10.8.1.26 tex2DLayeredGrad() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 179 10.8.1.27 tex2DLayeredGrad() for sparse CUDA arrays . . . . . . . . . . . . . . . . . . . . . . 179 10.8.1.28 texCubemap() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 179 10.8.1.29 texCubemapGrad() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 180 10.8.1.30 texCubemapLod() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 180 10.8.1.31 texCubemapLayered() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 180 10.8.1.32 texCubemapLayeredGrad() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 180 10.8.1.33 texCubemapLayeredLod() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 180 10.9 Surface Functions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 181 10.9.1 Surface Object API . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 181 10.9.1.1 surf1Dread() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 181 10.9.1.2 surf1Dwrite . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 181 10.9.1.3 surf2Dread() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 181 10.9.1.4 surf2Dwrite() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 182 10.9.1.5 surf3Dread() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 182 10.9.1.6 surf3Dwrite() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 182 10.9.1.7 surf1DLayeredread() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 182 10.9.1.8 surf1DLayeredwrite() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 183 10.9.1.9 surf2DLayeredread() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 183 10.9.1.10 surf2DLayeredwrite() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 183 10.9.1.11 surfCubemapread() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 183 10.9.1.12 surfCubemapwrite() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 184 10.9.1.13 surfCubemapLayeredread() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 184 10.9.1.14 surfCubemapLayeredwrite() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 184 10.10 Read-Only Data Cache Load Function . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 184 10.11 Load Functions Using Cache Hints . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 185 10.12 Store Functions Using Cache Hints . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 185 10.13 Time Function . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 185 10.14 Atomic Functions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 186 10.14.1 Arithmetic Functions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 188 10.14.1.1 atomicAdd() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 188 10.14.1.2 atomicSub() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 189 10.14.1.3 atomicExch() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 189 10.14.1.4 atomicMin() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 190 10.14.1.5 atomicMax() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 190 10.14.1.6 atomicInc() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 191 10.14.1.7 atomicDec() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 191 10.14.1.8 atomicCAS() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 191 10.14.1.9 __nv_atomic_exchange() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 192 10.14.1.10__nv_atomic_exchange_n() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 192 iv

Page 7

10.14.1.11__nv_atomic_compare_exchange() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 192 10.14.1.12__nv_atomic_compare_exchange_n() . . . . . . . . . . . . . . . . . . . . . . . . . . . 193 10.14.1.13__nv_atomic_fetch_add() and __nv_atomic_add() . . . . . . . . . . . . . . . . . . . 193 10.14.1.14__nv_atomic_fetch_sub() and __nv_atomic_sub() . . . . . . . . . . . . . . . . . . . . 194 10.14.1.15__nv_atomic_fetch_min() and __nv_atomic_min() . . . . . . . . . . . . . . . . . . . 194 10.14.1.16__nv_atomic_fetch_max() and __nv_atomic_max() . . . . . . . . . . . . . . . . . . . 194 10.14.2 Bitwise Functions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 195 10.14.2.1 atomicAnd() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 195 10.14.2.2 atomicOr() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 195 10.14.2.3 atomicXor() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 195 10.14.2.4 __nv_atomic_fetch_or() and __nv_atomic_or() . . . . . . . . . . . . . . . . . . . . . . 196 10.14.2.5 __nv_atomic_fetch_xor() and __nv_atomic_xor() . . . . . . . . . . . . . . . . . . . . 196 10.14.2.6 __nv_atomic_fetch_and() and __nv_atomic_and() . . . . . . . . . . . . . . . . . . . 196 10.14.3 Other atomic functions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 197 10.14.3.1 __nv_atomic_load() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 197 10.14.3.2 __nv_atomic_load_n() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 197 10.14.3.3 __nv_atomic_store() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 198 10.14.3.4 __nv_atomic_store_n() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 198 10.14.3.5 __nv_atomic_thread_fence() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 198 10.15 Address Space Predicate Functions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 199 10.15.1 __isGlobal() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 199 10.15.2 __isShared() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 199 10.15.3 __isConstant() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 199 10.15.4 __isGridConstant() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 199 10.15.5 __isLocal() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 199 10.16 Address Space Conversion Functions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 200 10.16.1 __cvta_generic_to_global() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 200 10.16.2 __cvta_generic_to_shared() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 200 10.16.3 __cvta_generic_to_constant() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 200 10.16.4 __cvta_generic_to_local() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 200 10.16.5 __cvta_global_to_generic() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 200 10.16.6 __cvta_shared_to_generic() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 201 10.16.7 __cvta_constant_to_generic() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 201 10.16.8 __cvta_local_to_generic() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 201 10.17 Alloca Function . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 201 10.17.1 Synopsis . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 201 10.17.2 Description . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 201 10.17.3 Example . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 202 10.18 Compiler Optimization Hint Functions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 202 10.18.1 __builtin_assume_aligned() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 202 10.18.2 __builtin_assume() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 202 10.18.3 __assume() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 203 10.18.4 __builtin_expect() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 203 10.18.5 __builtin_unreachable() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 203 10.18.6 Restrictions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 204 10.19 Warp Vote Functions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 204 10.20 Warp Match Functions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 205 10.20.1 Synopsis . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 205 10.20.2 Description . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 205 10.21 Warp Reduce Functions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 206 10.21.1 Synopsis . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 206 10.21.2 Description . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 206 10.22 Warp Shuffle Functions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 207 10.22.1 Synopsis . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 207 v

Page 8

10.22.2 Description . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 207 10.22.3 Examples . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 208 10.22.3.1 Broadcast of a single value across a warp . . . . . . . . . . . . . . . . . . . . . . . . 208 10.22.3.2 Inclusive plus-scan across sub-partitions of 8 threads . . . . . . . . . . . . . . . . 209 10.22.3.3 Reduction across a warp . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 209 10.23 Nanosleep Function . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 210 10.23.1 Synopsis . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 210 10.23.2 Description . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 210 10.23.3 Example . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 210 10.24 Warp Matrix Functions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 210 10.24.1 Description . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 211 10.24.2 Alternate Floating Point . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 213 10.24.3 Double Precision . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 213 10.24.4 Sub-byte Operations . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 213 10.24.5 Restrictions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 215 10.24.6 Element Types and Matrix Sizes . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 215 10.24.7 Example . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 217 10.25 DPX . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 217 10.25.1 Examples . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 218 10.26 Asynchronous Barrier . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 219 10.26.1 Simple Synchronization Pattern . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 219 10.26.2 Temporal Splitting and Five Stages of Synchronization . . . . . . . . . . . . . . . . . . . 219 10.26.3 Bootstrap Initialization, Expected Arrival Count, and Participation . . . . . . . . . . . . 220 10.26.4 A Barrier’s Phase: Arrival, Countdown, Completion, and Reset . . . . . . . . . . . . . . 221 10.26.5 Spatial Partitioning (also known as Warp Specialization) . . . . . . . . . . . . . . . . . . 222 10.26.6 Early Exit (Dropping out of Participation) . . . . . . . . . . . . . . . . . . . . . . . . . . . 224 10.26.7 Completion Function . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 224 10.26.8 Memory Barrier Primitives Interface . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 226 10.26.8.1 Data Types . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 226 10.26.8.2 Memory Barrier Primitives API . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 226 10.27 Asynchronous Data Copies . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 227 10.27.1 memcpy_async API . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 227 10.27.2 Copy and Compute Pattern - Staging Data Through Shared Memory . . . . . . . . . . 227 10.27.3 Without memcpy_async . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 228 10.27.4 With memcpy_async . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 229 10.27.5 Asynchronous Data Copies using cuda::barrier . . . . . . . . . . . . . . . . . . . . . 230 10.27.6 Performance Guidance for memcpy_async . . . . . . . . . . . . . . . . . . . . . . . . . . 231 10.27.6.1 Alignment . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 231 10.27.6.2 Trivially copyable . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 231 10.27.6.3 Warp Entanglement - Commit . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 231 10.27.6.4 Warp Entanglement - Wait . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 232 10.27.6.5 Warp Entanglement - Arrive-On . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 232 10.27.6.6 Keep Commit and Arrive-On Operations Converged . . . . . . . . . . . . . . . . . . 233 10.28 Asynchronous Data Copies using cuda::pipeline . . . . . . . . . . . . . . . . . . . . . . 233 10.28.1 Single-Stage Asynchronous Data Copies using cuda::pipeline . . . . . . . . . . . . 233 10.28.2 Multi-Stage Asynchronous Data Copies using cuda::pipeline . . . . . . . . . . . . 235 10.28.3 Pipeline Interface . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 240 10.28.4 Pipeline Primitives Interface . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 240 10.28.4.1 memcpy_async Primitive . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 241 10.28.4.2 Commit Primitive . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 241 10.28.4.3 Wait Primitive . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 241 10.28.4.4 Arrive On Barrier Primitive . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 242 10.29 Asynchronous Data Copies using the Tensor Memory Accelerator (TMA) . . . . . . . . . 242 10.29.1 Using TMA to transfer one-dimensional arrays . . . . . . . . . . . . . . . . . . . . . . . . 243 vi

Page 9

10.29.2 Using TMA to transfer multi-dimensional arrays . . . . . . . . . . . . . . . . . . . . . . . 246 10.29.2.1 Multi-dimensional TMA PTX wrappers . . . . . . . . . . . . . . . . . . . . . . . . . . 250 10.29.3 TMA Swizzle . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 252 10.29.3.1 Example ‘Matrix Transpose’ . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 252 10.29.3.2 The Swizzle Modes . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 256 10.30 Encoding a Tensor Map on Device . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 258 10.30.1 Device-side Encoding and Modification of a Tensor Map . . . . . . . . . . . . . . . . . . 259 10.30.2 Usage of a Modified Tensor Map . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 261 10.30.3 Creating a Template Tensor Map Value Using the Driver API . . . . . . . . . . . . . . . 262 10.31 Profiler Counter Function . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 263 10.32 Assertion . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 263 10.33 Trap function . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 264 10.34 Breakpoint Function . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 265 10.35 Formatted Output . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 265 10.35.1 Format Specifiers . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 265 10.35.2 Limitations . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 266 10.35.3 Associated Host-Side API . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 267 10.35.4 Examples . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 267 10.36 Dynamic Global Memory Allocation and Operations . . . . . . . . . . . . . . . . . . . . . . . 268 10.36.1 Heap Memory Allocation . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 269 10.36.2 Interoperability with Host Memory API . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 269 10.36.3 Examples . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 269 10.36.3.1 Per Thread Allocation . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 269 10.36.3.2 Per Thread Block Allocation . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 270 10.36.3.3 Allocation Persisting Between Kernel Launches . . . . . . . . . . . . . . . . . . . . 271 10.37 Execution Configuration . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 272 10.38 Launch Bounds . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 274 10.39 Maximum Number of Registers per Thread . . . . . . . . . . . . . . . . . . . . . . . . . . . . 276 10.40 #pragma unroll . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 276 10.41 SIMD Video Instructions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 277 10.42 Diagnostic Pragmas . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 278 10.43 Custom ABI Pragmas . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 279 11 Cooperative Groups 281 11.1 Introduction . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 281 11.2 What’s New in Cooperative Groups . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 281 11.2.1 CUDA 12.2 . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 281 11.2.2 CUDA 12.1 . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 282 11.2.3 CUDA 12.0 . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 282 11.3 Programming Model Concept . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 282 11.3.1 Composition Example . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 283 11.4 Group Types . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 284 11.4.1 Implicit Groups . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 284 11.4.1.1 Thread Block Group . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 284 11.4.1.2 Cluster Group . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 285 11.4.1.3 Grid Group . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 286 11.4.1.4 Multi Grid Group . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 287 11.4.2 Explicit Groups . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 287 11.4.2.1 Thread Block Tile . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 287 11.4.2.2 Coalesced Groups . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 290 11.5 Group Partitioning . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 292 11.5.1 tiled_partition . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 292 11.5.2 labeled_partition . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 293 11.5.3 binary_partition . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 293 vii

Page 10

11.6 Group Collectives . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 294 11.6.1 Synchronization . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 294 11.6.1.1 barrier_arrive and barrier_wait . . . . . . . . . . . . . . . . . . . . . . . . . . 294 11.6.1.2 sync . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 295 11.6.2 Data Transfer . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 295 11.6.2.1 memcpy_async . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 295 11.6.2.2 wait and wait_prior . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 297 11.6.3 Data Manipulation . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 298 11.6.3.1 reduce . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 298 11.6.3.2 Reduce Operators . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 300 11.6.3.3 inclusive_scan and exclusive_scan . . . . . . . . . . . . . . . . . . . . . . . . 302 11.6.4 Execution control . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 305 11.6.4.1 invoke_one and invoke_one_broadcast . . . . . . . . . . . . . . . . . . . . . . . 305 11.7 Grid Synchronization . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 306 11.8 Multi-Device Synchronization . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 307 12 CUDA Dynamic Parallelism 311 12.1 Introduction . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 311 12.1.1 Overview . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 311 12.1.2 Glossary . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 311 12.2 Execution Environment and Memory Model . . . . . . . . . . . . . . . . . . . . . . . . . . . . 312 12.2.1 Execution Environment . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 312 12.2.1.1 Parent and Child Grids . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 312 12.2.1.2 Scope of CUDA Primitives . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 313 12.2.1.3 Synchronization . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 313 12.2.1.4 Streams and Events . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 314 12.2.1.5 Ordering and Concurrency . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 314 12.2.1.6 Device Management . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 314 12.2.2 Memory Model . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 315 12.2.2.1 Coherence and Consistency . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 315 12.3 Programming Interface . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 317 12.3.1 CUDA C++ Reference . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 317 12.3.1.1 Device-Side Kernel Launch . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 317 12.3.1.2 Streams . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 318 12.3.1.3 Events . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 320 12.3.1.4 Synchronization . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 321 12.3.1.5 Device Management . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 321 12.3.1.6 Memory Declarations . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 321 12.3.1.7 API Errors and Launch Failures . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 322 12.3.1.8 API Reference . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 324 12.3.2 Device-side Launch from PTX . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 326 12.3.2.1 Kernel Launch APIs . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 326 12.3.2.2 Parameter Buffer Layout . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 327 12.3.3 Toolkit Support for Dynamic Parallelism . . . . . . . . . . . . . . . . . . . . . . . . . . . . 327 12.3.3.1 Including Device Runtime API in CUDA Code . . . . . . . . . . . . . . . . . . . . . . 327 12.3.3.2 Compiling and Linking . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 328 12.4 Programming Guidelines . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 328 12.4.1 Basics . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 328 12.4.2 Performance . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 329 12.4.2.1 Dynamic-parallelism-enabled Kernel Overhead . . . . . . . . . . . . . . . . . . . . . 329 12.4.3 Implementation Restrictions and Limitations . . . . . . . . . . . . . . . . . . . . . . . . . 329 12.4.3.1 Runtime . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 330 12.5 CDP2 vs CDP1 . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 332 12.5.1 Differences Between CDP1 and CDP2 . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 332 viii

Page 11

12.5.2 Compatibility and Interoperability . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 333 12.6 Legacy CUDA Dynamic Parallelism (CDP1) . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 333 12.6.1 Execution Environment and Memory Model (CDP1) . . . . . . . . . . . . . . . . . . . . . 333 12.6.1.1 Execution Environment (CDP1) . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 334 12.6.1.2 Memory Model (CDP1) . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 337 12.6.2 Programming Interface (CDP1) . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 340 12.6.2.1 CUDA C++ Reference (CDP1) . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 340 12.6.2.2 Device-side Launch from PTX (CDP1) . . . . . . . . . . . . . . . . . . . . . . . . . . . 348 12.6.2.3 Toolkit Support for Dynamic Parallelism (CDP1) . . . . . . . . . . . . . . . . . . . . 350 12.6.3 Programming Guidelines (CDP1) . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 351 12.6.3.1 Basics (CDP1) . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 351 12.6.3.2 Performance (CDP1) . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 352 12.6.3.3 Implementation Restrictions and Limitations (CDP1) . . . . . . . . . . . . . . . . . 353 13 Virtual Memory Management 357 13.1 Introduction . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 357 13.2 Query for Support . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 358 13.3 Allocating Physical Memory . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 358 13.3.1 Shareable Memory Allocations . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 359 13.3.2 Memory Type . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 360 13.3.2.1 Compressible Memory . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 360 13.4 Reserving a Virtual Address Range . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 360 13.5 Virtual Aliasing Support . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 361 13.6 Mapping Memory . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 362 13.7 Controlling Access Rights . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 362 13.8 Fabric Memory . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 363 13.8.1 Query for Support . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 363 13.9 Multicast Support . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 363 13.9.1 Query for Support . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 364 13.9.2 Allocating Multicast Objects . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 364 13.9.3 Add Devices to Multicast Objects . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 365 13.9.4 Bind Memory to Multicast Objects . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 365 13.9.5 Use Multicast Mappings . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 365 14 Stream Ordered Memory Allocator 367 14.1 Introduction . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 367 14.2 Query for Support . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 367 14.3 API Fundamentals (cudaMallocAsync and cudaFreeAsync) . . . . . . . . . . . . . . . . . . 368 14.4 Memory Pools and the cudaMemPool_t . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 369 14.5 Default/Implicit Pools . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 370 14.6 Explicit Pools . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 370 14.7 Physical Page Caching Behavior . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 371 14.8 Resource Usage Statistics . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 372 14.9 Memory Reuse Policies . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 372 14.9.1 cudaMemPoolReuseFollowEventDependencies . . . . . . . . . . . . . . . . . . . . . . . 373 14.9.2 cudaMemPoolReuseAllowOpportunistic . . . . . . . . . . . . . . . . . . . . . . . . . . . . 373 14.9.3 cudaMemPoolReuseAllowInternalDependencies . . . . . . . . . . . . . . . . . . . . . . . 373 14.9.4 Disabling Reuse Policies . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 374 14.10 Device Accessibility for Multi-GPU Support . . . . . . . . . . . . . . . . . . . . . . . . . . . . 374 14.11 IPC Memory Pools . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 375 14.11.1 Creating and Sharing IPC Memory Pools . . . . . . . . . . . . . . . . . . . . . . . . . . . . 375 14.11.2 Set Access in the Importing Process . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 376 14.11.3 Creating and Sharing Allocations from an Exported Pool . . . . . . . . . . . . . . . . . 376 14.11.4 IPC Export Pool Limitations . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 378 ix

Page 12

14.11.5 IPC Import Pool Limitations . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 378 14.12 Synchronization API Actions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 378 14.13 Addendums . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 379 14.13.1 cudaMemcpyAsync Current Context/Device Sensitivity . . . . . . . . . . . . . . . . . . 379 14.13.2 cuPointerGetAttribute Query . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 379 14.13.3 cuGraphAddMemsetNode . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 379 14.13.4 Pointer Attributes . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 379 15 Graph Memory Nodes 381 15.1 Introduction . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 381 15.2 Support and Compatibility . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 381 15.3 API Fundamentals . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 382 15.3.1 Graph Node APIs . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 382 15.3.2 Stream Capture . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 384 15.3.3 Accessing and Freeing Graph Memory Outside of the Allocating Graph . . . . . . . . 385 15.3.4 cudaGraphInstantiateFlagAutoFreeOnLaunch . . . . . . . . . . . . . . . . . . . . . . . . 387 15.4 Optimized Memory Reuse . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 388 15.4.1 Address Reuse within a Graph . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 389 15.4.2 Physical Memory Management and Sharing . . . . . . . . . . . . . . . . . . . . . . . . . 389 15.5 Performance Considerations . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 392 15.5.1 First Launch / cudaGraphUpload . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 392 15.6 Physical Memory Footprint . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 392 15.7 Peer Access . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 393 15.7.1 Peer Access with Graph Node APIs . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 393 15.7.2 Peer Access with Stream Capture . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 394 16 Mathematical Functions 395 16.1 Standard Functions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 395 16.2 Intrinsic Functions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 404 17 C++ Language Support 409 17.1 C++11 Language Features . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 409 17.2 C++14 Language Features . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 412 17.3 C++17 Language Features . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 412 17.4 C++20 Language Features . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 413 17.5 Restrictions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 413 17.5.1 Host Compiler Extensions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 413 17.5.2 Preprocessor Symbols . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 413 17.5.2.1 __CUDA_ARCH__ . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 413 17.5.3 Qualifiers . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 415 17.5.3.1 Device Memory Space Specifiers . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 415 17.5.3.2 __managed__ Memory Space Specifier . . . . . . . . . . . . . . . . . . . . . . . . . . 416 17.5.3.3 Volatile Qualifier . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 417 17.5.4 Pointers . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 418 17.5.5 Operators . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 418 17.5.5.1 Assignment Operator . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 418 17.5.5.2 Address Operator . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 418 17.5.6 Run Time Type Information (RTTI) . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 418 17.5.7 Exception Handling . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 418 17.5.8 Standard Library . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 419 17.5.9 Namespace Reservations . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 419 17.5.10 Functions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 420 17.5.10.1 External Linkage . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 420 17.5.10.2 Implicitly-declared and explicitly-defaulted functions . . . . . . . . . . . . . . . . . 420 x

Page 13

17.5.10.3 Function Parameters . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 421 17.5.10.4 Static Variables within Function . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 424 17.5.10.5 Function Pointers . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 425 17.5.10.6 Function Recursion . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 425 17.5.10.7 Friend Functions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 425 17.5.10.8 Operator Function . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 426 17.5.10.9 Allocation and Deallocation Functions . . . . . . . . . . . . . . . . . . . . . . . . . . 426 17.5.11 Classes . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 426 17.5.11.1 Data Members . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 426 17.5.11.2 Function Members . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 426 17.5.11.3 Virtual Functions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 426 17.5.11.4 Virtual Base Classes . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 427 17.5.11.5 Anonymous Unions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 427 17.5.11.6 Windows-Specific . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 427 17.5.12 Templates . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 428 17.5.13 Trigraphs and Digraphs . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 429 17.5.14 Const-qualified variables . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 429 17.5.15 Long Double . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 430 17.5.16 Deprecation Annotation . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 430 17.5.17 Noreturn Annotation . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 430 17.5.18 [[likely]] / [[unlikely]] Standard Attributes . . . . . . . . . . . . . . . . . . . . . . . . . . . 430 17.5.19 const and pure GNU Attributes . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 431 17.5.20 __nv_pure__ Attribute . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 431 17.5.21 Intel Host Compiler Specific . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 431 17.5.22 C++11 Features . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 432 17.5.22.1 Lambda Expressions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 432 17.5.22.2 std::initializer_list . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 433 17.5.22.3 Rvalue references . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 433 17.5.22.4 Constexpr functions and function templates . . . . . . . . . . . . . . . . . . . . . . 434 17.5.22.5 Constexpr variables . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 434 17.5.22.6 Inline namespaces . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 435 17.5.22.7 thread_local . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 436 17.5.22.8 __global__ functions and function templates . . . . . . . . . . . . . . . . . . . . . . 436 17.5.22.9 __managed__ and __shared__ variables . . . . . . . . . . . . . . . . . . . . . . . . . 437 17.5.22.10Defaulted functions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 438 17.5.23 C++14 Features . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 438 17.5.23.1 Functions with deduced return type . . . . . . . . . . . . . . . . . . . . . . . . . . . 439 17.5.23.2 Variable templates . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 440 17.5.24 C++17 Features . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 440 17.5.24.1 Inline Variable . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 440 17.5.24.2 Structured Binding . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 441 17.5.25 C++20 Features . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 441 17.5.25.1 Module support . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 441 17.5.25.2 Coroutine support . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 441 17.5.25.3 Three-way comparison operator . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 441 17.5.25.4 Consteval functions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 442 17.6 Polymorphic Function Wrappers . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 442 17.7 Extended Lambdas . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 445 17.7.1 Extended Lambda Type Traits . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 446 17.7.2 Extended Lambda Restrictions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 447 17.7.3 Notes on __host__ __device__ lambdas . . . . . . . . . . . . . . . . . . . . . . . . . . . . 457 17.7.4 *this Capture By Value . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 457 17.7.5 Additional Notes . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 460 17.8 Relaxed Constexpr (-expt-relaxed-constexpr) . . . . . . . . . . . . . . . . . . . . . . . . . . . 460 xi

Page 14

17.9 Code Samples . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 463 17.9.1 Data Aggregation Class . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 463 17.9.2 Derived Class . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 464 17.9.3 Class Template . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 464 17.9.4 Function Template . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 465 17.9.5 Functor Class . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 465 18 Texture Fetching 467 18.1 Nearest-Point Sampling . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 467 18.2 Linear Filtering . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 468 18.3 Table Lookup . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 469 19 Compute Capabilities 471 19.1 Feature Availability . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 471 19.2 Features and Technical Specifications . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 472 19.3 Floating-Point Standard . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 481 19.4 Compute Capability 5.x . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 482 19.4.1 Architecture . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 482 19.4.2 Global Memory . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 483 19.4.3 Shared Memory . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 483 19.5 Compute Capability 6.x . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 486 19.5.1 Architecture . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 486 19.5.2 Global Memory . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 486 19.5.3 Shared Memory . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 486 19.6 Compute Capability 7.x . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 487 19.6.1 Architecture . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 487 19.6.2 Independent Thread Scheduling . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 487 19.6.3 Global Memory . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 490 19.6.4 Shared Memory . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 490 19.7 Compute Capability 8.x . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 491 19.7.1 Architecture . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 491 19.7.2 Global Memory . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 492 19.7.3 Shared Memory . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 492 19.8 Compute Capability 9.0 . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 493 19.8.1 Architecture . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 493 19.8.2 Global Memory . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 493 19.8.3 Shared Memory . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 493 19.8.4 Features Accelerating Specialized Computations . . . . . . . . . . . . . . . . . . . . . . 494 19.9 Compute Capability 10.0 . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 494 19.9.1 Architecture . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 494 19.9.2 Global Memory . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 495 19.9.3 Shared Memory . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 495 19.9.4 Features Accelerating Specialized Computations . . . . . . . . . . . . . . . . . . . . . . 495 19.10 Compute Capability 12.0 . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 496 19.10.1 Architecture . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 496 19.10.2 Global Memory . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 496 19.10.3 Shared Memory . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 496 19.10.4 Features Accelerating Specialized Computations . . . . . . . . . . . . . . . . . . . . . . 497 20 Driver API 499 20.1 Context . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 501 20.2 Module . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 502 20.3 Kernel Execution . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 503 20.4 Interoperability between Runtime and Driver APIs . . . . . . . . . . . . . . . . . . . . . . . . 505 xii

Page 15

20.5 Driver Entry Point Access . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 505 20.5.1 Introduction . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 505 20.5.2 Driver Function Typedefs . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 506 20.5.3 Driver Function Retrieval . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 507 20.5.3.1 Using the Driver API . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 507 20.5.3.2 Using the Runtime API . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 508 20.5.3.3 Retrieve Per-thread Default Stream Versions . . . . . . . . . . . . . . . . . . . . . . 508 20.5.3.4 Access New CUDA features . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 509 20.5.4 Potential Implications with cuGetProcAddress . . . . . . . . . . . . . . . . . . . . . . . . 510 20.5.4.1 Implications with cuGetProcAddress vs Implicit Linking . . . . . . . . . . . . . . . 510 20.5.4.2 Compile Time vs Runtime Version Usage in cuGetProcAddress . . . . . . . . . . . 511 20.5.4.3 API Version Bumps with Explicit Version Checks . . . . . . . . . . . . . . . . . . . . 512 20.5.4.4 Issues with Runtime API Usage . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 513 20.5.4.5 Issues with Runtime API and Dynamic Versioning . . . . . . . . . . . . . . . . . . . 514 20.5.4.6 Issues with Runtime API allowing CUDA Version . . . . . . . . . . . . . . . . . . . . 515 20.5.4.7 Implications to API/ABI . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 515 20.5.5 Determining cuGetProcAddress Failure Reasons . . . . . . . . . . . . . . . . . . . . . . . 515 21 CUDA Environment Variables 517 22 Unified Memory Programming 525 22.1 Unified Memory Introduction . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 525 22.1.1 System Requirements for Unified Memory . . . . . . . . . . . . . . . . . . . . . . . . . . 526 22.1.2 Programming Model . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 528 22.1.2.1 Allocation APIs for System-Allocated Memory . . . . . . . . . . . . . . . . . . . . . 530 22.1.2.2 Allocation API for CUDA Managed Memory: cudaMallocManaged() . . . . . . . 531 22.1.2.3 Global-Scope Managed Variables Using __managed__ . . . . . . . . . . . . . . . . 532 22.1.2.4 Difference between Unified Memory and Mapped Memory . . . . . . . . . . . . . 533 22.1.2.5 Pointer Attributes . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 533 22.1.2.6 Runtime detection of Unified Memory Support Level . . . . . . . . . . . . . . . . . 534 22.1.2.7 GPU Memory Oversubscription . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 535 22.1.2.8 Performance Hints . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 535 22.2 Unified memory on devices with full CUDA Unified Memory support . . . . . . . . . . . . 538 22.2.1 System-Allocated Memory: in-depth examples . . . . . . . . . . . . . . . . . . . . . . . 538 22.2.1.1 File-backed Unified Memory . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 540 22.2.1.2 Inter-Process Communication (IPC) with Unified Memory . . . . . . . . . . . . . . 541 22.2.2 Performance Tuning . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 541 22.2.2.1 Memory Paging and Page Sizes . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 542 22.2.2.2 Direct Unified Memory Access from host . . . . . . . . . . . . . . . . . . . . . . . . 544 22.2.2.3 Host Native Atomics . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 546 22.2.2.4 Atomic accesses & synchronization primitives . . . . . . . . . . . . . . . . . . . . . 546 22.2.2.5 Memcpy()/Memset() Behavior With Unified Memory . . . . . . . . . . . . . . . . . 547 22.3 Unified memory on devices without full CUDA Unified Memory support . . . . . . . . . . 547 22.3.1 Unified memory on devices with only CUDA Managed Memory support . . . . . . . . 547 22.3.2 Unified memory on Windows or devices with compute capability 5.x . . . . . . . . . . 548 22.3.2.1 Data Migration and Coherency . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 548 22.3.2.2 GPU Memory Oversubscription . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 548 22.3.2.3 Multi-GPU . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 548 22.3.2.4 Coherency and Concurrency . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 549 23 Lazy Loading 557 23.1 What is Lazy Loading? . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 557 23.2 Lazy Loading version support . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 557 23.2.1 Driver . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 558 xiii

Page 16

23.2.2 Toolkit . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 558 23.2.3 Compiler . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 558 23.3 Triggering loading of kernels in lazy mode . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 558 23.3.1 CUDA Driver API . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 558 23.3.2 CUDA Runtime API . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 559 23.4 Querying whether Lazy Loading is Turned On . . . . . . . . . . . . . . . . . . . . . . . . . . . 559 23.5 Possible Issues when Adopting Lazy Loading . . . . . . . . . . . . . . . . . . . . . . . . . . . 559 23.5.1 Concurrent Execution . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 560 23.5.2 Allocators . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 560 23.5.3 Autotuning . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 560 24 Extended GPU Memory 561 24.1 Preliminaries . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 561 24.1.1 EGM Platforms: System topology . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 562 24.1.2 Socket Identifiers: What are they? How to access them? . . . . . . . . . . . . . . . . . 562 24.1.3 Allocators and EGM support . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 562 24.1.4 Memory management extensions to current APIs . . . . . . . . . . . . . . . . . . . . . . 562 24.2 Using the EGM Interface . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 563 24.2.1 Single-Node, Single-GPU . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 563 24.2.2 Single-Node, Multi-GPU . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 563 24.2.2.1 Using VMM APIs . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 564 24.2.2.2 Using CUDA Memory Pool . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 564 24.2.3 Multi-Node, Single-GPU . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 565 25 Notices 567 25.1 Notice . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 567 25.2 OpenCL . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 568 25.3 Trademarks . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 568 xiv

Page 17

CUDA C++ Programming Guide, Release 12.8 CUDA C++ Programming Guide The programming guide to the CUDA model and interface. Changes in Version 12.8 ▶ Added section TMA Swizzle Contents 1

Page 18

CUDA C++ Programming Guide, Release 12.8 2 Contents

Page 19

Chapter 1. The Benefits of Using GPUs The Graphics Processing Unit (GPU)1 provides much higher instruction throughput andmemory band- width than the CPUwithin a similar price and power envelope. Many applications leverage these higher capabilities to run faster on the GPU than on the CPU (see GPU Applications). Other computing de- vices, like FPGAs, are also very energy efficient, but offermuch less programming flexibility than GPUs. This difference in capabilities between the GPU and the CPU exists because they are designed with different goals in mind. While the CPU is designed to excel at executing a sequence of operations, called a thread, as fast as possible and can execute a few tens of these threads in parallel, the GPU is designed to excel at executing thousands of them in parallel (amortizing the slower single-thread performance to achieve greater throughput). The GPU is specialized for highly parallel computations and therefore designed such thatmore transis- tors are devoted to data processing rather than data caching and flow control. The schematic Figure 1 shows an example distribution of chip resources for a CPU versus a GPU. Figure 1: The GPU Devotes More Transistors to Data Processing Devoting more transistors to data processing, for example, floating-point computations, is beneficial for highly parallel computations; the GPU can hidememory access latencies with computation, instead 1 The graphics qualifier comes from the fact that when the GPU was originally created, two decades ago, it was designed as a specialized processor to accelerate graphics rendering. Driven by the insatiable market demand for real-time, high-definition, 3D graphics, it has evolved into a general processor used for many more workloads than just graphics rendering. 3

Page 20

CUDA C++ Programming Guide, Release 12.8 of relying on large data caches and complex flow control to avoid long memory access latencies, both of which are expensive in terms of transistors. In general, an application has amix of parallel parts and sequential parts, so systems are designedwith a mix of GPUs and CPUs in order to maximize overall performance. Applications with a high degree of parallelism can exploit this massively parallel nature of the GPU to achieve higher performance than on the CPU. 4 Chapter 1. The Benefits of Using GPUs