GPUNvidia CUDA 编程高级教程——支持点对点访问的多 GPU

Posted 2022-12-07 从善若水

博主未授权任何人或组织机构转载博主任何原创文章，感谢各位对原创的支持！
博主链接

本人就职于国际知名终端厂商，负责modem芯片研发。
在5G早期负责终端数据业务层、核心网相关的开发工作，目前牵头6G算力网络技术标准研究。


博客内容主要围绕：
       5G/6G协议讲解
       算力网络讲解（云计算，边缘计算，端计算）
       高级C语言讲解
       Rust语言讲解

支持点对点访问的多 GPU

通用虚拟地址空间

CUDA 使用通用虚拟地址 (UVA) 空间。在 UVA 空间中，CPU 和 GPU 上的所有通过 CUDA 分配的空间（包括cudaMalloc和cudaMallocHost）都可确保享有唯一的虚拟地址。例如，您可以使用cudaMallocHost或cudaHostAlloc分配固定的主机内存，并在设备代码中直接获取其地址（同时固定了虚拟到物理的地址转换，这样 GPU 就不需要与 CPU 的内存管理单元对话）。

在 UVA 范例中，CUDA 知道给定的地址属于哪台设备，因为UVA的构造方法可以确保系统不会为不同设备分配的空间使用相同的地址。

注意：上图描述了 GPU 通过 PCIe 连接，但是在 UVA 受到支持时，它也可以通过 NVLink 或 NVSwitch 工作。

直接点对点内存访问

UVA 还支持直接点对点内存访问，有时也叫 GPUDirect Peer-to-Peer (P2P) 。当多个 GPU 连接到同一个 PCI-e 树或通过 NVLINK 互连时，GPU Direct P2P才是可以使用的。它与 UVA 是截然不同的概念，但由 UVA 促进实现。

启用直接点对点内存访问

除了一些例外情况（取决于系统 PCIe、NVLink 或 NVSwitch 拓扑），一个 GPU 可以直接读取和写入同一服务器上的另一个 GPU 的地址。我们使用 CUDA API 调用 cudaDeviceCanAccessPeer()，来检查是否可以在给定的配置下这么做。假设可以这样做，我们要在程序的开头使用 cudaDeviceEnablePeerAccess() 启用这个点对点访问功能。

int this_device = 0;
int peer_device = 1;

cudaSetDevice(this_device);

int can_access_peer;

cudaDeviceCanAccessPeer(&can_access_peer, this_device, peer_device);

if (can_access_peer) 
    cudaDeviceEnablePeerAccess(peer_device, 0); // Note: `0` is the required value passed to this 2nd positional argument which is being reserved for future use.

代码实现

我们在应用中尝试一下。我们的策略是每个线程都更新相同的点击计数器，而不是每个 GPU 都有一个计数器。我们将把这个计数器任意放置在 GPU 0 上。这样一来，应用程序看起来会更像最初的单 GPU 的情况，因为我们不再需要为每个可用的 GPU 分配和复制内存。另一方面，至少对于此应用而言，这种方法会增加计数器上可能发生的原子操作的碰撞次数。

#include <iostream>
#include <curand_kernel.h>

#define N 1024*1024

__global__ void calculate_pi(int* hits, int device) 
    int idx = threadIdx.x + blockIdx.x * blockDim.x;

    // 初始化随机数状态（网格中的每个线程不得重复）
    int seed = device;
    int offset = 0;
    curandState_t curand_state;
    curand_init(seed, idx, offset, &curand_state);

    // 在 (0.0, 1.0] 内生成随机坐标
    float x = curand_uniform(&curand_state);
    float y = curand_uniform(&curand_state);

    // 如果这一点在圈内，增加点击计数器
    if (x * x + y * y <= 1.0f) 
        atomicAdd(hits, 1);
    



int main(int argc, char** argv) 
    // 启动 GPU 0
    cudaSetDevice(0);

    int device_count;
    cudaGetDeviceCount(&device_count);

    // 分配主机和设备值
    int* hits;
    hits = (int*) malloc(sizeof(int));

    int* d_hits;
    cudaMalloc((void**) &d_hits, sizeof(int));

    // 初始化点击次数并复制到设备
    *hits = 0;
    cudaMemcpy(d_hits, hits, sizeof(int), cudaMemcpyHostToDevice);

    // 检查每台设备都能访问其对等设备。
    // 如果可以，请继续并启用该访问。

    for (int dev = 0; dev < device_count; ++dev) 
        cudaSetDevice(dev);
        for (int peer = 0; peer < device_count; ++peer) 
            if (peer != dev) 
                int can_access_peer;
                cudaDeviceCanAccessPeer(&can_access_peer, dev, peer);

                if (can_access_peer) 
                    cudaDeviceEnablePeerAccess(peer, 0);
                 else 
                    std::cout << "Device " << dev << " could not access peer " << peer << std::endl;
                    return -1;
                
            
        
    

    // 启动核函数进行计算
    int threads_per_block = 256;
    int blocks = (N / device_count + threads_per_block - 1) / threads_per_block;

    // 先启动所有核函数，以支持异步执行
    // 然后在所有设备上同步。
    for (int i = 0; i < device_count; ++i) 
        cudaSetDevice(i);
        calculate_pi<<<blocks, threads_per_block>>>(d_hits, i);
    

    for (int i = 0; i < device_count; ++i) 
        cudaSetDevice(i);
        cudaDeviceSynchronize();
    

    // 将最终结果复制回主机
    cudaMemcpy(hits, d_hits, sizeof(int), cudaMemcpyDeviceToHost);

    // 计算 pi 的最终值
    float pi_est = (float) *hits / (float) (N) * 4.0f;

    // 打印结果
    std::cout << "Estimated value of pi = " << pi_est << std::endl;
    std::cout << "Error = " << std::abs((M_PI - pi_est) / pi_est) << std::endl;

    // 清理
    free(hits);
    cudaFree(d_hits);

运行结果

Estimated value of pi = 3.14072
Error = 0.000277734
CPU times: user 30.8 ms, sys: 6.3 ms, total: 37.1 ms
Wall time: 2.41 s

---