CUDA 12.4官方文档 - 英文

CUDA C++ Programming Guide

Release 12.4

NVIDIA

Mar 02, 2024

Contents

1 The Benets 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.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 Workow . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 22

6.1.1.1 Oine Compilation . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 22

6.1.1.2 Just-in-Time Compilation . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 22

6.1.2 Binary Compatibility . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 22

6.1.3 PTX Compatibility . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 23

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.7 Memory Synchronization Domains . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 44

6.2.7.1 Memory Fence Interference . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 44

6.2.7.2 Isolating Trac 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.9.2 Device Selection . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 80

6.2.9.3 Stream and Event Behavior . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 80

6.2.9.4 Peer-to-Peer Memory Access . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 81

6.2.9.5 Peer-to-Peer Memory Copy . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 82

6.2.10 Unied Virtual Address Space . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 82

6.2.11 Interprocess Communication . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 83

6.2.12 Error Checking . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 83

6.2.13 Call Stack . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 84

6.2.14 Texture and Surface Memory . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 84

6.2.14.1 Texture Memory . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 84

6.2.15.3 SLI Interoperability . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 102

6.2.16 External Resource Interoperability . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 103

6.2.16.1 Vulkan Interoperability . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 103

6.2.16.2 OpenGL Interoperability . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 111

6.2.16.3 Direct3D 12 Interoperability . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 112

6.2.16.4 Direct3D 11 Interoperability . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 118

6.2.16.5 NVIDIA Software Communication Interface Interoperability (NVSCI) . . . . . . . . 126

6.3 Versioning and Compatibility . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 132

6.4 Compute Modes . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 133

6.5 Mode Switches . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 134

8.2 Maximize Utilization . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 139

8.2.1 Application 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 Speciers . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 161

10.1.1 __global__ . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 161

10.1.2 __device__ . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 161

10.1.3 __host__ . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 162

10.1.4 Undened behavior . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 162

10.1.5 __noinline__ and __forceinline__ . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 162

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, oat, 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.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.10 tex2DGrad() for sparse CUDA arrays . . . . . . . . . . . . . . . . . . . . . . . . . . . 176

10.8.1.11 tex2DLod() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 176