CUDA C编程手册: 编程接口（四）

• CUDA C 运行时
• * 多设备系统
  • * 设备枚举
    • 设备选取
    • 流与事件的行为
    • 点对点内存访存
  • 统一虚拟内存地址空间
  • 进程间通信
  • 错误校验
  • 回调

CUDA C 运行时

多设备系统

设备枚举

一个主机系统可以拥有多个设备。下列代码演示了如何枚举这些设备， 查询它们的属性以及确定激活的设备数量。

int deviceCount;
cudaGetDeviceCount(&deviceCount);
int device;
for (device = 0; device < deviceCount; ++device) {
    cudaDeviceProp deviceProp;
    cudaGetDeviceProperties(&deviceProp, device);
    printf("Device %d has compute capability %d.%d.\n",
    device, deviceProp.major, deviceProp.minor);
}

设备选取

主机线程可以通过cudaSetDevice()在任何时刻设置它想利用的设备。设备内存的分配和核函数的启动都会在当前设置的设备上；流与事件的创建也是发生在当前所设置的设备上。如果没有显式的调用设置设备的函数，则默认的当前设备是0。

size_t size = 1024 * sizeof(float);
cudaSetDevice(0); // Set device 0 as current
float* p0;
cudaMalloc(&p0, size); // Allocate memory on device 0
MyKernel<<<1000, 128>>>(p0); // Launch kernel on device 0
cudaSetDevice(1); // Set device 1 as current
float* p1;
cudaMalloc(&p1, size); // Allocate memory on device 1
MyKernel<<<1000, 128>>>(p1); // Launch kernel on device 1

流与事件的行为

如果一个流并不在相对应的设备上，在其中启动核函数则会会失败。但是内存的拷贝则总是会成功。

cudaSetDevice(0); // Set device 0 as current
cudaStream_t s0;
cudaStreamCreate(&s0); // Create stream s0 on device 0
MyKernel<<<100, 64, 0, s0>>>(); // Launch kernel on device 0 in s0
cudaSetDevice(1); // Set device 1 as current
cudaStream_t s1;
cudaStreamCreate(&s1); // Create stream s1 on device 1
MyKernel<<<100, 64, 0, s1>>>(); // Launch kernel on device 1 in s1
// This kernel launch will fail:
MyKernel<<<100, 64, 0, s0>>>(); // Launch kernel on device 1 in s0

cudaEventRecord()对于输入事件和输入流并不在相应的设备上则会调用失败。
cudaEventElapsedTime()对于输入的两个时间在不同的设备上会发生调用失败。
cudaEventSynchronize()和cudaEventQuery()对于及时输入事件并不关联当前设备也能成功。
cudaStreamWaitEvent()在输入流和输入事件关联不同的设备情况下也能成功调用。因此它常被用来同步不同设备。
每个设备都有自己的默认流， 所以设备上的默认流中的命令执行时无序的或者与其他设备上的默认流中的命令并发执行。

点对点内存访存

当应用程序以64位的进程运行，Tesla系列中计算力高于2.0的设备之间可以互相访问对方的内存。当cudaDeviceCanAccessPeer()返回True时则表示设备之间支持点对点的内存访存。 内存的点对点访存这一特性必须通过cudaDeviceEnablePeerAccess()来激活。在激活non-NVSwitch的系统上，每个设备可以支持系统最大允许的8个点连接。所有的设备使用同一个统一地址空间。因此同一个地址可以被不同设备使用以访问同一块内存。

cudaSetDevice(0); // Set device 0 as current
float* p0;
size_t size = 1024 * sizeof(float);
cudaMalloc(&p0, size); // Allocate memory on device 0
MyKernel<<<1000, 128>>>(p0); // Launch kernel on device 0
cudaSetDevice(1); // Set device 1 as current
cudaDeviceEnablePeerAccess(0, 0); // Enable peer-to-peer access
// with device 0
// Launch kernel on device 1
// This kernel launch can access memory on device 0 at address p0
MyKernel<<<1000, 128>>>(p0);

当所有设备使用统一内存空间时， 可以使用常规的内存拷贝函数来进行跨设备的内存拷贝。否则，需要使用 cudaMemcpyPeer()、cudaMemcpyPeerAsync()、cudaMemcpy3DPeer() 和cudaMemcpy3DPeerAsync()来进行设备之间内存的拷贝。

cudaSetDevice(0); // Set device 0 as current
float* p0;
size_t size = 1024 * sizeof(float);
cudaMalloc(&p0, size); // Allocate memory on device 0
cudaSetDevice(1); // Set device 1 as current
float* p1;
cudaMalloc(&p1, size); // Allocate memory on device 1
cudaSetDevice(0); // Set device 0 as current
MyKernel<<<1000, 128>>>(p0); // Launch kernel on device 0
cudaSetDevice(1); // Set device 1 as current
cudaMemcpyPeer(p1, 1, p0, 0, size); // Copy p0 to p1
MyKernel<<<1000, 128>>>(p1); // Launch kernel on device 1

设备之间的内存拷贝在它之前的所有命令完成之前不会开始，并且在拷贝完成之前其后的命令也不会开始执行。
与正常的流行为一致的是，设备之间的异步内存拷贝会与其他流之间的拷贝与核函数进行overlap。

需要指出的是，如果点对点的访存通过接口的调用激活之后，点对点的内存拷贝从此不再需要借助主机端内存来缓存，因此速度会更加快速。

统一虚拟内存地址空间

对于计算力高于2.0的设备，当应用程序是一个64位的进程时，主机和所有设备可以使用同一个地址空间。所有的通过CUDA API调用进行的主机内存分配和设备内存分配都在同一个虚拟地址范围内。这就会：

• 主机端内存和设备端内存使用统一地址空间，能够通过接口cudaPointerGetAttributes()来获取其原始位置。
• 当进行设备端内存的拷贝时，cudaMemcpy*()的参数cudaMemcpyKind可以被设置成cudaMemcpyDefault()来通过指针判断其位置。这个方法同样适用于并不是通过CUDA API分配的主机端内存，只要当前的设备使用了统一内存寻址。
• 通过cudaHostAlloc()的内存分配可以在开启统一地址空间的设备之间自动地移植，也能够直接被核函数直接使用。

应用程序可以通过设备属性中的unifiedAddressing来判断是否使用了统一地址空间。结果为1表示已经启用，否则未启用。

进程间通信

任何被主机端线程创建的设备内存指针或者事件句柄都可以被同进程中的其他线程使用。但并不能被其他进程使用， 也不能被其他进程中的线程使用。为了在进程之间共享设备内存指针，一个应用必须使用进程间的通信API（IPC API）, 具体可以查看参考手册。进程间的通信API在LINUX系统上只支持64位进程与设备计算力高于2.0的设备。需要记住的是， IPC API 并不支持cudaMallocManaged分配的内存。使用这个API， 一个应用程序使用cudaIpcGetMemHandle()可以得到一个含有给定设备内存指针的IPC句柄， 使用标准IPC机制将其传递给其他进程， 然后使用cudaIpcOpenMemHandle()来从IPC句柄中恢复一个有效的设备指针。事件句柄也可以通过类似的方法来进行共享。

一个使用IPC API通信的场景就是： 当一个主进程生成了一批输入数据，为了将这些数据在副进程中使用而不产生数据的拷贝或者重新生成， 则可以使用ICP 来实现这种需求。

使用CUDA IPC 的应用程序，必须使用同样的CUDA 驱动和运行时来进行编译，连接以及运行。

错误校验

所有的运行时函数都会返回衣蛾错误代码。但对于异步函数， 返回的错误代码可能代表发生在设备上的任何异步错误，因为函数实在设备执行完任务前进行返回的。错误代码智能表示发生在主机代码执任务之前， 如参数检查。如果任何一个异步错误发生，它可能被之后无任何关系的运行函数所报告。

唯一检测异步错误的办法就是在相关异步函数调用的只有使用cudaDeviceSynchronize()来进行同步并同时检验这个函数返回的错误代码。

运行时为每个主机线程保存了一个初始化值为cudaSuccess错误变量，当每次有错误发生时这个变量都会被重写。cudaPeekAtLastError()可以取回这个错误代码。cudaGetLastError()也可以获取这个变量值并会将之重置为初始值。

核的启动并不会返回任何的错误代码，所以之前的两个函数需要在每个核函数启动之后立即进行调用，这样才能有效的检测每次启动是否会发生错误。为了保证每次返回的错误代码都是由当前核所造成的，应该在每次核启动前将错误代码变得置为初始值cudaSuccess。核的启动是异步的，因此为了检测异步错误，必须要在核与核之间进行同步，然后才调用上述接口来进行错误检验。

注意， cudaStreamQuery()和cudaEventQuery()返回的cudaErrorNotReady并不会被认为是一个错误。因此，两个获取错误代码信息的API也不会对此进行报告。

回调

对于计算力大于2.0的设备， 调用堆栈的大小可以通过cudaDeviceGetLimit()和cudaDeviceSetLimit()来进行获取和设置。当调用堆栈溢出的时候，核的调用就会失败， 同时会在使用CUDA调试器（cuda-gbd / Nsight）返回一个堆栈溢出错误或者一个未知的启动错误。