1. GPU 有完善的内存管理的机制,会强制结束任何违反内存访问规则的进程,但是无法阻止应用程序的继续执行,因而,错误处理函数非常重要。
1 static void HandleError( cudaError_t err, 2 const char *file, 3 int line ) { 4 if (err != cudaSuccess) { 5 printf( "%s in %s at line %d ", cudaGetErrorString( err ), 6 file, line ); 7 exit( EXIT_FAILURE ); 8 } 9 } 10 #define HANDLE_ERROR( err ) (HandleError( err, __FILE__, __LINE__ ))
值得注意的事,这种处理方式中,并没有处理动态申请的内存,这是由main()函数结束之后由操作系统负责处理的部分。这部分规则是,
a. 就算没有free(),main()结束后也是会自动释放malloc()的内存的,这里监控者是操作系统,设计严谨的操作系统会登记每一块给每一个应用程序分配的内存,这使得它能够在应用程序本身失控的情况下仍然做到有效地回收内存。你可以试一下在TaskManager里强行结束你的程序,这样显然是没有执行程序自身的free()操作的,但内存并没有发生泄漏。
b. free()的用处在于实时回收内存。如果你的程序很简单,那么你不写free()也没关系,在你的程序结束之前你不会用掉很多内存,不会降低系统性能;而你的程序结束之后,操作系统会替你完成这个工作。但你开始开发大型程序之后就会发现,不写free()的后果是很严重的。很可能你在程序中要重复10k次分配10M的内存,如果每次使用完内存后都用free()释放,你的程序只需要占用10M内存就能运行;但如果你不用free(),那么你的程序结束之前就会吃掉100G的内存。这其中当然包括绝大部分的虚拟内存,而由于虚拟内存的操作是要读写磁盘,因此极大地影响系统的性能。你的系统很可能因此而崩溃。
c. 任何时候都为每一个malloc()写一个对应的free()是一个良好的编程习惯。这不但体现在处理大程序时的必要性上,更体现在程序的优良的风格和健壮性上。毕竟只有你自己的程序知道你为哪些操作分配了哪些内存以及什么时候不再需要这些内存。因此,这些内存当然最好由你自己的程序来回收。
2.
1 add<<<blockPerGrid, threadPerGrid>>> ( void )
int i=blockIdx.x
int i=threadIdx.x
这里面,尖括号里面的参数,第一个表示一个Grid里面线程块的大小,第二个表示一个线程块里面线程的大小。
(这部分有待验证,较老的规则)-----线程块数量上限为:65535. 每个线程块中的线程的数量也有限制,这个数值为 maxThreadPerBlock ,一般是512个值。
1 int i=threadIdx.x+blockIdx.x*blockDim.x
3. 共享内存和同步
1 __shared__ float cache [threadPerBlock];
上述共享内存缓存区,
同步线程:
1 __syncthreads();
保证每一个程序都已经执行完了该语句之前的程序。
4. GPU的计算瓶颈不在于芯片的数学计算吞吐量,而在于芯片的内存带宽。输入的数据无法维持较高的计算速率。GPU的内存种类包括:全局内存,共享内存,常量内存(64KB)。
5. 常量内存,在GPU中,数据不可修改的内存空间。其生成与共享内存相似,标识符为:__constant__ ,在声明常量内存的时候,需要指明常量内存的大小。
1 __constant__ Sphere* s[Num]; 2 cudaMemcpyToSymble(s, temp_s, Num*sizeof(Sphere));
注意上面不需要使用cudaMalloc和cudaFree()。其访问限制为只读。
优点:a.对常量内存的连续读写不额外的消耗时间; b. 对常量内存的单次操作将会广播到其他的邻近的区域,这将会节约15次的读取时间。
6. 线程束(warp)
包含一定数量的线程(通常是32个),捆绑在一起,步调一致的执行程序,在不同的数据集上执行相同的指令。读取常量内存的时候,GPU不仅会缓存这些数据,而且将这些常量数据广播到半个warp(大约是16个线程),然而,如果,不同的线程所操作的数据是不一致的,这将会导师常量内存的读取速度比全局变量的读取速度更慢。
7. 使用事件来测试性能。
使用cuda的事件API来测试某一段cuda代码的运行时间,注意,在创建事件之后,一定要相应的销毁事件。
这里可以统计使用某一段代码进行优化之后的效果,特别是当使用了常量内存的时候。常量内存在某些情况下,特比的节省事件,作者统计算例优化百分之五十之上。当所有的线程访问相同 的变量的时候。
cudaEvent_t start, stop; cudaEventCreat(&start); cudaEventCreate(&srop); cudaEventRecord(start,0); //执行一段时间的GPU代码 cudaEventRecord(stop,0); cudaEventSynchronize(stop); //保证在sop之前,所有的GPU事件已经完成。 float elapsedTime; cudaEventElapsedTime(&elapsedTime, start, stop); cudaEventDestroy(start); cudaEvemtDestory(stop); //销毁事件
8. 注意在有可能在某些GPU仍然为计算完成时,CPU已经开始计算下一段代码。
9. 纹理内存(texture Memory)——另外一种形式的只读内存。