首页 > 代码库 > Cuda learn record one

Cuda learn record one

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\n", cudaGetErrorString( err ),
 6                 file, line );
 7         exit( EXIT_FAILURE );
 8     }
 9 }
10 #define HANDLE_ERROR( err ) (HandleError( err, __FILE__, __LINE__ ))

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)——另外一种形式的只读内存。

 

Cuda learn record one