在 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 作为起点。我会这样做以进行原型设计。但是请参阅最后关于带宽瓶颈的我的 cmets。
由于您没有提供有关多次运行点积的外循环的足够详细信息,因此我没有尝试对此做任何事情。
// 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];
}