使用 P2P 特性在 GPU 之间传输、读写数据。
▶ 源代码。包括 P2P 使用前的各项检查，设备之间的数据互拷，主机和设备之间数据传输和相互访问。
#include <stdlib.h>
#include <stdio.h>
#include <cuda_runtime.h>
#include "device_launch_parameters.h"
#include <helper_cuda.h>
#include <helper_functions.h> #define MAX_GPU_COUNT 64 __global__ void SimpleKernel(float *src, float *dst)
{
const int idx = blockIdx.x * blockDim.x + threadIdx.x;
dst[idx] = src[idx] * 2.0f;
} inline bool IsGPUCapableP2P(cudaDeviceProp *pProp)
{
#ifdef _WIN32
return (bool)(pProp->tccDriver ? true : false);
#else
return (bool)(pProp->major >= );
#endif
} int main(int argc, char **argv)
{
printf("\n\tStarting\n", argv[]); // 检查是否使用 64 位操作系统环境
if (sizeof(void*) != )
{
printf("\n\tError for program only supported with 64-bit OS and 64-bit target\n");
return EXIT_WAIVED;
} // 找到头两块计算能力不小于 2.0 的设备
int gpu_n;
cudaGetDeviceCount(&gpu_n);
printf("\n\tDevice count: %d\n", gpu_n);
if (gpu_n < )
{
printf("\n\tError for two or more GPUs with SM2.0 required\n");
return EXIT_WAIVED;
} cudaDeviceProp prop[MAX_GPU_COUNT];
int gpuid[MAX_GPU_COUNT], gpu_count = ;
printf("\n\tShow device\n");// 展示所有设备
for (int i=; i < gpu_n; i++)
{
cudaGetDeviceProperties(&prop[i], i);
if ((prop[i].major >= )
#ifdef _WIN32
&& prop[i].tccDriver// Windows 系统还要求有 Tesla 计算集群驱动
#endif
)
gpuid[gpu_count++] = i;
printf("\n\tGPU%d = \"%15s\" ---- %s\n", i, prop[i].name, (IsGPUCapableP2P(&prop[i]) ? "YES" : "NO"));
}
if (gpu_count < )
{
printf("\n\tError for two or more GPUs with SM2.0 required\n");
#ifdef _WIN32
printf("\nOr for TCC driver required\n");
#endif
cudaSetDevice();
return EXIT_WAIVED;
} // 寻找测试设备
int can_access_peer, p2pCapableGPUs[];
p2pCapableGPUs[] = p2pCapableGPUs[] = -;
printf("\n\tShow combination of devices with P2P\n");// 展示所有能 P2P 的设备组合
for (int i = ; i < gpu_count - ; i++)
{
for (int j = i + ; j < gpu_count; j++)
{
cudaDeviceCanAccessPeer(&can_access_peer, gpuid[i], gpuid[j]);
if (can_access_peer)
{
printf("\n\tGPU%d (%s) <--> GPU%d (%s) : %s\n", gpuid[i], prop[gpuid[i]].name, gpuid[j], prop[gpuid[j]].name);
if (p2pCapableGPUs[] == -)
p2pCapableGPUs[] = gpuid[i], p2pCapableGPUs[] = gpuid[j];
}
}
}
if (p2pCapableGPUs[] == - || p2pCapableGPUs[] == -)
{
printf("\n\tError for P2P not available among GPUs\n");
for (int i=; i < gpu_count; i++)
cudaSetDevice(gpuid[i]);
return EXIT_WAIVED;
} // 使用找到的设备进行测试
gpuid[] = p2pCapableGPUs[];
gpuid[] = p2pCapableGPUs[];
printf("\n\tEnabling P2P between GPU%d and GPU%d\n", gpuid[], gpuid[]); // 启用 P2P
cudaSetDevice(gpuid[]);
cudaDeviceEnablePeerAccess(gpuid[], );
cudaSetDevice(gpuid[]);
cudaDeviceEnablePeerAccess(gpuid[], ); // 检查设备是否支持同一可视地址空间 (Unified Virtual Address Space，UVA)
if (!(prop[gpuid[]].unifiedAddressing && prop[gpuid[]].unifiedAddressing))
printf("\n\tError for GPU not support UVA\n");
return EXIT_WAIVED; // 申请内存
const size_t buf_size = * * * sizeof(float);
printf("\n\tAllocating buffers %iMB\n", int(buf_size / / ));
cudaSetDevice(gpuid[]);
float *g0;
cudaMalloc(&g0, buf_size);
cudaSetDevice(gpuid[]);
float *g1;
cudaMalloc(&g1, buf_size);
float *h0;
cudaMallocHost(&h0, buf_size); cudaEvent_t start_event, stop_event;
int eventflags = cudaEventBlockingSync;
float time_memcpy;
cudaEventCreateWithFlags(&start_event, eventflags);
cudaEventCreateWithFlags(&stop_event, eventflags);
cudaEventRecord(start_event, ); for (int i=; i<; i++)
{
// GPU 互拷
// UVA 特性下 cudaMemcpyDefault 直接根据指针（属于主机还是设备）来确定拷贝方向
if (i % == )
cudaMemcpy(g1, g0, buf_size, cudaMemcpyDefault);
else
cudaMemcpy(g0, g1, buf_size, cudaMemcpyDefault);
}
cudaEventRecord(stop_event, );
cudaEventSynchronize(stop_event);
cudaEventElapsedTime(&time_memcpy, start_event, stop_event);
printf("\n\tcudaMemcpy: %.2fGB/s\n", (100.0f * buf_size) / (1024.0f * 1024.0f * 1024.0f * (time_memcpy / 1000.0f))); for (int i=; i<buf_size / sizeof(float); i++)
h0[i] = float(i % );
cudaSetDevice(gpuid[]);
cudaMemcpy(g0, h0, buf_size, cudaMemcpyDefault); const dim3 threads(, );
const dim3 blocks((buf_size / sizeof(float)) / threads.x, ); // 使用 GPU1 读取 GPU0 的全局内存数据，计算并写入 GPU1 的全局内存
printf("\n\tRun kernel on GPU%d, reading data from GPU%d and writing to GPU%d\n", gpuid[], gpuid[], gpuid[]);
cudaSetDevice(gpuid[]);
SimpleKernel<<<blocks, threads>>>(g0, g1);
cudaDeviceSynchronize(); // 使用 GPU0 读取 GPU1 的全局内存数据，计算并写入 GPU0 的全局内存
printf("\n\tRun kernel on GPU%d, reading data from GPU%d and writing to GPU%d\n", gpuid[], gpuid[], gpuid[]);
cudaSetDevice(gpuid[]);
SimpleKernel<<<blocks, threads>>>(g1, g0);
cudaDeviceSynchronize(); // 检查结果
cudaMemcpy(h0, g0, buf_size, cudaMemcpyDefault);
int error_count = ;
for (int i=; i<buf_size / sizeof(float); i++)
{
if (h0[i] != float(i % ) * 2.0f * 2.0f)
{
printf("\n\tResult error at %i: gpu[i] = %f, cpu[i] = %f\n", i, h0[i], (float(i%)*2.0f*2.0f));
if (error_count++ > )
break;
}
} // 关闭 P2P
cudaSetDevice(gpuid[]);
cudaDeviceDisablePeerAccess(gpuid[]);
cudaSetDevice(gpuid[]);
cudaDeviceDisablePeerAccess(gpuid[]); // 回收工作
cudaFreeHost(h0);
cudaSetDevice(gpuid[]);
cudaFree(g0);
cudaSetDevice(gpuid[]);
cudaFree(g1);
cudaEventDestroy(start_event);
cudaEventDestroy(stop_event);
for (int i=; i<gpu_n; i++)
cudaSetDevice(i);
printf("\n\t%s!\n",error_count?"Test failed": "Test passed"); getchar();
return ;
}
▶ 输出结果
只有一台设备，暂无法进行测试
▶ 涨姿势：
● P2P 要求：至少两台计算能力不低于 2.0 的设备，并支持同一可视内存空间特性；计算环境不低于 CUDA 4.0；Windows 安装 Tesla 计算集群驱动。
● 使用P2P的关键步骤
// 检查两台设备之间是否能使用 P2P
int can_access_peer;
cudaDeviceCanAccessPeer(&can_access_peer, gpuid[i], gpuid[j])); // 启用 P2P
cudaSetDevice(gpuid[i]);
cudaDeviceEnablePeerAccess(gpuid[j], );
cudaSetDevice(gpuid[j];
cudaDeviceEnablePeerAccess(gpuid[i], ); // 设备间传输数据
cudaMemcpy(g1, g0, buf_size, cudaMemcpyDefault); // 关闭 P2P
cudaSetDevice(gpuid[i]);
cudaDeviceDisablePeerAccess(gpuid[i]);
cudaSetDevice(gpuid[j]);
cudaDeviceDisablePeerAccess(gpuid[j]); // cuda_runtime_api.h
extern __host__ cudaError_t CUDARTAPI cudaDeviceCanAccessPeer(int *canAccessPeer, int device, int peerDevice); extern __host__ cudaError_t CUDARTAPI cudaDeviceEnablePeerAccess(int peerDevice, unsigned int flags); extern __host__ cudaError_t CUDARTAPI cudaDeviceDisablePeerAccess(int peerDevice);
● 其他代码中的定义
// helper_string.h
#define EXIT_WAIVED 2