OpenCL 内核中计算精度发生不可预测的变化

Question

大家下午好！

我有以下

OpenCL kernel code

：

double calcA(double2 point1, double2 point2);

void __kernel mainProgram()
{
    long long index = get_global_id(0);

    if (index == 0)
    {
        double2 p1   = (double2)(0.891949793450334649, 0.513485940430910115);
        double2 p2   = (double2)(0.891949846176460226, 0.513485940430959964);

        double dist = calcA(p1, p2);

        printf("NV A        = %.18e\n", dist);
    }
}


double calcA(double2 p1, double2 p2)
{
    double cos1 = cos(p1[0]);
    double cos2 = cos(p2[0]);
    double sin1 = sin(p1[0]);
    double sin2 = sin(p2[0]);
    double cosDelta = cos(p2[1] - p1[1]);

    double A_h1 = cos1 * sin2;
    double A_h2 = sin1 * cos2 * cosDelta;
    double A = A_h1 - A_h2;

    // printf("NV A_h1     = %.18e\n", A_h1);

    return A;
}

这是更大代码的一部分，仅演示了根本问题。在此代码中，如您所见，在中央内核函数内部，调用了一个本地函数来计算

A

的某个值（分三步）并将其打印在屏幕上。 在主机程序中，并行执行相同的计算并打印在屏幕上以比较获得的值。 问题在于，在

A

上计算出的

GPU

的值，精度高达

9 decimal places

:

CPU A = 5.27261255 7671862987e-08
NV  A = 5.27261255 9188061227e-08

但是，如果在 calcD 函数中打印中间变量之一（

A_h1

或

A_h2

）（例如，取消注释一行），则 GPU 开始输出

A

的绝对正确值：

CPU A = 5.272612557671862987e-08
NV  A = 5.272612557671862987e-08

我注意到我直接使用

<CL/opengl.hpp>

库，对

RTX 4080

（

CUDA 12.4

，驱动程序

551.78

）执行计算。 我还注意到，我看到过像使用编译标志

--fmad=false

这样的提示，但是 OpenCL 规范中没有这样的标志。

非常感谢您的帮助！

Answer 1

打印中间值时，GPU 在您编码时执行数学计算，并获得准确的最终结果。当不打印中间值时，这里的乘法和减法由编译器收缩为一个 fma/mad ，这可能会使用降低的精度：

double A = fma(cos1, sin2, -sin1 * cos2 * cosDelta); 

您可以使用传递给

-cl-opt-disable

的 clBuildProgram

编译器标志

禁用此行为。

请注意，Nvidia Ada GPU 上的双精度 (FP64) 性能非常糟糕，甚至很可能低于 CPU 的 FP64 性能。我建议仅当您可以摆脱单精度（FP32）时才使用 GPU。