单个GPU具有强大的并行计算的能力，当把多个GPU同时用来执行同一个任务的时候，计算的性能将会得到更大的提升。本文在两块GPU上实现大数据量的向量点积运算。主要步骤为：
1. 获得设备数量；
2. 任务分配；
3. 为每个任务创建一个线程；
4. 启动每个线程进行运算；
5. 合并每个GPU得到的结果。
程序代码：
主程序：main.cpp
#include "main.h"#include <stdio.h>extern "C" void runDotProduct(float *dev_a, float *dev_b, float *dev_partial_c, int size);void* worker(void *pvoidData){ GPUPlan *plan = (GPUPlan*) pvoidData; HANDLE_ERROR(cudaSetDevice(plan->deviceID)); int size = plan->size; float *a, *b, c, *partial_c; float *dev_a, *dev_b, *dev_partial_c; a = plan->a; b = plan->b; partial_c = (float*)malloc(blockPerGrid*sizeof(float)); HANDLE_ERROR(cudaMalloc((void**)&dev_a, size*sizeof(float))); HANDLE_ERROR(cudaMalloc((void**)&dev_b, size*sizeof(float))); HANDLE_ERROR(cudaMalloc((void**)&dev_partial_c, blockPerGrid*sizeof(float))); HANDLE_ERROR(cudaMemcpy(dev_a, a, size*sizeof(float), cudaMemcpyHostToDevice)); HANDLE_ERROR(cudaMemcpy(dev_b, b, size*sizeof(float), cudaMemcpyHostToDevice)); runDotProduct(dev_a, dev_b, dev_partial_c, size); HANDLE_ERROR(cudaMemcpy(partial_c, dev_partial_c, blockPerGrid*sizeof(float), cudaMemcpyDeviceToHost)); c = 0; for (int i=0; i<blockPerGrid; i++) { c += partial_c[i]; } HANDLE_ERROR(cudaFree(dev_a)); HANDLE_ERROR(cudaFree(dev_b)); HANDLE_ERROR(cudaFree(dev_partial_c)); free(partial_c); plan->returnValue = c; return 0;}int main(){ //on two GPUs int i; int deviceCount; HANDLE_ERROR(cudaGetDeviceCount(&deviceCount)); if (deviceCount < 2) { printf("No more than 2 device with compute 1.0 or greater." "only %d devices found", deviceCount); return 0; } float *a = (float*)malloc(sizeof(float)*N); HANDLE_NULL(a); float *b = (float*)malloc(sizeof(float)*N); HANDLE_NULL(b); for (i=0; i<N; i++) { a[i] = i; b[i] = i * 2; } GPUPlan plan[2]; plan[0].deviceID = 0; plan[0].size = N/2; plan[0].a = a; plan[0].b = b; plan[1].deviceID = 1; plan[1].size = N/2; plan[1].a = a + N/2; plan[1].b = b + N/2; cudaEvent_t start, stop; HANDLE_ERROR(cudaEventCreate(&start)); HANDLE_ERROR(cudaEventCreate(&stop)); float elapsedTime; HANDLE_ERROR(cudaEventRecord(start)); CUTThread mythread1 = start_thread((CUT_THREADROUTINE)worker, &plan[0]); CUTThread mythread2 = start_thread((CUT_THREADROUTINE)worker, &plan[1]); //worker(&plan[1]); end_thread(mythread1); end_thread(mythread2); HANDLE_ERROR(cudaEventRecord(stop)); HANDLE_ERROR(cudaEventSynchronize(stop)); HANDLE_ERROR(cudaEventElapsedTime(&elapsedTime, start, stop)); printf("Computing by 2 GPUs finished in %3.1f <ms>\n", elapsedTime); printf("value calculated: %f\n", plan[0].returnValue + plan[1].returnValue); HANDLE_ERROR(cudaEventDestroy(start)); HANDLE_ERROR(cudaEventDestroy(stop)); free(a); free(b); // on one GPU float *host_a; float *host_b; float *partial_c; host_a = (float*)malloc(N*sizeof(float)); host_b = (float*)malloc(N*sizeof(float)); partial_c = (float*)malloc(blockPerGrid*sizeof(float)); for (int i=0; i<N; i++) { host_a[i] = i; host_b[i] = 2 * i; } float *dev_a, *dev_b, *dev_partial_c; HANDLE_ERROR(cudaMalloc((void**)&dev_a, N*sizeof(float))); HANDLE_ERROR(cudaMalloc((void**)&dev_b, N*sizeof(float))); HANDLE_ERROR(cudaMalloc((void**)&dev_partial_c, blockPerGrid*sizeof(float))); HANDLE_ERROR(cudaMemcpy(dev_a, host_a, N*sizeof(float), cudaMemcpyHostToDevice)); HANDLE_ERROR(cudaMemcpy(dev_b, host_b, N*sizeof(float), cudaMemcpyHostToDevice)); cudaEvent_t start1, stop1; HANDLE_ERROR(cudaEventCreate(&start1)); HANDLE_ERROR(cudaEventCreate(&stop1)); HANDLE_ERROR(cudaEventRecord(start1)); runDotProduct(dev_a, dev_b, dev_partial_c, N); HANDLE_ERROR(cudaEventRecord(stop1)); HANDLE_ERROR(cudaEventSynchronize(stop1)); HANDLE_ERROR(cudaEventElapsedTime(&elapsedTime, start1, stop1)); printf("Computing by one GPU finished in %3.1f <ms>\n", elapsedTime); HANDLE_ERROR(cudaMemcpy(partial_c, dev_partial_c, blockPerGrid*sizeof(float), cudaMemcpyDeviceToHost)); float res = 0; for (int i=0; i<blockPerGrid; i++) { res += partial_c[i]; } printf("value calculated: %f\n", res); HANDLE_ERROR(cudaEventDestroy(start1)); HANDLE_ERROR(cudaEventDestroy(stop1)); HANDLE_ERROR(cudaFree(dev_a)); HANDLE_ERROR(cudaFree(dev_b)); HANDLE_ERROR(cudaFree(dev_partial_c)); free(host_a); free(host_b); free(partial_c); return 0;}核函数：kernel.cu
#define imin(a,b) (a<b?a:b)extern const int N = 33 * 1024 * 1024;extern const int threadsPerBlock = 256;extern const int blockPerGrid = imin(32, (N+threadsPerBlock-1)/threadsPerBlock);__global__ void dotProduct(float *a, float *b, float *c, int N){ __shared__ float cache[threadsPerBlock]; int tid = blockDim.x * blockIdx.x + threadIdx.x; int cacheIdx = threadIdx.x; float temp = 0; while (tid < N) { temp += a[tid] * b[tid]; tid += blockDim.x * gridDim.x; } cache[cacheIdx] = temp; __syncthreads(); int i = blockDim.x /2; while (i != 0) { if (cacheIdx < i) { cache[cacheIdx] += cache[cacheIdx+i]; } __syncthreads(); i /= 2; } if (cacheIdx == 0) { c[blockIdx.x] = cache[0]; }}extern "C" void runDotProduct(float *dev_a, float *dev_b, float *dev_partial_c, int size){ dotProduct<<<blockPerGrid, threadsPerBlock>>>(dev_a, dev_b, dev_partial_c, size);}本文试图将同样的数据在单个GPU上计算，比较计算时间来突出多GPU在计算性能上的提升。但实际情况是多GPU的计算时间却比单GPU更长。初步考虑是觉得核函数太简单，使得GPU执行的性能提升不足以弥补设备分配以及线程调度等带来的开销。所以多GPU也许更适合在大量复杂计算的场景下使用~