常量内存：
常量内存用于保存在核函数执行期间不会发生变化的数据，在变量面前添加 __constant__ 修饰符：
__constant__ Sphere s[SPHERES];
cudaMemcpyToSymbol(s, temp_s, sizeof(Sphere)*SPHERES);
这个特殊版本的cudaMemcpy()用于将主机内存复制到GPU上的常量内存。
从常量内存读取相同的数据可以节约内存带宽，主要原因：
（1）对常量内存的单次读操作可以广播到其他的“邻近”线程，这将节约15次读取操作
（2）常量内存的数据将缓存起来，因此对相同地址的连续读操作将不会产生而额外的内存通信量。
解释：
如果在半线程束中的每个线程都从常量内存的相同地址上读取数据，那么GPU只会产生一个读取请求并在随后将数据广播到每个线程。如果从常量内存中读取大量的数据，那么这种方式产生的内存流量只是使用全局内存的1/16.
另外由于这块内存的内容是不会发生变化的， 因此硬件将主动把这个常量数据缓存在GPU上。在第一次从常量内存的某个地址上读取后，当其他的半线程束请求同一个地址时，将命中缓存，这同样减少了额外的内存流量。
然而，使用常量内存也可能对性能产生负面影响。如果半线程束中的所有16个线程需要访问常量内存中不同的数据，那么这个16次不同的读取操作会被串行化，从而需要16倍的时间发出请求。但如果从全局内存中读取，那么这些请求会同时发出。这种情况下，从常量内存读取就慢于从全局内存中读取。
事件：
CUDA的事件本质上是一个GPU时间戳，这个时间戳是在用户指定的时间点上记录的。应该将cudaEventRecord()视为一条记录当前时间的语句， 并且把这条语句放入GPU的未完成队列中。因此直到GPU执行完了在调用cudaEventRecord（）之前的所有语句时，事件才会被记录下来。为了安全的读取stop值，需要告诉CPU在某个事件上同步，这个函数就是cudaEventSynchronize().当该函数返回时，代表stop事件之前的所有GPU工作已完成，stop可以安全读取。
由于CUDA事件是直接在GPU上实现的，因此不适用于同时包含设备代码和主机代码的混合代码计时，也就是说如果试图通过CUDA事件对核函数和设备内存复制之外的代码进行计时，将得到不可靠的结果。
cudaEvent_t start,stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, 0);
//在GPU上执行一些工作
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
float elapseTime;
cudaEventElapsedTime(&elapsedTime, start, stop);
cudaEventDestroy(start);
cudaEventDestroy(stop);