__float128加法在CPU与CUDA上结果不一致的原因

我用以下CUDA程序测试__float128类型的加法：

// nvcc -arch=sm_100 test_fp128_add.cu -o test_fp128_add

#include <cstdio>
#include <cuda_runtime.h>

__global__ void add_fp128(__float128 a,
                          __float128 b,
                          __float128* out)
{
  *out = a + b;
}

int main()
{
    __float128 a = 16.0Q;   // Q for quad precision
    __float128 b = 32.0Q;

    __float128 *d_out;
    __float128 h_out;

    cudaMalloc(&d_out, sizeof(__float128));

    add_fp128<<<1,1>>>(a, b, d_out);

    cudaMemcpy(&h_out, d_out, sizeof(__float128), cudaMemcpyDeviceToHost);

    printf("GPU: %.2f + %.2f = %.2f\n", (double)a, (double)b, (double)h_out);

    printf("CPU: %.2f + %.2f = %.2f\n", (double)a, (double)b, (double)(a + b));

    return 0;
}

运行后控制台输出：

GPU: 16.00 + 32.00 = 0.00
CPU: 16.00 + 32.00 = 48.00

问题原因 #

• 老旧架构无__float128硬件支持：你编译时指定的-arch=sm_100属于早期Tesla架构（如G80系列），这类GPU完全没有128位浮点运算的硬件单元，编译器无法生成有效的设备端__float128运算代码，导致加法操作未正确执行，结果无效。
• 核函数参数传递限制：早期CUDA架构的核函数参数传递不支持__float128类型，主机端的参数无法正确传递到设备端，进一步导致运算数值错误。
• __float128的最低支持要求：从Volta架构（sm_70）开始，CUDA才提供部分__float128硬件支持，且需要配合CUDA 9及以上版本的工具包，仅部分运算（如乘加）能得到硬件加速。

修复建议 #

• 编译时指定支持__float128的GPU架构，例如：nvcc -arch=sm_70 test_fp128_add.cu -o test_fp128_add
• 升级CUDA Toolkit到9.0及以上版本
• 注意：CUDA没有原生支持__float128的格式化输出，当前用(double)强制转换会丢失精度，若需完整打印__float128数值，需自行实现字符串转换函数。

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