CUDA流与事件：并发处理与性能优化

### CUDA流与事件：并发处理与性能优化 在CUDA编程中，流和事件是实现高效并发处理和性能优化的重要工具。下面将详细介绍如何利用这些工具来提高GPU计算性能。 #### 1. CUDA事件计时 CUDA事件可用于对GPU操作进行计时。CUDA驱动通过检查一个32位的值来实现`cuEventQuery()`和`cuEventSynchronize()`函数。此外，GPU还能写入一个64位的计时器值，该值源自基于GPU的高分辨率时钟。 使用CUDA事件计时具有以下优点： - 受系统事件（如页面错误或中断）的干扰较小，因为使用的是基于GPU的时钟。 - 计算时间戳之间经过时间的函数在所有操作系统上都是可移植的。 不过，CUDA事件最好用于针对性地调整内核或其他GPU密集型操作，而不是向用户报告绝对时间。 在使用CUDA事件进行计时时，建议将事件记录在NULL流中，以消除计时操作的歧义。例如： ```cpp CUDART_CHECK( cudaEventRecord( startEvent, NULL ) ); // synchronous memcpy – invalidates CUDA event timing CUDART_CHECK( cudaMemcpy( deviceIn, hostIn, N*sizeof(int) ); CUDART_CHECK( cudaEventRecord( stopEvent, NULL ) ); ``` #### 2. 并发复制与内核处理 CUDA应用程序需要通过PCI Express总线传输数据，这为并发执行主机 - 设备内存传输和内核处理提供了性能优化的机会。根据阿姆达尔定律，使用多个处理器可实现的最大加速比为： \[ \text{Speedup} = \frac{1}{r_s + \frac{r_p}{N}} \] 其中，$r_s + r_p = 1$，$N$是处理器的数量。在并发复制和内核处理的情况下，“处理器数量”指的是GPU中的自主硬件单元数量，即一个或两个复制引擎加上执行内核的SM。 理论上，对于具有一个复制引擎的GPU，如果程序能使SM和复制引擎完美重叠，并且数据传输和处理时间相等，那么性能可以提高2倍。但在实际应用中，需要仔细评估这种并发处理是否对应用程序有益。对于极端传输受限或极端计算受限的应用程序，重叠传输和计算可能不会带来明显的性能提升。 ##### 2.1 concurrencyMemcpyKernel.cu程序 该程序用于演示如何实现并发内存复制和内核执行，并确定是否值得这样做。程序中的`AddKernel()`是一个“伪工作”内核，通过`cycles`参数控制其运行时间： ```cpp __global__ void AddKernel( int *out, const int *in, size_t N, int addValue, int cycles ) { for ( size_t i = blockIdx.x*blockDim.x+threadIdx.x; i < N; i += blockDim.x*gridDim.x ) { volatile int value = in[i]; for ( int j = 0; j < cycles; j++ ) { value += addValue; } out[i] = value; } } ``` 程序中使用两个例程来测量`AddKernel()`的性能： - `TimeSequentialMemcpyKernel()`：按顺序执行输入数据复制到GPU、调用`AddKernel()`和将输出数据从GPU复制回主机的操作。 - `TimeConcurrentOperations()`：分配多个CUDA流，并并行执行主机 - 设备内存复制、内核处理和设备 - 主机内存复制操作。 `TimeSequentialMemcpyKernel()`函数使用四个CUDA事件分别计时主机 - 设备内存复制、内核处理和设备 - 主机内存复制的时间，并报告总时间： ```cpp bool TimeSequentialMemcpyKernel( float *timesHtoD, float *timesKernel, float *timesDtoH, float *timesTotal, size_t N, const chShmooRange& cyclesRange, int numBlocks ) { cudaError_t status; bool ret = false; int *hostIn = 0; int *hostOut = 0; int *deviceIn = 0; int *deviceOut = 0; const int numEvents = 4; cudaEvent_t events[numEvents]; for ( int i = 0; i < numEvents; i++ ) { events[i] = NULL; CUDART_CHECK( cudaEventCreate( &events[i] ) ); } cudaMallocHost( &hostIn, N*sizeof(int) ); cudaMallocHost( &hostOut, N*sizeof(int) ); cudaMalloc( &deviceIn, N*sizeof(int) ); cudaMalloc( &deviceOut, N*sizeof(int) ); for ( size_t i = 0; i < N; i++ ) { hostIn[i] = rand(); } cudaDeviceSynchronize(); for ( chShmooIterator cycles(cyclesRange); cycles; cycles++ ) { printf( "." ); fflush( stdout ); cudaEventRecord( events[0], NULL ); cudaMemcpyAsync( deviceIn, hostIn, N*sizeof(int), cudaMemcpyHostToDevice, NULL ); cudaEventRecord( events[1], NULL ); AddKernel<<<numBlocks, 256>>>( deviceOut, deviceIn, N, 0xcc, *cycles ); cudaEventRecord( events[2], NULL ); cudaMemcpyAsync( hostOut, deviceOut, N*sizeof(int), cudaMemcpyDeviceToHost, NULL ); cudaEventRecord( events[3], NULL ); cudaDeviceSynchronize(); cudaEventElapsedTime( timesHtoD, events[0], events[1] ); cudaEventElapsedTime( timesKernel, events[1], events[2] ); cudaEventElapsedTime( timesDtoH, events[2], events[3] ); cudaEventElapsedTime( timesTotal, events[0], events[3] ); timesHtoD += 1; timesKernel += 1; timesDtoH += 1; timesTotal += 1; } ret = true; Error: for ( int i = 0; i < numEvents; i++ ) { cudaEventDestroy( events[i] ); } cudaFree( deviceIn ); cudaFree( deviceOut ); cudaFreeHost( hostOut ); cudaFreeHost( hostIn ); return ret; } ``` 在EC2的cg1.4xlarge实例上，当`cycles`值从4到64时，各操作的时间（单位：ms）如下表所示： | CYCLES | HTOD | KERNEL | DTOH | TOTAL | | ---- | ---- | ---- | ---- | ---- | | 4 | 89.19 | 11.03 | 82.03 | 182.25 | | 8 | 89.16 | 17.58 | 82.03 | 188.76 | | 12 | 89.15 | 24.10 | 82.03 | 195.28 | | 16 | 89.15 | 30.57 | 82.03 | 201.74 | | 20 | 89.14 | 37.03 | 82.03 | 208.21 | | 24 | 89.16 | 43.46 | 82.03 | 214.65 | | 28 | 89.16 | 49.90 | 82.03