并发写入相同的全局内存位置
我有几个块，每个块都有一个大小为512的共享内存数组中的整数。如何检查每个块中的数组是否包含零作为元素？
我正在做的是创建一个驻留在全局内存中的数组。 此数组的大小取决于块的数量，并初始化为0.因此，如果共享内存数组包含零，则每个块写入a[blockid] = 1 。
我的问题是当我在一个块中同时写入多个线程时。 也就是说，如果共享内存中的数组包含多个零，那么多个线程将写入a[blockid] = 1 。 这会产生任何问题吗？
换句话说， 如果2个线程将完全相同的值写入全局内存中完全相同的数组元素 ，那会不会有问题？
在CUDA执行模型中，无法保证从同一块中的线程到同一全局内存位置的每个同时写入都将成功。 至少有一个写操作可以工作，但编程模型不能保证将发生多少写事务，或者如果执行多个事务，它们将以何种顺序发生。
如果这是一个问题，那么更好的方法（从正确的角度来看），就是每个块只有一个线程进行全局写操作。 您可以使用primefaces设置的共享内存标志或还原操作来确定是否应设置该值。 您选择哪个可能取决于可能存在多少个零。 零越多，减少的吸引力就越大。 CUDA包含warp level __any()和__all()运算符，可以在几行代码中构建一个非常有效的布尔减少。
对于CUDA程序，如果warp中的多个线程写入同一位置，则该位置将被更新，但未定义位置更新的次数（即，多少次实际写入串行）并且未定义 哪个线程将写最后一次（即哪个线程将赢得比赛）。
对于计算能力为2.x的设备，如果warp中的多个线程写入同一地址，则只有一个线程实际执行写操作， 该线程未定义。
从CUDA C编程指南第F.4.2节：
如果由warp执行的非primefaces指令写入warp的多个线程的全局内存中的相同位置，则只有一个线程执行写操作，而哪个线程执行它是未定义的。
有关详细信息，另请参阅指南的第4.1节。
换句话说，如果写入给定位置的所有线程都写入相同的值，那么它是安全的。
是的，这将是一个称为Race Condition的问题。
您应该考虑通过process Semaphores synchronizing对全局数据的访问
虽然不是互斥锁或信号量，但CUDA确实包含可用于序列化对给定代码段或存储器位置的访问的同步原型。 通过__syncthreads()函数，您可以创建一个屏障，以便任何给定的线程在命令调用点处阻塞，直到给定块中的所有线程都执行了__syncthreads()命令。 这样，您可以希望序列化对内存位置的访问，并避免两个线程需要同时写入同一内存位置的情况。 唯一的警告是所有线程必须在某个时刻执行__syncthreads() ，否则最终会出现死锁情况。 因此，不要将调用置于某些条件if语句中，其中某些线程可能永远不会执行该命令。 如果您确实像这样处理问题，则需要为最初不调用__syncthreads()的线程做出一些规定，以便稍后调用该函数以避免死锁。