CUDA内核在发布模式下会产生不同的结果

Question

我正在CUDA中测试一些代码（我是CUDA的新手，这是我的第一个应用程序）。到目前为止，我已经在CUDA中获得了与在CPU上串行运行代码相同的结果。我使用的是Visual Studio 2010，生成配置为"调试"。但是，一旦我将构建配置更改为"发布"，我就会开始得到错误的结果。我还没能使用英伟达论坛，因为它们现在已经关闭了。有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:%dn", nx,ny);
if(!(nx< N && ny <N))
    return;
//printf("nx:%d ny:%dn", 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.download.nvidia.com/assets/cuda/files/NVIDIA-CUDA-Floating-Point.pdf

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

GPU提供FMAD和FMA（融合乘法-加法）运算，将浮点乘法与相关浮点加法组合为一个运算。这有助于提高性能，因为组合操作通常需要与其每个组成部分相似的时间。然而，任何一种组合运算的取整行为都不同于使用两种单独取整运算：

单精度FMAD（计算能力<2.0）截断乘法的结果，然后根据IEEE-754舍入将最终加法的结果舍入为最接近或偶数。相比之下，FMA（计算能力上的单精度>=2.0，双精度）计算未舍入的双宽乘积，将第三个操作数添加到该乘积上，并根据IEEE-754四舍五入将最终和舍入为最接近或偶数。由于这种单一舍入，FMA提供的平均精度优于使用两个单独舍入运算。FMA操作在2008版本的IEEE-754浮点标准中指定。

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

如果怀疑精度问题，我通常建议将其与更高精度的参考实现进行比较（例如，基于四精度或双精度技术的参考实现），以准确评估CPU和GPU上的误差，而不是使用CPU版本作为参考（因为CPU结果也会受到舍入误差的影响）。