CUDA C++ Best Practices Guide笔记二

覆盖范围：第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.