OpenCL内核计算精度出现不可预测变化的技术求助

我有一段从大型代码中提取的OpenCL内核测试片段，用于复现核心精度问题：

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值并打印，主机端并行执行相同计算作对比：

• 默认情况下，GPU计算的A值仅精确到9位小数，与CPU结果存在差异：

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

• 若取消calcA函数中中间变量（如A_h1）的打印注释，GPU会输出与CPU完全一致的精确结果：

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

环境信息 #

• 使用<CL/opengl.hpp>库
• 硬件：RTX 4080
• 软件：CUDA 12.4，驱动551.78

已尝试的解决方法 #

• 尝试使用--fmad=false编译标志，但该标志不属于OpenCL规范，无效
• 按建议使用-cl-opt-disable标志，仅在本测试示例中有效，但在主程序中仅略微改变不精确值，未达到预期精度

内容的提问来源于stack exchange，提问作者Roman Zuev