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