GPU 高性能编程 CUDA : 常量内存与事件

常量内存：

CUDA C 可以使用全局内存，共享内存和常量内存

常量内存用于保存在核函数执行期间不会发生变化的数据，64kb

常量内存的声明方法和共享内存是类似的，在变量前加上 __constant__ 修饰符，常量内存不需要先定一个指针，然后使用 cudaMalloc() 来为指针分配 GPU 内存，当我们将其修改为常量内存的时候，不在需要调用 cudaMalloc() 函数或者 cudaFree() ，而是在编译的时候为其提交一个固定的大小，使用 cudaMemcpyToSymbol() 来将数据复制到常量内存中。

常量内存把变量的访问限制为只读，与全局内存中读取数据相比，从常量内存中读取相同的数据可以节约内存带宽：

• 对常量内存的单词读操作可以广播到其他的“邻近”线程，这将节约15次读取操作
• 常量内存的数据将缓存起来，因此对相同的地址的连续读操作将不会产生额外的内存通信量

Warp（线程束）是指一个包含32个线程的集合。在程序中的每一行，这个线程束中的每个线程都将在不同的数据上执行相同的指令

在处理常量内存时，NVIDIA 硬件将把单次内存读取操作广播到每个半线程束，即16个线程， 如果半个线程束的每个线程都从常量内存中读取数据，那么 GPU 只会产生一次读取请求并在随后将数据传播到每个线程。这将大约减少百分之九十四的带宽，由于这块内存时不会发生变化的，所以硬件将主动把这个常量数据缓存在 GPU 上，在第一次从常量内存的某个地址上读取后，当其他半线程束请求同一个地址时，那么将命中缓存，同样减少了额外的内存流量

只有当16个线程每次都需要相同的数据请求时，才值得将这个读取操作广播到16个线程，如果半线程束中所有16个线程需要访问常量内存中不同的数据，那么这16次不同的读取操作将会被串行化，需要16倍的时间，但如果从全局内存中读取，那么这些请求会同时发出，这样，全局就快于常量内存

使用事件来测试性能：

cudaEvent_t start,stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start,0);
//在 GPU 上执行一些工作
cudaEventRecord(stop,0);     

为了克服 GPU 和 CPU 的异步计算造成计时的困难问题，使用 cudaEventSynchronize( stop ) ，来阻塞后面的语句，知道 GPU 执行到达 stop 事件，当 cudaEventSynchronize 返回时，我们知道 stop 事件之前的所有 GPU 工作已经完成了

float elapsedTime;
cudaEventElapsedTime(&elapsedTime,start,stop);
cudaEventDestroy(start);
cudaEventDestroy(stop);

cudaEventElapsedTime() 是一个工具函数，用来计算两个事件之间经历的时间，第一个参数是浮点变量的地址，单位为毫秒。当使用完时间后需要销毁他们，使用 cudaEventDestroy() ，每一个 cudaEventVCreate() 都对应一个 cudaEventDestroy()