标签:
1.shared memory
__shared__ 声明为共享内存,将会保存在共享内存中
2.constant memory
__constant__ 声明为常量内存,将会保存在常量内存中,常量内存是只读内存,声明时要静态的分配空间
将数据从CPU拷贝到常量内存中时用cudaMemcpyToSymbol,例如cudaMemcpyToSymbol( s, temp_s,sizeof(Sphere) * SPHERES)
常量内存带来性能提升的原因:
1.对常量内存的单次读操作可以广播到临近线程,将节约15次读操作
2.常量内存的数据将缓存起来,对相同地址的连续读操作将不会产生额外的内存通信量
当处理常量内存时,NVIDIA硬件将单次内存读操作广播到每个半线程束(线程束中线程的一半),如果半线程束中的每个线程都从常量内存的相同地址上读取数据,那么使用常量内存产生的内存流量将会是使用全局内存的1/16。但是当所有的线程读取不同的地址时,会降低性能,因为若半线程束中的16个线程访问常量内存中的不同数据时,这16次不同的读取操作将会被串行化,从而需要16倍的时间来发出请求,但在全局内存中将会同时发出请求。
标签:
原文地址:http://www.cnblogs.com/shrimp-can/p/5052305.html