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