[attach]3338[/attach]
请问有可能实现如图所示的原子操作吗？例如：atomicCAS(&a, a, atomicCAS(&b, b, c)); ? 貌似不能保证整体是原子操作。。
如果不能实现，能不能有其他的替代方法？
楼主您好，您的问题可以分解为：
（1）是否我必须需要原子性的(a=>b, b=>c)的变换.
（2）如果我需要原子性的操作，我该怎么办。
但您不能直接限定死必须“双地址的原子操作”，这样会限制论坛为您分析问题和解答问题的途径的，希望您能接受。
回到您的问题，先分析您是否需要原子性完成此过程：
如果您的这些值都是固定的序列，那么您可以简单的进行滑窗操作，
例如说，
您有地址p, 指向的内容是p[0]=a,p[1]=b,p[2]=c (a,b,c是一些已知常量）
而现在您需要让它们变成p[0]=b,p[1]=c,
您可以简单的将p向后滑动一个元素，这样同时p指向的[0], [1]甚至更多元素都同时发生了改变。（当然，根据您的情况，您也许需要对p进行atomicAdd来+1个元素大小)
上文是说将2个元素改变，但无需1个原子性的2次赋值的过程的例子。
然后我们考虑下，如果您的算法需要原子操作保护此2次赋值您该怎么办。
最简单的思路就是让当前操作的线程在操作此2步的时候，别的线程不能干扰。
那么您可以实现一个锁（本文不说如何具体实现，只是给您提供解决思路的）
或者，
如果在一个block内的这种2步操作（例如shared上的），您可以简单的暴力暂停其他线程的执行，例如：
__syncthreads();
if (我是需要干活的线程)
{
2步赋值;
}
__syncthreads();这样。
以及还可能有其他的方式，这些都是通用的思路和例子，
希望能引起您的思考，花道。
感谢周末深夜来访。
感谢版主大神的答复！
我的程序中a,b,c在global memory中的位置是不连续的，因此滑窗操作不能实现。
另外今天尝试了一下critical section，但是出现一个新的问题：在critical section中的两次写入操作似乎不能保证完成之后再解锁，即写入操作的延迟导致这个critical section无效。无奈加上__threadfence()之后，结果程序还是不对，并且效率下降得厉害。。
最后还是不得已修改算法避免使用这种操作。