CUDA C++ 编程指引：其他运行时特性 | CUDA

本系列参考自 CUDA C++ Programming Guide。

初始化

运行时库没有显式的初始化函数，在调用第一个函数时会自动初始化（除了错误处理和版本管理函数）。初始化时，会为每个设备生成一个所有主机线程可见的上下文。当主机端调用了 cudaDeviceReset() 函数，则会销毁掉该主机线程正在操作的设备的上下文。

进程间通信

线程间通讯，可以很方便的通过共享内存的变量来实现。然而进程间通讯不行。

为了在进程间共享设备端内存的指针或者事件，必须使用 IPC API。IPC API 只支持 64 位程序，并且要求设备计算能力 ≥2.0。通过 IPC 中的 cudaIpcGetMemHandle()，可以得到设备内存指针的 IPC 句柄。该句柄可以通过标准的 IPC 机制（interprocess shared memory or files）传递到另一个进程，再使用 cudaIpcOpenMemHandle() 解码得到该进程可以使用的设备内存指针。事件的共享也是如此。

使用 IPC API 的一个例子是，单个主进程生成一批输入数据，使数据可用于多个从属进程，而无需重新生成或复制。

使用 CUDA IPC 进行进程间通信的应用程序应使用相同的 CUDA 驱动程序和运行时库进行编译、链接和运行。

错误检查

所有的运行时函数都返回错误码，但对于异步函数，由于会在任务结束前返回，因此错误码不能报告异步调用的错误；错误码只报告在任务执行之前的错误，典型的错误比如参数有效性；如果异步调用出错，错误将会在后面某个无关的函数调用中出现。

唯一能够检查异步调用出错的方式是通过在异步调用函数后面使用 cudaDeviceSynchronize() 同步（或使用前面介绍的其它同步机制），然后检查 cudaDeviceSynchronize() 的返回值。

运行时库为每个主机线程维护着一个初始化为 cudaSuccess 的错误变量，每次错误发生（可以是参数不正确或异步错误）时，该变量会被错误码重写。该变量不会被直接调用，但可以被 cudaPeekAtLastError() 和 cudaGetLastError() 访问到。不同的是，cudaGetLastError() 在返回这一变量的同时，会把它重置为 cudaSuccess。

内核函数不返回任何错误码，所以应当在内核函数启动后立刻调用 cudaGetLastError() 或 cudaPeekAtLastError() 检测内核函数启动前错误。为保证返回的错误值不是由于内核函数启动之前的错误导致的，可以通过在内核函数启动前调用 cudaGetLastError() 保证运行时错误变量在内核函数启动前被设置为 cudaSuccess。内核函数是异步的，因此为了检测异步错误，应用必须在内核函数启动和 cudaGetLastError() 或 cudaPeekAtLastError() 之间同步。

另外需要注意的是，cudaStreamQuery() 和 cudaEventQuery() 这类函数，有可能会返回 cudaErrorNotReady。但这不被认为是错误，因此不会被 cudaPeekAtLastError() 和 cudaGetLastError() 捕获到。

调用栈

在计算能力 ≥2.0 的设备上，调用栈的长度可以使用 cudaDeviceGetLimit() 查询，使用 cudaDeviceSetLimit() 设置。

当调用栈上溢时，如果通过 CUDA 调试器（cuda-gdb，Parallel Nsight）运行，内核会因为栈上溢失败，否则会出现无法确定的启动（unspecified launch）错误。

计算模式

NVIDIA 的设备可以设置三种计算模式：

• 默认模式：多个主机线程可同时使用设备（使用运行时调用 cudaSetDevice()，或使用驱动 API 时将关联到设备的上下文作为当前上下文）。
• 互斥进程计算模式：在系统的所有进程中，一个设备上只能由一个进程创建 CUDA 上下文。一旦创建成功后，该进程的所有线程都可以使用该设备，而其他进程则不行。
• 禁止模式：无法对设备建立 CUDA 上下文，即不允许任何主机线程使用设备。

正常情况下，如果程序没有调用 cudaSetDevice()，则会默认使用 0 号设备。但是如果 0 号设备被置成禁止模式，亦或是被其他进程所专属，则会在其他设备上创建上下文并使用。可以向 cudaSetValidDevices() 函数输入一个设备列表，函数会在第一个可以使用的设备上创建上下文。

Pascal 及以上架构（计算能力 ≥6.0），支持计算抢占（指令级的优先级调度），不再是以线程块为 SM 的最小调度单位，而是以指令为最小调度单位，且具有优先级。这意味着具有冗长 kernel 的线程块不再会占据太多的计算资源，或是发生 timeout。但是这也有缺点：当多个进程创建了上下文时，以往基于线程块的调度不会造成太多的上下文切换，但现在的指令级调度则会造成很多的上下文切换，降低效率。

如果函数 cudaDeviceGetAttribute() 返回 cudaDevAttrComputePreemptionSupported 则设备支持计算抢占。

应用可检查 computeMode 属性以查询设备的计算模式。

参考

CUDA C++ Programming Guide
《CUDA C Programming Guide》(《CUDA C 编程指南》)导读
《CUDA 并行程序设计：GPU 编程指南》