5

我是多 GPU 编程的新手，我对多 GPU 计算有一些疑问。例如，让我们以点积为例。我正在运行一个创建 2 个大数组 A[N] 和 B[N] 的 CPU 线程。由于这些数组的大小，我需要将其点积的计算拆分为 2 个 GPU，均为 Tesla M2050（计算能力 2.0）。问题是我需要在我的 CPU 线程控制的 do-loop 内多次计算这些点积。每个点积都需要前一个点积的结果。我读过关于创建 2 个不同的线程来分别控制 2 个不同的 GPU（如 cuda 中的示例所述），但我不知道如何在它们之间同步和交换数据。还有其他选择吗？我真的很感谢任何形式的帮助/示例。在此先感谢！

4

2 回答 2

7

在 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，你也不会在设备之间获得并发，因为推力::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];
}
于 2012-03-05T05:25:04.220 回答
1

要创建多个线程，您可以使用 OpenMP 或 pthreads。要执行您所说的操作，您似乎需要创建并启动两个线程（omp 并行部分或 pthread_create），让每个线程完成其计算部分并将其中间结果存储在单独的进程范围变量中（回想一下，全局变量在进程的线程之间自动共享，因此原始线程将能够看到两个派生线程所做的更改）。要让原始线程等待其他线程完成，请同步（使用全局屏障或线程连接操作）并在两个衍生线程完成后将结果合并到原始线程中（如果您

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

于 2012-03-04T14:16:46.267 回答