CUDA工作记录3

这次主要记录cuda线程的同步问题

1. 全局变量同步
   当不同block同时读写一个全局变量时，很容易出现前面线程修改而后面线程不知道的情况，这时候的计算结果往往是错误的，因为cuda为了加速计算会把一些内存较小的值存在L1缓存里面，第一个线程计算后写入的值不会更新其他block的线程的L1缓存。
   解决办法就是使用原子操作atomicMax，atomicCAS…等，cuda中的原子操作是作用在L2缓存上的，会直接绕过L1。
   另外还有volatile关键字，这个也会让变量的读写绕过L1。

2. __threadfence()
   在我实际工作中会发现即使使用原子操作或者volatile关键字，还是会出现意外情况。比如cuda的CAS锁

 int* lock;
 cudaMalloc(&lock,sizeof(int));
 cudaMemset(lock,0,sizeof(int));
 bool loop = true;
 while( loop ){
 int mark = atomicCAS(lock , 0 , 1);
 if(mark == 0)
 {
 loop = false;
 // do ...
 atomicExch(lock,0);
 }
 }

有时候这里会出现死锁，我分析是因为第一个线程lock的解锁没有同步到其他的线程上，导致其它线程mark的值一直为1，无法跳出while循环。
解决方法是在修改lock的值后添加__threadfence()函数。使得lock的修改能及时被其它线程知道。

 int* lock;
 cudaMalloc(&lock,sizeof(int));
 cudaMemset(lock,0,sizeof(int));
 bool loop = true;
 while( loop ){
 int mark = atomicCAS(lock , 0 , 1);
 if(mark == 0)
 {
 loop = false;
 // do ...
 atomicExch(lock,0);
 }
 }
 __threadfence();