CUDA浮点数添加给出错误的答案（与CPU浮点运算相比）

Question

我是CUDA的新手。我在使用cuda查找浮点矢量的点prod时，遇到cuda中的浮点加法问题。本质上，下面是简单的内核。我正在使用-arch = sm_50，所以基本思路是让thread_0添加向量a的值。

__global__ void temp(float *a, float *b, float *c) {

if (0 == threadIdx.x && blockIdx.x == 0 && blockIdx.y ==0 ) {
    float xx = 0.0f;
        for (int i = 0; i < LENGTH; i++){
            xx += a[i];
        }
        *c = xx;
    }
}

当我用1000个元素的1.0初始化'a'时，我得到了1000.00的期望结果

但是当我用1.1初始化'a'时，我应该得到1100.00xx，但是我却得到了1099.989014。cpu的实现仅产生1100.000024

我试图了解这里的问题！:-(

我什至试图计算向量中1.1个元素的数量，并得出1000的期望值。我什至使用了atomicAdd，但仍然遇到相同的问题。

如果有人可以在这里帮助我，将不胜感激！

最好

编辑：这里最大的担心是CPU结果与GPU结果之间的差异！我知道浮点数可以减去一些小数点。但是GPU错误非常重要！:-(

Answer 1

使用 IEEE-754 浮点表示法不可能精确地表示 1.1。正如 @RobertCrovella 在他的评论中提到的，在 CPU 上执行的计算不使用与 GPU 相同的 IEEE-754 设置。

事实上，浮点形式的 1.1 存储为0x3F8CCCCD = ，即 1.10000002384185。对 1000 个元素执行求和，最后一位在路由中丢失，第一次加法一位，四位后两位，依此类推，直到 1000 个元素后的 10 位。根据舍入模式，您可能会截断后半部分的 10 位操作，因此最终求和为 0x3F8CCC00，即 1.09997558。

CUDA除以1000的结果是0x3F8CCC71，这与32位计算一致。

在 CPU 上编译时，根据优化标志，您可能会使用快速数学，它使用内部寄存器精度。如果不指定向量寄存器，则可以使用 80 位精度的 x87 FPU。在这种情况下，计算将以浮点形式读取 1.1，即 1.10000002384185，使用更高的精度将其相加 1000 倍，因此在舍入中不会丢失任何位，结果为 1100.00002384185，并显示 1100.000024，这是其舍入到最接近的显示。

根据编译标志，CPU 上的实际等效计算可能需要强制执行 32 位浮点算术，这可以使用adds来完成例如 SSE2 指令集的

您还可以使用/fp:选项或-mfpmath与编译器一起使用并探索发布的指令。在这种情况下组装指令fadd是 80 位精度加法。

所有这些都与 GPU 浮点精度无关。这是对 IEEE-754 规范和传统 x87 FPU 行为的一些误解。