覆盖范围:第8章至
计时:
CUDA过程和核函数的执行可以通过CPU或GPU计时器进行计时。
CPU计时:
任何CPU计时器都可以被用于CUDA的计时。
很多CUDA API是异步执行的——调用后会立即返回,无论执行完成与否。故需要在启动和终止CPU计时器前,用cudaDeviceSynchronize()
来同步CPU线程和GPU。
上述的同步函数不适合用于计时在非默认流的代码。
GPU计时:
使用流程:建立start、stop event -> 记录start event -> 记录stop event -> 同步event -> 计算时差 -> 销毁两个event
cudaEvent_t start, stop; float time; cudaEventCreate(&start); cudaEventCreate(&stop); cudaEventRecord( start, 0 ); kernel<<<grid,threads>>> ( d_odata, d_idata, size_x, size_y, NUM_REPS); cudaEventRecord( stop, 0 ); cudaEventSynchronize( stop ); cudaEventElapsedTime( &time, start, stop ); cudaEventDestroy( start ); cudaEventDestroy( stop );
注意时差的单位是毫秒,分辨率大致是半个微秒。
Page-locked memory:
这种内存空间可以在Host和Device的传输中速度最快。用cudaHostAlloc()
函数分配这种内存空间。
对之前分配过的内存,使用cudaHostRegister()
能将该快内存注册成pinned memory。
不应过度使用,因为这种内存空间在全局角度是相对稀缺的资源,而什么量是过量难以事先预知。
(名词解释:stream是在device上顺序进行的一组操作。)
通过指定不同流,可以让核函数和内存拷贝过程并行执行。这种策略常用于整个数据能被分成小块,让每个核函数处理一个小块的情境下。

size=N*sizeof(float)/nStreams; for (i=0; i<nStreams; i++) { offset = i*N/nStreams; cudaMemcpyAsync(a_d+offset, a_h+offset, size, dir, stream[i]); kernel<<<N/(nThreads*nStreams), nThreads, 0, stream[i]>>>(a_d+offset); }
零拷贝:令GPU线程直接访问Host内存。需要mapped pinned memory。
Page-locked memory mapping is enabled by calling cudaSetDeviceFlags()
with cudaDeviceMapHost
.
Mapped pinned host memory allows you to overlap CPU-GPU memory transfers with computation while avoiding the use of CUDA streams. But since any repeated access to such memory areas causes repeated CPU-GPU transfers, consider creating a second area in device memory to manually cache the previously read host memory data.