CUDA内核结果在释放模式下不同

Question

我在CUDA中测试了一些代码（我是CUDA的新手，这是我的第一个申请）。到目前为止，我已经在CUDA中取得了成果，这与通过在CPU上串行运行代码而相同的结果。我正在使用Visual Studio 2010，构建配置是调试。但是一旦我将构建配置更改为“释放”，我开始获得错误的结果。我无法在此刻使用NVIDIA论坛。拥有CUDA体验的人能指出这个问题。代码如下

__global__ void MyKernel(int *Nptr,int *deltaptr, double *gravityptr, double *separationptr, double *fconptr, double *xForce, double *yForce, double *zForce,
double *xPos, double *yPos, double *zPos )
{
int N = *Nptr;
int delta= *deltaptr;
double gravity= *gravityptr;
double separation = *separationptr;
double fcon = *fconptr;

double len=0.0;
double r12X =0.0;
double r12Y =0.0;
double r12Z =0.0;
double PE=0.0;


int nx = blockDim.x * blockIdx.x + threadIdx.x;//use this place of nx
//int ny = blockDim.x * blockIdx.x + threadIdx.y;//use this place of ny
int ny = blockDim.y * blockIdx.y + threadIdx.y;
//printf("nx:%d ny:%d\n", nx,ny);

if(!(nx< N && ny <N))
    return;
//printf("nx:%d ny:%d\n", nx,ny);


xForce[nx*N+ny] = 0.0;
yForce[nx*N+ny] = -gravity;
zForce[nx*N+ny] = 0.0;

int lowerValuedx = maxOnDevice(nx-delta,0);
int upperValuedx=minOnDevice(nx+delta+1,N);
for(int dx=lowerValuedx; dx<upperValuedx;dx++)
{
    int lowerValuedy=maxOnDevice(ny-delta,0);
    int upperValuedy=minOnDevice(ny+delta+1,N);
    for(int dy=lowerValuedy; dy<upperValuedy;dy++)
    {
        len=sqrt((double)((nx-dx)*(nx-dx)+(ny-dy)*(ny-dy)) ) *separation;
        bool condition = ny!=dy;
        bool condition1 = nx!=dx;

        //if (nx!=dx || ny!=dy)
        if (condition || condition1)
        {
            r12X = xPos[dx*N+dy] - xPos[nx*N+ny];
            r12Y = yPos[dx*N+dy] - yPos[nx*N+ny];
            r12Z = zPos[dx*N+dy] - zPos[nx*N+ny];
            xForce[nx*N+ny] = xForce[nx*N+ny] +fcon*normxOnDevice(r12X,r12Y,r12Z)*(magOnDevice(r12X,r12Y,r12Z)-len);
            yForce[nx*N+ny]= yForce[nx*N+ny] +fcon*normyOnDevice(r12X,r12Y,r12Z)*(magOnDevice(r12X,r12Y,r12Z)-len);
            zForce[nx*N+ny]= zForce[nx*N+ny] +fcon*normzOnDevice(r12X,r12Y,r12Z)*(magOnDevice(r12X,r12Y,r12Z)-len);


        }
    }
}   

.

}

感谢

Answer 1

对于CPU和GPU结果之间的数值差异，以及在GPU上的调试和发布构建之间存在数值差异并不罕见。这并不意味着一组结果是不正确的，但是一个可能比另一个更准确。请参阅NVIDIA的以下白皮书，讨论可能导致数值差异的各种机制：

http://developer.dowload。 nvidia.com/assets/cuda/files/nvidia-cuda-floation-point.pdf

您可以检查nvcc flag -fmad= false是否消除了所看到的差异，这表明这些差异是由于FMA / FMAD合并，因此可能是无害的。

GPU提供FMAD和FMA（熔融乘法添加）操作，该操作将浮点数与从属浮点添加到单个操作中。这有助于性能，因为组合操作通常需要与每个组成部分中的每一个相似的时间。但是，组合操作的舍入行为与使用两个单独舍入的操作不同：

单精度FMAD（计算能力<2.0）截断乘法的结果，然后根据IEEE-754圆形到最接近的或偶数舍入的最终添加结果。通过对比度FMA（Compute能力的单精度>= 2.0，双精度）计算Untryed，双宽的产品，将第三个操作数添加到该操作数，并根据IEEE-754圆形到最接近或最终的最终和-甚至。由于这种单曲，FMA提供的平均精度优于使用两个单独舍入的操作。 FMA操作在2008年版本的IEEE-754浮点标准中指定。

默认情况下，对于发布构建，CUDA编译器积极生成合并操作（FMAD，FMA）以实现最佳性能。换句话说，编译器默认为-fmad= true，允许编译器合并浮点乘法并添加。通过指定-fmad= false，禁止乘法和添加的合并，这通常可以通过CPU结果提供更大的一致性，因为大多数CPU不提供FMA操作。显然，禁用合并操作的使用对性能产生负面影响，因此-fmad= false主要有用作为理智检查。

在怀疑准确性问题的情况下，我通常建议比较与更高精度的参考实现（例如，基于四边形或双级技术）进行比较，以准确地评估CPU和GPU上的错误，而不是使用CPU版本作为参考（由于CPU结果也受圆截止误差的影响）。