由于G8x硬件不支持原子操作，实现的原子函数如下：
这是我看文章的源代码：device void addData256( volatile unsigned int *s_WarpHist, unsigned int data, unsigned int threadTag )
{ unsigned int count;
do{ count = s_WarpHist[data] & 0x07FFFFFFU;
count = threadTag | (count + 1);
s_WarpHist[data] = count;
}while(s_WarpHist[data] != count);
}
这个程序共享存储器的组织形式：行：WARP_N（每个线程块的warp个数） 列：灰度个数（每行记录一个warp线程指向像素的灰度的分布， 每个warp有自己的灰度分布计数数组，这里warp是怎么实现原子操作的呢？）。原文的解释是这样的，但是还是不懂，尤其是do…while循环中的第二行“位或”，这里得出的count代表什么了，还有这里的volatile有什么作用，红字具体指的是什么操作？请版主解惑：
addData256() is the core of the 256-bin histogram implementation. Let’s describe its logic in detail. According to data value (lying within 0 … 255 range), read from global memory, each warp thread must increment corresponding value in the s_WarpHistp array – a “frame”(row) within s_Hist array, corresponding to current warp. Each warp thread reads current warp counter s_WarpHist[data], corresponding to data value, then locally increments, tags it by warp-local thread ID (equal to threadIdx.x % 32), and writes it back to the same s_WarpHist[data] position. In case each warp thread received unique data values (from global memory), there are no collisions at all and no additional actions need to be done. Otherwise, when two or more threads collide on the same bin counter, the hardware performs shared memory write combining, resulting in acceptance of the tagged counter from one thread and rejection from all other pending threads. After the write attempt each thread queries the shared memory count value (the same s_WarpHist[data]) and owing to the tag decides whether its pending increment made its way to shared memory. If true, it becomes idle (masked out by hardware) until the entire warp is done (all the collisions are resolved). Otherwise, some other thread has submitted its increment into s_WarpHist, and current thread needs to grab the new counter value and perform the same actions. Since each warp is isolated and warp threads are always synchronized we do not rely on warp scheduling order (which is undefined). Not more than after 32 loop iterations all the warp threads submit their increments into s_WarpHist.