形象地介绍一下我的问题：一条绳子长L，上面有很多只蚂蚁，蚂蚁的速度是随时间变化的V(t)，位置也是随时间变化的。我想统计1000个dt后，蚂蚁在绳子上的分布情况
初始条件：蚂蚁在绳子上均匀分布。给定初始速度，初始蚂蚁数为np
边界条件：蚂蚁位置超过绳子两端就掉下去了，不予以统计。
算法：蚂蚁的位置x存储在x[np]内，经过时间dt后蚂蚁位置有了新的变化：x(t+dt)=x(t)+v(t+dt)dt,
蚂蚁位置超过绳子一端就掉下去了，即x[i]>L或<0;即x[i]没用了v[i]也没用了，
我将L划分为nc个网格，每个网格在初始时刻有nnp个蚂蚁，每个block对应一个网格，每个block里有nnp个线程
现在，我用每条线程对应一个蚂蚁，进行时间上的循环：
for(t=0;t<1000dt;t+=dt)
{ move<<<nc, nnp>>>(); //统计新时刻的速度和位置
adjust<<<nc, nnp>>>(); //统计绳子上还剩的蚂蚁
}
我的问题出在adjust函数上，经过move后，有的蚂蚁掉下去了，有的蚂蚁跑到其他网格去了，一个线程对应一个蚂蚁，
int i = threadIdx.x + blockDim.x*blockIdx.x ;
if(a[i]>L||a[i]<0)
然后呢……如果将 i 线程return掉，下次运行adjust那么 i 线程不是又恢复了吗？
ps:不是要各位大侠帮我写代码，只是有个合理的并行思想就行，问题的关键就是怎么处理跑到其他网格的蚂蚁（原来网格的蚂蚁数减 1，现在所在的网格的蚂蚁数加1，另外还要移动该蚂蚁的数据在数组中的位置），和掉下去的蚂蚁……
求教……
大概说一下我的想法：
1：我觉得如果你每个block按照初始状态的分布计算一个网格的话，随着蚂蚁的爬动，网格之间的交互可能不好处理，所以我觉得您的block干脆和网格无关好了。假定有1024只蚂蚁，那么随便划分为2个还是4个block都可以，然后每个时间步，更新一次位置，把位置的数据保存起来即可，这样无论蚂蚁跑到哪里，你的block之间都无需做交换。
2：对于边界条件的处理，因为你是一个thread对应一个蚂蚁，那么但凡掉下去的蚂蚁，以后将不再处理，此时可以用判断+return杀掉此线程。
大致如上，供您参考。
欢迎莅临CUDAZONE，祝您新年编码愉快~
楼主你好。
似乎你需要的是对每一个点, 计算v(t)在t上的积分，但这个过程中途又加入了限制条件：任何时候，计算到的位置值不满足条件，就不必继续计算下去了。
而您要求的结果，看您的直接意思是求出每个点的位置值。这很好办。
但看您字里行间的隐藏的意思，是要求给出一个位置值的区间，求出其中对应的点。
不知道哪一种是您的真实心理想法。
我就说一下看上去的表现您的意图吧，这个好办，可以按点（蚂蚁）分配线程，每个线程读取当前位置，如果位置合法，那么加上dt时间内的偏移，回写当前位置。
即无需考虑您的block对应区间。这样安排可能简单点。您觉得呢？
此外，如果您有隐含的意思，例如上文所说的，欢迎跟帖表达。
其实根据你的原来代码，里面有2个kernel, 后者多少算是你在统计每个区间上的点。
如果直接按照这个说，您的“ 然后呢……如果将 i 线程return掉，下次运行adjust那么 i 线程不是又恢复了吗？” — 那么其实也没事，下次依然在adjust里测试这个边界条件，超界的不予统计，继续让他return掉即可。
统计每个CELL里面有多少个，可以用atomicAdd，每次统计前将CELL COUNT的数组置0就好了，当然原子操作的性能不是很好，不过目前来看除了原子操作，似乎没有别的方法可以了！
如果不考虑其他因素，前两位版主的算法确是最简便的，但诚如横扫千军版主所说，有一部分是我没有说出来的，就是我还要计算每个网格内的 所有属于此网格的 蚂蚁位置的总和，只怪我没有说清楚…… 所以我就用每个block对应一个网格了，如果加上上述条件，是不是只有tianyuan08版主所说的atomicadd这一种方案了？
其实和前面不矛盾的，如果你保存了每个蚂蚁的位置，自然能反求出它在哪个网格里面。所以实现你的目的是不成问题的。
您可以在更新完之后单独写个kernel实现此功能，也可以在计算的时候直接就更新了。
所以，现在剩下的问题就是，如果多线程对某个计数的量进行累加，这个自然需要使用原子操作。
思路大致如上，具体实现细节，您再想想看。
祝您编码愉快~
嗯嗯。那依然不必“每个block对应一个网格”。
你没能仔细领会tianyuan08在楼上的精神，他让你直接上atomicAdd, 是从你的点（蚂蚁）的角度出发的。每个线程负责一个点，然后判断归属，例如归于ragion X, 那么可以直接atomic(®ion, 1)。（而不是你从区域的角度出发)。这样，只要线程数目能覆盖到每个点上即可。无需考虑block是否有它的含义。
至于有无好于atomicAdd的分类统计方法，我建议你先不要考虑，先正确实现了你的算法，然后速度不满意再考虑也不迟。
ok！彻底明白啦！！感谢各位版主！:handshake
客气客气，欢迎您常来论坛转转，讨论问题。
顺便转达横扫千军斑竹的回谢！
祝您新年编码愉快~