1

我正在 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);


        }
    }
}

}

谢谢

4

1 回答 1

2

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 主要用作健全性检查。

Where accuracy issues are suspected, I generally recommend comparison with a higher-precision reference implementation (e.g. one based on quad precision or double-double techniques) to accurately assess the error on both CPU and GPU, rather than using the CPU version as a reference (as the CPU results are also affected by round-off error).

于 2012-07-27T02:21:48.127 回答