CUDA浮点加法给出了错误的答案(与CPU浮动操作相比)

2023年7月15日 256次阅读

我是CUDA的新手.我正在使用cuda找到浮动向量的点刺,我在cuda中遇到了浮点加法问题.本质上,以下是简单的内核.我正在使用-arch = sm_50

所以基本思路是让thread_0添加vector 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

我想知道这里的问题是什么！ 🙁

我甚至试图计算a矢量中的1.1个元素的数量,并且预计会产生1000个元素.我甚至使用atomicAdd,但我仍有同样的问题.

如果有人能帮帮我,我将非常感激！

最好

编辑：
这里最大的担忧是CPU结果与GPU结果的差异！我知道浮点数可能会被一些小数点所取消.但GPU错误非常重要！ 🙁

最佳答案 不可能完全使用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.在那次出现时,计算将在float中读取1.1,其为1.10000002384185,使用更高的精度将其添加1000次,因此在舍入中不会丢失任何位,导致1100.00002384185,并且显示1100.000024,这是其到最近显示的舍入.

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

您还可以使用编译器和/fp:选项或-mfpmath一起使用,并浏览已发布的指令.在那种情况下,汇编指令fadd是80位精度加法.

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

点赞