0 
#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调试器是没有问题的。原因可能是使用调试器运行很慢并且不会导致数据竞争。

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

4

1 回答 1

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;
于 2019-12-16T12:48:01.430 回答