11.3 使用多个GPU

在11.2节中，我们指出设备要么是集成的GPU，要么是独立的GPU，其中前者是构建在系统芯片组中的，而后者通常是位于PCIE槽中的扩展卡。然而，在越来越多的系统中同时包含了集成GPU和独立GPU，这意味着它们拥有多个支持GUDA的处理器。NVIDIA同样也会不断地推出包含多个GPU的产品，例如GeForce GTX 295。虽然GeForce GTX 295在物理上占用一个扩展槽，但在CUDA应用程序看来则是两个独立的GPU。而且，用户还可以将多个GPU添加到独立的PCIE槽上，通过NVIDIA的SLI技术将它们桥接。这些趋势导致的结果之一就是，在运行CUDA应用程序的系统中，包含多个图形处理器逐渐成为一种常见情况。由于CUDA应用程序是可高度并行化的，如果可以使系统中的每个CUDA设备都实现最大的吞吐量，那么无疑是最好的情形。因此，我们来看看如何实现这个目标。

为了避免学习新的示例，我们将把点积应用程序修改为使用多个GPU。为了降低编码难度，我们将在一个结构中把计算点积所需的全部数据都相加起来。你马上将看到为什么这会降低难度。

struct DataStruct {
    int deviceID;
    int size;
    float *a;
    float *b;
    float returnValue;
};

这个结构包含了在计算点积时使用的设备标识，以及输入缓冲区的大小和指向两个输入缓冲区的指针a和b。最后，它还包含了一个成员用于保存a和b的点积运算结果。

要使用N个GPU，我们首先需要准确地知道N值是多少。因此，在应用程序的开头调用CUDAGetDeviceCount()，从而判断在系统中安装了多少个支持CUDA的处理器。

int main(void) { int deviceCount; HANDLE_ERROR( CUDAGetDeviceCount( &deviceCount ) ); if (deviceCount < 2) { printf( "We need at least two compute 1.0 or greater " "devices, but only found  $\% \mathrm{d}\backslash \mathrm{n}^{\prime \prime}$  , deviceCount ); return 0; }

这个示例只是说明多个GPU的使用，因此如果系统只有一个CUDA设备，程序将退出（这并不是因为出现了什么错误）。显然，这种做法并不是最佳的。为了使问题尽可能简单，我们将为输入缓冲区分配标准的主机内存，并按照之前的方式来填充它们。

float \*a = (float\*)malloc( sizeof(float) \* N);   
HANDLE_NULL(a);   
float \*b  $=$  (float\*)malloc( sizeof(float) \* N);   
HANDLE_NULL(b);   
//用数据填充主机内存   
for(int  $\mathrm{i} = 0$  ：  $\mathrm{i} <   \mathrm{N}$  ；  $\mathrm{i + + })$  { a[i]  $=$  i; b[i]  $=$  i\*2;   
1

现在，我们就可以进一步深入研究多GUP代码。当通过CUDA运行时API来使用个多GPU时，要意识到每个GPU都需要由一个不同的CPU线程来控制。由于之前只使用了单个GPU，因此不需要担心这个问题。我们将多线程代码的大部分复杂性都移入到辅助代码文件book.h中。在精简了代码后，我们需要做的就是填充一个结构来执行计算。虽然在系统中可以有任意数量的GPU，但为了简单起见，在这里只使用两个：

DataStruct data[2];  
data[0].deviceID = 0;  
data[0].size = N/2;  
data[0].a = a;  
data[0].b = b;  
data[1].deviceID = 1;  
data[1].size = N/2;  
data[1].a = a + N/2;  
data[1].b = b + N/2;

我们将其中一个DataStruct变量传递给辅助函数start_thread()。此外，还将一个函数指针传递给了start_thread()，新创建的线程将调用这个函数，在这个示例中的线程函数为routine()。函

数start_thread()将创建一个新线程，这个线程将调用routine()，并将DataStruct变量作为参数传递进去。在应用程序的默认线程中也将调用routine()（因此只多创建了一个线程）。

CUTThread thread = start_thread( routine, &(data[0]) ); routine(&(data[1]);

通过调用end_thread()，主应用程序线程将等待另一个线程执行完成。

end_thread( thread );

由于这两个线程都在main()的这个位置上执行完成，因此可以安全地释放内存并显示结果。

free(a);   
free(b);   
printf("Value calculated:  $\% f\backslash n$  ", data[0].回报Value  $^+$  data[1].回报Value）; return 0;

注意，我们要将每个线程的计算结果相加起来。这是点积归约运算的最后一步。在其他算法中，可能还需要其他的步骤将这些结果合并起来。事实上，在某些应用程序中，这两个GPU可能在不同的数据集上执行不同的代码。为了简单起见，在我们点积运算示例中不考虑这些情况。

由于这里的点积函数与之前版本中的点积函数是相同的，因此在本节中不介绍它。但是，我们要注意routine()函数的代码。在声明routine()时指定该函数带有一个void参数，并返回void，这样在start_thread()部分代码保持不变的情况下可以任意实现线程函数。虽然这个想法确实不错，但这却是C中回调函数的标准过程。

void\*routine(void\*pvoidData){ DataStruct \*data  $=$  (DataStruct\*)pvoidData; HANDLE_ERROR( CUDASetDevice( data->deviceID));

每个线程都调用CUDASetDevice()，并且每个线程都传递一个不同的ID给这个函数。因此，我们知道每个线程都将使用一个不同的GPU。这些GPU可能有着相同的性能，例如双GPU的 GeForce GTX 295，或者有着不同的性能，例如在系统中同时包含一个集成GPU和一个独立GPU时。对于应用程序来说，这些细节并不重要，但它们对你来说或许是需要注意的。特别是，如果在启动核函数时需要某个最低的计算功能集版本，或者如果希望在系统的多个GPU之间实现负载均衡，那么这些细节将非常有用。如果各个GPU的性能不同，那么你需要对计算进行划分，从而使每个GPU都执行基本相等的时间。然而，在这个示例中，我们不需要担心这些细节问题。

除了调用CUDASetDevice()来指定希望使用的CUDA设备外，routine()的实现非常类似于11.1.1节中的malloc_test()。我们为输入数据和临时计算结果分别分配了内存，随后调用CUDAMemcpy()将每个输入数组复制到GPU。

int size = data->size;  
float *a, *b, c, *partial_c;  
float *dev_a, *dev_b, *dev_partial_c;  
// 在CPU上分配内存  
a = data->a;  
b = data->b;  
partial_c = (float*)malloc( blocksPerGrid*sizeof(float));

// 在GPU上分配内存  
HANDLE_ERROR(udaMalloc((void**)&dev_a, size*sizeof(float)));  
HANDLE_ERROR(udaMalloc((void**)&dev_b, size*sizeof(float)));  
HANDLE_ERROR(udaMalloc((void**)&dev_partial_c, blocksPerGrid*sizeof(float)));

// 将数组 'a' 和 'b' 复制到GPU上  
HANDLE_ERROR(udaMemcpy(dev_a, a, size*size(float),udaMemcpyHostToDevice());  
HANDLE_ERROR(udaMemcpy(dev_b, b, size*size(float),udaMemcpyHostToDevice());

然后，我们启动点积核函数，复制回计算结果，并且结束CPU上的操作。

dot<<<blocksPerGrid,threadsPerBlock>>>(size,dev_a,dev_b, dev_partial_c);

// 将数组 'c' 从GPU复制回CPU  
HANDLE_ERROR(udaMemcpy( partial_c, dev_partial_c, blocksPerGrid*size(float),udaMemcpyDeviceToHost));

//结束CPU上的操作  
c = 0;  
for (int i = 0; i < blocksPerGrid; i++) {  
    c += partial_c[i];  
}

和通常一样，接下来释放GPU内存，并在DataStruct结构的returnValue成员中返回计算好的点积结果。

HANDLE_ERROR(udaFree(dev_a)）;   
HANDLE_ERROR(udaFree(dev_b)）;   
HANDLE_ERROR(udaFree(dev_partial_c))；   
//释放CPU侧的内存 free( partial_c); data->returnValue  $= c$  return0;

因此，除了主机线程的管理问题外，使用多个GPU的程序并不比使用单个GPU的程序要复杂。如果使用我们的辅助代码来创建一个线程，并在线程上执行一个函数，那么将极大地降低编码的复杂性。如果你有自己的线程库，那么也可以在程序中使用它们。你只需要记住，每个GPU都有自己的线程，除此之外其他的使用方式都是类似的。