多GPU Cuda计算

Question

我是一个多GPU编程的新手，我有一些关于多GPU计算的问题。举例来说，让我们看看点积产品的例子。我正在运行一个创建2个大数组A [N]和B [N]的CPU线程。由于这些阵列的大小，我需要将他们的点积计算分割成两个GPU，两个都是Tesla M2050（计算能力2.0）。问题是我需要在由我的CPU线程控制的do循环内多次计算这些点积。每个点积都需要前一个的结果。我已经读过关于创建2个不同的线程来分别控制2个不同的GPU（如cuda上所描述的那样），但我不知道如何同步和交换它们之间的数据。有另外一种选择吗？我真的很感激任何帮助/ example.Thanks提前！

Answer 1

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

因此，从CUDA 4.0开始，您可以从单线程主机程序轻松编程多个GPU。 Here are some slides I presented last year about this。

编程多个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]; 
} 

Answer 2

要创建多个线程，可以使用OpenMP或pthreads。要做你正在谈论的事情，看起来你需要创建和启动两个线程（omp parallel section或pthread_create），让每个线程完成它的计算部分并将其中间结果存储在单独的process-wIDE变量中（回想一下，全局变量在一个进程的线程之间自动共享，因此原始线程将能够看到两个衍生线程所做的更改）。为了让原始线程等待其他线程完成，在完成两个衍生线程之后（使用全局屏障或线程连接操作）进行同步并将结果合并到原始线程中（如果将数组拆分为一半并且通过乘以相应的元素计算点积，并在两半上执行全局总和减少，只需要添加来自两个衍生线程的两个中间结果）。

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