cuda：多个线程访问同一个全局变量答案

【问题标题】：cuda:multiple threads access the same global variablecuda：多个线程访问同一个全局变量
【发布时间】：2019-12-16 06:23:42
【问题描述】：

#define dimG 16
#define dimB 64

// slovebyGPU
__global__ void SloveStepGPU(float* X, float* Y, int * iCons, int* jCons, int * dCons, float* wCons, int cnt, float c)
{
    int id = blockDim.x * blockIdx.x + threadIdx.x;
    for (int i = id; i<cnt; i += dimG*dimB) {

        int I = iCons[i];
        int J = jCons[i];
        int d = dCons[i];
        float wc = 1.0f*wCons[i]*c;

        if (wc > 1.0)wc = 1.0;

        float XI = atomicAdd(&(X[I]), 0);
        float XJ = atomicAdd(&(X[J]), 0);
        float YI = atomicAdd(&(Y[I]), 0);
        float YJ = atomicAdd(&(Y[J]), 0);
        float pqx = XI - XJ;
        float pqy = YI - YJ;
        float mag = sqrtf(pqx*pqx + pqy*pqy);
        float r = 1.0f*(d - mag) / 2;
        float mx = wc * r * pqx / (mag + eps);
        float my = wc * r * pqy / (mag + eps);
        if (d == 1) {
            atomicAdd(&(X[I]), mx);
            atomicAdd(&(Y[I]), my);
        }
        atomicAdd(&(X[J]), -mx);
        atomicAdd(&(Y[J]), -my);
}

在这段代码中，我知道 X、Y 可能存在数据竞争。我之前的想法是：允许读取的XI、XJ、YI、YJ可能不是最新的数据。但是我发现在数据竞争的过程中，可能会导致XI、XJ、YI、YJ读取随机内存值。也就是说，内存访问冲突。即使我在读写过程中加了一个锁，我仍然得到相同的结果。只有当我减小 dimB 和 dimG 的大小以使几乎没有数据竞争时，我才能得到正确的结果。有什么解决办法吗？

我在windows + vs2015 + cuda9.1环境下使用64位编译。

但是，我在linux下使用同样的代码，没有发现任何问题。

windows下使用nsight cuda调试器没有问题。原因可能是使用调试器运行很慢，不会导致数据争用。

-------更新行----- 删除其他代码

【问题讨论】：

能否添加设置iCons、jCons和dCons的代码？
感谢您的回答，我更新了问题。您认为该错误是由未定义的行为或初始化问题引起的，对吗？我查了一下，我认为原因是数据竞争。但是即使我使用原子函数，我仍然会遇到同样的错误。这个错误只出现在 Windows 上，而不会出现在 Linux 上。
我对@987654325@ 的元素很感兴趣，以了解I 和J 的值，因为您将它们用作X 和Y 的索引，而我不知道这些指数可能是什么。如果iCons 的所有元素都为零，那么您有理由参加比赛。如果它们大于 X 的大小，那么您可能有未定义的行为。
&(X[I]) 和X+I 有区别吗？当你用原子方法包装比赛时，为什么要报告比赛也许是有原因的？
iCons,jCons,dCons,wCons 描述了一些约束。iCons ,jCons 是一个索引。它的大小不会超过GN.X，Y 是一个随机生成的 0-1 值。可以肯定的是，XY、iCons等应该是没有问题的，因为它们都在CPU求解功能中。我更新了CPU slove功能。

标签： c++ windows cuda gpu data-race

【解决方案1】：

问题出现在这个if (d == 1)，我把if换成设备函数fminf,fmaxf等解决问题。估计是进入了同一个warp的分支，存在数据竞争，部分进程被挂起，导致奇怪的问题。

if (d == 1) {
            atomicAdd(&(X[I]), mx);
            atomicAdd(&(Y[I]), my);
        }

到

float fd = fmaxf(2.0f - d, 0.0f);
X[I] += fd * 1.0f * mx;
Y[I] += fd * 1.0f * my;

【讨论】：