好久不写博客了,慢慢的感觉自己懒散好多,废话不说,看看关于原子操作吧
关于CUDA的书一本接着一本,算来看了好多本,在做CT重建中FDK,也算是有个结果了吧,不过后期优化还没有做,想看看关于CUDA的一些simple,那就从第一个开始吧,结果第一个是关于CUDA的动态并行的,也就是递归调用的,项目名字叫cdpSimplePrint,调试结果分析,发现原子操作的结果并不是自己以前认为的。
例子中:
s_uid = atomicAdd(&g_uids, 1);
其中变量声明:__shared__ int s_uid;__device__ int g_uids = 0;
我觉得s_uid返回的是s_uid=g_uids+1;只不过将线程串行,强制高速缓存刷新,跟踪,也确实有volite之类的字样,可结果却是s_uid等于0!
这让我困惑了许久,最后在CUDA_C_Programming_Guide中查到了这么一段:
int atomicAdd(int* address, int val);
unsigned int atomicAdd(unsigned int* address,
unsigned int val);
unsigned long long int atomicAdd(unsigned long long int* address,
unsigned long long int val);
float atomicAdd(float* address, float val);
reads the 32-bit or 64-bit word old located at the address address in global or shared
memory, computes (old + val), and stores the result back to memory at the same
address. These three operations are performed in one atomic transaction. The function
returns old.
好吧,原来搞了半天,我没有搞清这个函数具体怎么来使用的