这一系列的文章是CUDA5.5样例代码的阅读笔记，每篇文章针对某一特定的样例代码。
本篇文章所涉及的项目是simpleMultiGPU。
项目所在的位置：0_Simple/simpleMultiGPU
源程序的分析
首先我们查看该程序的内核函数（kernel function），以了解程序中GPU所完成的核心功能：
__global__ static void reduceKernel(float *d_Result, float *d_Input, int N) {
const int tid = blockIdx.x * blockDim.x + threadIdx.x;
const int threadN = gridDim.x * blockDim.x;
float sum = 0;
for (int post = tid; pos < N; pos += threadN)
sum += d_Input[pos];
d_Result[tid] = sum;
}
可以看出，threadN代表的是当前GPU中运行的总的线程数，每个线程所完成的工作，是统计数组中所有下标index满足index % threadN == tid的元素的累加和。
然后我们定位到main函数，我们首先查看main函数的前数十行：
// Solver config
TGPUplan plan[MAX_GPU_COUNT];
// GPU reduction results
float h_SumGPU[MAX_GPU_COUNT];
float sumGPU;
double sumCPU, diff;
int i, j, gpuBase, GPU_N;
const int BLOCK_N = 32;
const int THREAD_N = 256;
const int ACCUM_N = BLOCK_N * THREAD_N;
printf("Starting simpleMultiGPU\n");
checkCudaErrors(cudaGetDeviceCount(&GPU_N));
if (GPU_N > MAX_GPU_COUNT)
GPU_N = MAX_GPU_COUNT;
printf("CUDA-capable device count: %i\n", GPU_N);
可以看到，程序首先建立了MAX_GPU_COUNT个类型为TGPUplan的plan，MAX_GPU_COUNT是程序中定义的一个常量：
const int MAX_GPU_COUNT = 32;
而TGPUplan结构体的引入主要是为了方便GPU计算数据的管理，TGPUplan结构体的定义如下：
typedef struct {
// Host-side input data
int dataN;
float *h_Data;
// Partial sum for this GPU
float *h_Sum;
// Device buffers
float *d_Data, *d_Sum;
// Reduction copied back from GPU
float *h_Sum_from_device;
// Stream for asynchronous command execution
cudaStream_t stream;
} TGPUplan;
结构体中各个成员变量的作用如注释所示。
BLOCK_N代表一个GPU网格（Grid）中的Block数目，THREAD_N代表一个Block中的Thread数。
cudaGetDeviceCount()函数用于获取系统中可用GPU数目。
main函数接下来的代码如下：
printf("Generating input data...\n\n");
// Subdividing input data across GPUs
// Get data sizes for each GPU
for (i = 0; i < GPU_N; ++i) {
planpi].dataN = DATA_N / GPU_N;
}
// Take into account "odd" data sizes
for (i = 0; i < DATA_N % GPU_N; ++i) {
plan[i].dataN++;
}
// Assign data ranges to GPUs
gpuBase = 0;
for (i = 0; i < GPU_N; ++i) {
plan[i].h_Sum = h_SumGPU + i;
gpuBase += plan[i].dataN;
}
// Create streams for issuing GPU command asynchronously and allocate memory (GPU and System page-locked)
for (i = 0; i < GPU_N; ++i) {
checkCudaErrors(cudaSetDevice(i));
checkCudaErrors(cudaStreamCreate(&plan[i].stream));
// Allocate memory
checkCudaErrors(cudaMalloc((void **)&plan[i].d_Data, plan[i].dataN * sizeof(float)));
checkCudaErrors(cudaMalloc((void **)&plan[i].d_Sum, ACCUM_N * sizeof(float)));
checkCudaErrors(cudaMallocHost((void **)&plan[i].h_Sum_from_device, ACCUM_N * sizeof(float)));
checkCudaErrors(cudaMallocHost((void **)&plan[i].h_Data, plan[i].dataN * sizeof(float)));
for (j = 0; j < plan[i].dataN; ++j)
plan[i].h_Data[j] = (float)rand() / (float)RAND_MAX;
}
总的数据量大小是DATA_N，它也是一个常量，在main函数之前定义。这段代码首先将数组中的所有数据尽可能均匀地分配给可利用的每一个GPU，然后为数组每一个元素产生一个随机小数：
for (j = 0; j < plan[i].dataN; ++j)
plan[i].h_Data[j] = (float)rand() / (float)RAND_MAX;
同时这段代码也完成了内存分配工作。值得注意的是cudaMallocHost()函数，该函数将在主机端分配一块锁页内存(page-locked)，这一类型的内存专门用于cuda异步数据拷贝。（数据拷贝和内核启动可以并行进行）
main函数中剩下的代码的功能就是运算和比较结果正确性了，运算部分的代码如下：
// Start timing and compute on GPU(s)
printf("Computing with %d GPUs...\n", GPU_N);
StartTimer();
// Copy data to GPU, launch the kernel and copy data back. All asynchronously
for (i = 0; i < GPU_N; ++i) {
// Set device
checkCudaErrors(cudaSetDevice(i));
// Copy input data from CPU
checkCudaErrors(cudaMemcpyAsync(plan[i].d_Data, plan[i].h_Data, plan[i].dataN * sizeof(float), cudaMemcpyHostToDevice, plan[i].stream));
// Perform GPU computations
reduceKernel<<<BLOCK_N, THREAD_N, 0, plan[i].stream>>> (plan[i].d_sum, plan[i].d_Data, plan[i].dataN);
getLastCudaError("reduceKernel() execution failed.\n");
// Read back GPU results
checkCudaErrors(cudaMemcpyAsync(plan[i].h_Sum_from_device, plan[i].d_Sum, ACCUM_N * sizeof(float), cudaMemcpyDeviceToHost, plan[i].stream));
}
// Process GPU results
for (i = 0; i < GPU_N; ++i) {
float sum;
// Set device
checkCudaErrors(cudaSetDevice(i));
// Wait for all operations to finish
cudaStreamSynchronize(plan[i].stream);
// Finalize GPU reduction for current subvector
sum = 0;
for (j = 0; j < ACCUM_N; ++j)
sum += plan[i].h_Sum_from_device[j];
*(plan[i].h_Sum) = (float)sum;
// Shut down this GPU
checkCudaErrors(cudaFreeHost(plan[i].h_Sum_from_device));
checkCudaErrors(cudaFree(plan[i].d_Sum));
checkCudaErrors(cudaFree(plan[i].d_Data));
checkCudaErrors(cudaStreamDestroy(plan[i].stream));
}
sumGPU = 0;
for (i = 0; i < GPU_N; ++i)
sumGPU += h_SumGPU[i];
printf(" GPU Processing time: %f (ms)\n\n", GetTimer());
每个GPU都分别执行它们的内核函数，得到ACCUM_N个部分和，然后在CPU端将这些部分和叠加起来，便得到了该GPU上计算所得的部分和h_Sum.
有两个函数值得注意：
1、getLastCudaError()：该函数在helper_functions.h头文件中定义，主要功能是封装cudaGetLastError()函数并打印该函数的执行诊断错误信息。
2、cudaStreamSynchronize()：该函数在指定流上执行同步，等待指定流执行完成之后再继续执行后面的代码。
CPU端的结果比较代码这里就不再细说了~有兴趣的同学可以自己查看剩余部分的代码~