9.5_单线程多GPU方案

9.5 单线程多GPU方案

当使用CUDA运行时时，一个单线程的应用程序可以驱动多个GPU。CPU线程通过调用CUDASetDevice()来指定欲控制的GPU。这一模式语句用在代码清单9-1中，对源和目标GPU进行切换。该模式不仅适用于点对点内存复制，也适用于9.5.2节中描述的N-体的单线程多GPU实现。在驱动程序API中，CUDA维护一个当前上下文栈，使子程序可以很容易地改变和重新恢复调用者的当前上下文。

9.5.1 当前上下文栈

驱动程序API的应用程序可以使用当前上下文栈管理当前上下文：cuCtxPushCurrent()函数使得一个新上下文成为当前上下文，并把它压入栈的顶部；cuCtxPopCurrent()弹出当前上下文，并恢复前一个当前上下文。代码清单9-2给出驱动程序API版本的chMemcpyPeerToPeer()，它使用cuCtxPopCurrent()和cuCtxPushCurrent()在两个上下文之间执行点对点的内存复制。

当前上下文栈最初在CUDA 2.2中引入。当时CUDA运行时和驱动程序API不能在同一应用程序中配合使用。该限制已在后续版本中放宽。

代码清单9-2 chMemcpyPeerToPeer(驱动程序API版本)

CResult
chMemcpyPeerToPeer(
void *_dst, CUcontext dstContext, int dstDevice,
const void *_src, CUcontext srcContext, int srcDevice,
size_t N)
{
CUresult status;
CUdeviceptr dst = (CUdeviceptr) (intptr_t)_dst;
CUdeviceptr src = (CUdeviceptr) (intptr_t)_src;
int stagingIndex = 0;
while(N) {
size_t thisCopySize = min(N, STAGING_BUFFER_SIZE);
CUDA_CHECK( cuCtxPushCurrent(srcContext) );
CUDA_CHECK( cuStreamWaitEvent(NULL, g_events[dstDevice][stagingIndex], 0);
CUDA_CHECK( cuMemcpyDtoHAsync(g_hostBuffers[stagingIndex],
src,
thisCopySize,
NULL) );
CUDA_CHECK( cuEventRecord(g_events[srcDevice][stagingIndex],
0) );
CUDA_CHECK( cuCtxPopCurrent(&srcContext) );
CUDA_CHECK( cuCtxPushCurrent(srcContext) );
CUDA_CHECK( cuStreamWaitEvent(NULL, g_events[srcDevice][stagingIndex],
0) );
CUDA_CHECK( cuMemcpyHtoDAsync(dst, g_hostBuffers[stagingIndex],
thisCopySize,
NULL) );
CUDA_CHECK( cuEventRecord(g_events[srcDevice][stagingIndex],
0) );
CUDA_CHECK( cuCtxPopCurrent(&dstContext) );
dst += thisCopySize;
src += thisCopySize;
N -= thisCopySize;
stagingIndex = 1 - stagingIndex;
}
// Wait until both devices are done
CUDA_CHECK( cuCtxPushCurrent(srcContext) );
CUDA_CHECK( cuCtxSynchronize()) ;
CUDA_CHECK( cuCtxPopCurrent(&srcContext) );
CUDA_CHECK( cuCtxPushCurrent.dstContext) ;
CUDA_CHECK( cuCtxSynchronize()) ;
CUDA_CHECK( cuCtxPopCurrent(&dstContext) );
Error:
return status;
}

9.5.2 N-体问题

N-体计算问题（在第14章中详细描述）以 O(N2)O(N^{2}) 的时间复杂度计算N个作用力。其输出可独立地计算。在k个GPU构成的系统上，我们的多GPU实现把计算任务分割成k个部分。

我们的实现均假设所有GPU是相同的，所以我们可以均匀地划分计算任务。针对那些采用GPU性能不均匀的应用程序，或者那些工作量运行时较难预测的应用程序，可以更精细地划分计算任务，并让主机代码从一个任务队列中提交工作片段到GPU。

代码清单9-3对代码清单14-3作了修改。它有两个额外的参数（一个基索引base和作用力的子数组的大小为n），来计算N-体问题输出数组的一个子集。这个__device__函数被一个声明为__global__的外层内核调用。它组织成这种方式是为了重用代码而不发生链接错误。如果函数声明为__global__，链接器会产生一个符号冗余的错误[1]。

代码清单9-3 N-体问题的内核函数（多GPU配置）

inline _device _void   
ComputeNBodyGravitation_Shared.MultiGPU( float \*force, float \*posMass, float softeningSquared, size_t base, size_t n, size_t N) { float4 \*posMass4  $=$  (float4 \*) posMass; extern _shared _float4 shPosMass[]; for ( int m = blockIdx.x\*blockDim.x + threadIdx.x; m < n; m += blockDim.x\*gridDim.x ) { size_t i = base+m; float acc[3]  $= \{0\}$  ; float4 myPosMass  $=$  posMass4[i]; #pragma unroll 32 for ( int j = 0; j < N; j ++ blockDim.x ) { shPosMass [threadIdx.x]  $=$  posMass4[j+threadIdx.x]; _syncthreads(); for ( size_t k = 0; k < blockDim.x; k++) { float fx, fy, fz; float4 bodyPosMass  $=$  shPosMass[k]; bodyBodyInteraction{ &fx, &fy, &fz, myPosMass.x, myPosMass.y, myPosMass.z, bodyPosMass.x, bodyPosMass.y, bodyPosMass.z, bodyPosMass.w, softeningSquared ); acc[0]  $+=$  fx; acc[1]  $+=$  fy; acc[2]  $+=$  fz; } _syncthreads(); } force[3*m+0]  $=$  acc[0]; force[3*m+1]  $=$  acc[1]; force[3*m+2]  $=$  acc[2]; }

代码清单9-4给出了N-体问题的单线程多GPU版本的主机代码[2]。dptrPosMass和dptrForce数组为每个GPU的输入及输出数组跟踪设备指针（GPU的最大数量在nbody.h中声明为常数；默认是32）。跟分派工作到CUDA流的过程相似，函数在计算的不同阶段使用不同的循环：第一个循环为每个GPU分配并填充输入数组；第二个循环启动内核和输出数据的异步复制；第三个循环依次在每个GPU调用

CUDADeviceSynchronize()函数。这样组织函数可以最大化CPU/GPU重叠。在第一个循环中，当CPU忙着为第i个GPU分配内存之际，前i个GPU（第0～i-1个GPU）能够执行主机到设备的异步内存复制操作。如果内核启动和设备到主机的异步复制在第一个循环中调用同步CUDAAlloc()，会降低性能，因为它们要与当前GPU进行同步。

代码清单9-4 N-体问题的主机代码（单线程多GPU配置）

float   
ComputeGravitation-multiGPU_singlethread( float \*force, float \*posMass, float softeningSquared, size_t N   
}   
{udaError_t status; float ret  $= 0$  .of; float \*dptrPosMass[g_maxGPUs]; float \*dptrForce[g_maxGPUs]; chTimerTimestamp start, end; chTimerGetTime(&start); memset(dptrPosMass,0,sizeof(dptrPosMass)); memset(dptrForce,0,sizeof(dptrForce)); size_t bodiesPerGPU  $= \mathbb{N}$  /g_numGPUs; if（(0！  $=$  N%g_numGPUs）|（g_numGPUs>g_maxGPUs））{ return 0.of; } // kick off the asynchronous memcpy's - overlap GPUs pulling // host memory with the CPU time needed to do the memory // allocations. for (int i  $= 0$  ;i< g_numGPUs;i++) { CUDASetDevice(i);
cudamalloc(&ptrPosMass[i],4*N*size(float));   
cudamalloc(&ptrForce[i],3*bodiesPerGPU*size(float));   
cudamemcpyAsync( dptrPosMass[i], g_hostAOS_PosMass, 4*N*size(float), CUDAMemcpyHostToDevice);   
}   
for (int i = 0; i < g_numGPUs; i++) {   
    CUDASetDevice(i);}   
    ComputeNBodyGravitation_Shared_device<< 300,256,256*size(float4) >>(); dptrForce[i], dptrPosMass[i], softeningSquared, i*bodiesPerGPU, bodiesPerGPU, N);   
    CUDAMemcpyAsync( g_hostAOS_Force+3*bodiesPerGPU*i, dptrForce[i], 3*bodiesPerGPU*size(float), CUDAMemcpyDeviceToHost );   
}   
// Synchronize with each GPU in turn.   
for (int i = 0; i < g_numGPUs; i++) { CUDASetDevice(i); CUDADeviceSynchronize(); } chTimerGetTime (&end); ret = chTimerElapsedTime (&start, &end) * 1000.of;   
Error:   
    for (int i = 0; i < g_numGPUs; i++) { CUDAFree(dptrPosMass[i]); CUDAFree(dptrForce[i]); } return ret;

[1] 这是一个传统的解决方法。CUDA5.0新增的链接器，使__global__函数可以编译为静态链接库并链接到应用程序。
[2] 为版面美观，这里的错误检查代码已被移除。

Previous9.4_多GPU间同步
Next9.6_多线程多GPU方案