多GPU的Cuda计算

Question

多GPU的Cuda计算

5

我是多GPU编程的新手，对于多GPU计算有一些问题。例如，我们来看一下点积的例子。我正在运行一个创建了两个大数组A[N]和B[N]的CPU线程。由于这些数组的大小，我需要将它们的点积计算分成两个GPU（均为Tesla M2050（计算能力2.0））。问题是，我需要在由CPU线程控制的循环内多次计算这些点积。每个点积都需要前一个点积的结果。我已经阅读了有关创建控制2个不同GPU的2个不同线程的信息（如CUDA by example中所述），但我不知道如何在它们之间进行同步和交换数据。是否还有其他替代方案？我真的非常感谢任何形式的帮助/示例。提前致谢！

- chemeng

2个回答

1

要创建多个线程，您可以使用OpenMP或pthreads。为了做你所说的事情，似乎需要创建并启动两个线程（omp parallel section或pthread_create），让每个线程完成其计算的一部分，并将其中间结果存储在单独的进程范围变量中（请记住，全局变量会自动在进程的线程之间共享，因此原始线程将能够看到由两个生成的线程所做的更改）。为了使原始线程等待其他线程完成，需要同步（使用全局屏障或线程加入操作）并在两个生成的线程完成后在原始线程中合并结果（如果您将数组分成两半并通过乘以相应元素并对一半执行全局求和约简来计算点积，则只需要将两个生成的线程的两个中间结果相加即可）。

您还可以使用MPI或fork，在这种情况下，通信可以通过类似于网络编程的方式进行...管道/套接字或通过（阻塞）发送和接收进行通信和同步。

- Patrick87

这个实现难道不会极大地降低我的应用程序的加速吗？由于频繁的GPU-CPU-CPU-GPU通信..我看到了一些关于属于不同设备的并发流的东西，可能会帮助我解决问题，但我找不到一个有用的例子。 - chemeng

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- harrism · Accepted Answer

在CUDA 4.0之前，多GPU编程需要使用多线程CPU编程。这可能会很具有挑战性，特别是当你需要在线程或GPU之间进行同步和/或通信时。如果你的所有并行性都在GPU代码中，那么使用多个CPU线程可能会增加软件的复杂性，并未能进一步提高GPU的性能。

因此，从CUDA 4.0开始，你可以轻松地从单线程主机程序中编写多个GPU。这可以简单地实现多GPU编程：

这里是我去年关于这方面的演示文稿。

int numDevs = 0;
cudaGetNumDevices(&numDevs);
...
for (int d = 0; d < numDevs; d++) {
    cudaSetDevice(d);
    kernel<<<blocks, threads>>>(args);
}

对于您提到的点积的具体示例，您可以使用thrust::inner_product作为起点进行原型制作。我会在原型制作阶段这样做。但请注意我的评论，关于带宽瓶颈的问题。

由于您没有提供有关运行多次点积的外部循环的足够细节，我没有尝试对其进行任何处理。

// assume the deviceIDs of the two 2050s are dev0 and dev1.
// assume that the whole vector for the dot product is on the host in h_data
// assume that n is the number of elements in h_vecA and h_vecB.

int numDevs = 0;
cudaGetNumDevices(&numDevs);
...
float result = 0.f;
for (int d = 0; d < numDevs; d++) {
    cudaSetDevice(d);
    device_vector<float> vecA(h_vecA + d*(n/d), h_vecA + (d+1)*(n/d)-1);
    device_vector<float> vecB(h_vecB + d*(n/d), h_vecB + (d+1)*(n/d)-1);
    result += thrust::inner_product(vecA.begin(), vecA.end(), vecB.begin(), 0.f);
}

我承认如果n不是numDevs的偶数倍，上面的索引就不正确了，但我会把修复留给读者作为练习。:)

这很简单，是一个很好的开始。先让它工作，然后再优化。

一旦你让它工作了，如果你在设备上所做的全部都是点积计算，你会发现你受到带宽限制 - 主要是由PCI-e造成的，而且你也无法在设备之间实现并发，因为thrust::inner_product由于读回结果而是同步的。所以你可以使用cudaMemcpyAsync( device_vector构造函数将使用cudaMemcpy)。但更容易、更高效的方法是使用"零拷贝"——直接访问主机内存(也在上述多GPU编程演示中讨论过)。由于你所做的只是一次读取每个值并将其添加到总和中(并行重用在共享内存复制中发生)，所以你最好直接从主机读取它，而不是从主机复制到设备，然后在内核中从设备内存中读取它。此外，你还需要在每个GPU上异步启动内核，以确保最大并发性。

你可以像这样做:

int bytes = sizeof(float) * n;
cudaHostAlloc(h_vecA, bytes, cudaHostAllocMapped | cudaHostAllocPortable);
cudaHostAlloc(h_vecB, bytes, cudaHostAllocMapped | cudaHostAllocPortable);
cudaHostAlloc(results, numDevs * sizeof(float), cudaHostAllocMapped | cudaHostAllocPortable);
// ... then fill your input arrays h_vecA and h_vecB


for (int d = 0; d < numDevs; d++) {
    cudaSetDevice(d);
    cudaEventCreate(event[d]));
    cudaHostGetDevicePointer(&dptrsA[d], h_vecA, 0);
    cudaHostGetDevicePointer(&dptrsB[d], h_vecB, 0);
    cudaHostGetDevicePointer(&dresults[d], results, 0);
}

...

for (int d = 0; d < numDevs; d++) {
    cudaSetDevice(d);
    int first = d * (n/d);
    int last   = (d+1)*(n/d)-1;
    my_inner_product<<<grid, block>>>(&dresults[d], 
                                      vecA+first, 
                                      vecA+last, 
                                      vecB+first, 0.f);
    cudaEventRecord(event[d], 0);
}

// wait for all devices
float total = 0.0f;
for (int d = 0; d < devs; d++) {
    cudaEventSynchronize(event[d]);
    total += results[numDevs];
}