CUDA: Различия в точности при суммировании удваивается на хосте и устройстве? - PullRequest
Новая
       27

CUDA: Различия в точности при суммировании удваивается на хосте и устройстве?

0 голосов
/ 20 апреля 2020

Только что начал экспериментировать с CUDA, и я наткнулся на контрольно-пропускной пункт: (

Я пытаюсь суммировать несколько двойных переменных с высокой точностью, но результат потрясающе отличается (после десятичной запятой), когда выполняя добавление на устройстве против хоста.Я проверил следующий пример кода с помощью nv cc: 

#include <stdio.h>

using namespace std;

/// Kernel
__global__ void SolveRE(double x, double y, double z)
{
    printf("From GPU: %.15lf\n", (x*x + y*y + z*z));
}

int main ()
{
    // Define values
    double x = 3051774.263732617720962;
    double y = 7168331.130877199582756;
    double z = 983.741348489015081;

    // Call kernel
    SolveRE<<<1, 1>>>(x, y, z);
    cudaDeviceSynchronize();

    // Print from host
    printf("From CPU: %.15lf\n", (x*x + y*y + z*z));

    return 0;
}

Когда я запускаю вышеуказанный код, я получаю следующий вывод: 

From GPU: 60698298326430.984375000000000
From CPU: 60698298326430.992187500000000

Я предполагаю, что это расхождение как-то связано с точностью операций с плавающей запятой? Или я что-то упускаю? Любой способ привести результаты в соответствие друг с другом?

Буду признателен за любой совет Спасибо всем!

Ответы [ 2 ]

2 голосов
/ 21 апреля 2020

Как уже говорилось в комментариях, причиной root является генерация команды FMAD на GPU. У меня нет доступа к системе, чтобы проверить, какой результат будет сгенерирован, если инструкции FMA3 / FMA4 использовались для использования инструкций FMADD на стороне процессора. В таких условиях результат может быть другим.

Однако вам нужно быть предельно осторожным при рассмотрении подобных примеров, потому что то, что вы видите, не обязательно связано с различиями между ЦП и ГП .

Чтобы показать, что я имею в виду, посмотрите на упрощенную (и более сложную) версию вашего кода хоста: 

#include <stdio.h>

using namespace std;

double func(double x, double y, double z, bool docalc=true)
{
    if (docalc) {
        return x*x + y*y + z*z;
    } else {
        double xp = 3051774.263732617720962;
        double yp = 7168331.130877199582756;
        double zp =     983.741348489015081;
        return xp*xp + yp*yp + zp*zp;
    }
}

int main ()
{
    // Define values
    double x = 3051774.263732617720962;
    double y = 7168331.130877199582756;
    double z =     983.741348489015081;

    // Print from host
    double val = func(x, y, z);
    double valp = func(x, y, z, false);
    printf("From CPU: %.15lf\n", val);
    printf("From CPU: %.15lf\n", valp);

    return 0;
}

скомпилировано, выполнено и дизассемблировано следующим образом: 

$ g++ -std=c++11 -O2 -o doubleh doubleh.cc
$./doubleh
From CPU: 60698298326430.992187500000000
From CPU: 60698298326430.992187500000000
$ objdump -d doubleh

....
0000000000000560 <main>:
 560:   48 83 ec 08             sub    $0x8,%rsp
 564:   48 8d 35 09 02 00 00    lea    0x209(%rip),%rsi        # 774 <_IO_stdin_used+0x4>
 56b:   bf 01 00 00 00          mov    $0x1,%edi
 570:   c5 fb 10 05 10 02 00    vmovsd 0x210(%rip),%xmm0        # 788 <_IO_stdin_used+0x18>
 577:   00 
 578:   b8 01 00 00 00          mov    $0x1,%eax
 57d:   e8 be ff ff ff          callq  540 <__printf_chk@plt>
 582:   c5 fb 10 05 fe 01 00    vmovsd 0x1fe(%rip),%xmm0        # 788 <_IO_stdin_used+0x18>
 589:   00 
 58a:   48 8d 35 e3 01 00 00    lea    0x1e3(%rip),%rsi        # 774 <_IO_stdin_used+0x4>
 591:   bf 01 00 00 00          mov    $0x1,%edi
 596:   b8 01 00 00 00          mov    $0x1,%eax
 59b:   e8 a0 ff ff ff          callq  540 <__printf_chk@plt>
 5a0:   31 c0                   xor    %eax,%eax
 5a2:   48 83 c4 08             add    $0x8,%rsp
 5a6:   c3                      retq   
 5a7:   66 0f 1f 84 00 00 00    nopw   0x0(%rax,%rax,1)
 5ae:   00 00

Результат - правильный результат, что неудивительно. Но обратите внимание, что в сборке нет вызова функции и арифметики с плавающей запятой c. Компилятор встроил и оптимизировал оба вызова функций и заменил их двумя константами времени компиляции. Таким образом, результат, который вы видите, хотя и является правильным, не был вычислен во время выполнения и ничего не говорит о том, что на самом деле будет делать ЦП, если ваш код был запущен.

Итак, давайте go для GPU: 

#include <stdio.h>

using namespace std;
__device__ __host__ double func(double x, double y, double z, bool docalc=true)
{
    if (docalc) {
        return x*x + y*y + z*z;
    } else {
        double xp = 3051774.263732617720962;
        double yp = 7168331.130877199582756;
        double zp =     983.741348489015081;
        return xp*xp + yp*yp + zp*zp;
    }
}

/// Kernel
__global__ void SolveRE(double x, double y, double z, bool docalc=true)
{
    double val = func(x,y,z, docalc);
    printf("From GPU: %.15lf\n", val);
}

int main ()
{
    // Define values
    double x = 3051774.263732617720962;
    double y = 7168331.130877199582756;
    double z =     983.741348489015081;

    // Call kernel
    SolveRE<<<1, 1>>>(x, y, z);
    SolveRE<<<1, 1>>>(x, y, z, false);
    cudaDeviceSynchronize();

    // Print from host
    double val = func(x, y, z);
    double valp = func(x, y, z, false);
    printf("From CPU: %.15lf\n", val);
    printf("From CPU: %.15lf\n", valp);

    return 0;
}

Снова скомпилировано, запущено и разобрано: 

$ nvcc -std=c++11 -arch=sm_52 --fmad=true -o double double.cu
$ ./double
From GPU: 60698298326430.984375000000000
From GPU: 60698298326430.992187500000000
From CPU: 60698298326430.992187500000000
From CPU: 60698298326430.992187500000000

$ cuobjdump -ptx double

.version 6.4
.target sm_52
.address_size 64


.visible .entry _Z7SolveREdddb(
.param .f64 _Z7SolveREdddb_param_0,
.param .f64 _Z7SolveREdddb_param_1,
.param .f64 _Z7SolveREdddb_param_2,
.param .u8 _Z7SolveREdddb_param_3
)
{
.local .align 8 .b8 __local_depot0[8];
.reg .b64 %SP;
.reg .b64 %SPL;
.reg .pred %p<2>;
.reg .b16 %rs<3>;
.reg .b32 %r<2>;
.reg .f64 %fd<8>;
.reg .b64 %rd<5>;


mov.u64 %SPL, __local_depot0;
cvta.local.u64 %SP, %SPL;
ld.param.f64 %fd1, [_Z7SolveREdddb_param_0];
ld.param.f64 %fd2, [_Z7SolveREdddb_param_1];
ld.param.f64 %fd3, [_Z7SolveREdddb_param_2];
add.u64 %rd1, %SP, 0;
add.u64 %rd2, %SPL, 0;
ld.param.s8 %rs1, [_Z7SolveREdddb_param_3];
and.b16 %rs2, %rs1, 255;
setp.eq.s16 %p1, %rs2, 0;
mul.f64 %fd4, %fd2, %fd2;
fma.rn.f64 %fd5, %fd1, %fd1, %fd4;
fma.rn.f64 %fd6, %fd3, %fd3, %fd5;
selp.f64    %fd7, 0d42CB9A36414ECF7F, %fd6, %p1;
st.local.f64 [%rd2], %fd7;
mov.u64 %rd3, $str;
cvta.global.u64 %rd4, %rd3;

    {
.reg .b32 temp_param_reg;

    .param .b64 param0;
st.param.b64    [param0+0], %rd4;
.param .b64 param1;
st.param.b64    [param1+0], %rd1;
.param .b32 retval0;
call.uni (retval0), 
vprintf, 
(
param0, 
param1
);
ld.param.b32    %r1, [retval0+0];


    }
    ret;
}

Вы видите, что графический процессор выдает два разных результата! Компилятор графического процессора сделал именно то, что сделал хост-компилятор - оптимизировал вычисление с плавающей запятой и заменил его константой, если он определил, что это можно сделать безопасно.

Отключение инструкций FMAD устраняет разницу и изменяет код устройства сгенерированный компилятором: 

$ nvcc -std=c++11 -arch=sm_52 --fmad=false -o double double.cu
$ cuobjdump -ptx double


.version 6.4
.target sm_52
.address_size 64



.visible .entry _Z7SolveREdddb(
.param .f64 _Z7SolveREdddb_param_0,
.param .f64 _Z7SolveREdddb_param_1,
.param .f64 _Z7SolveREdddb_param_2,
.param .u8 _Z7SolveREdddb_param_3
)
{
.local .align 8 .b8 __local_depot0[8];
.reg .b64 %SP;
.reg .b64 %SPL;
.reg .pred %p<2>;
.reg .b16 %rs<3>;
.reg .b32 %r<2>;
.reg .f64 %fd<10>;
.reg .b64 %rd<5>;


mov.u64 %SPL, __local_depot0;
cvta.local.u64 %SP, %SPL;
ld.param.f64 %fd1, [_Z7SolveREdddb_param_0];
ld.param.f64 %fd2, [_Z7SolveREdddb_param_1];
ld.param.f64 %fd3, [_Z7SolveREdddb_param_2];
add.u64 %rd1, %SP, 0;
add.u64 %rd2, %SPL, 0;
ld.param.s8 %rs1, [_Z7SolveREdddb_param_3];
and.b16 %rs2, %rs1, 255;
setp.eq.s16 %p1, %rs2, 0;
mul.rn.f64 %fd4, %fd1, %fd1;
mul.rn.f64 %fd5, %fd2, %fd2;
add.rn.f64 %fd6, %fd4, %fd5;
mul.rn.f64 %fd7, %fd3, %fd3;
add.rn.f64 %fd8, %fd6, %fd7;
selp.f64    %fd9, 0d42CB9A36414ECF7F, %fd8, %p1;
st.local.f64 [%rd2], %fd9;
mov.u64 %rd3, $str;
cvta.global.u64 %rd4, %rd3;

    {
.reg .b32 temp_param_reg;

    .param .b64 param0;
st.param.b64    [param0+0], %rd4;
.param .b64 param1;
st.param.b64    [param1+0], %rd1;
.param .b32 retval0;
call.uni (retval0), 
vprintf, 
(
param0, 
param1
);
ld.param.b32    %r1, [retval0+0];


    }
    ret;
}


$ ./double
From GPU: 60698298326430.992187500000000
From GPU: 60698298326430.992187500000000
From CPU: 60698298326430.992187500000000
From CPU: 60698298326430.992187500000000

Ничто из этого не меняет того факта, что источником разницы в вашем примере является генерация команды GPU FMAD, но cavaet emptor , компиляторы намного умнее чем вы думаете, и они делают вещи, которые могут легко опровергнуть сравнения, такие как в вашем вопросе.

1 голос
/ 22 апреля 2020

Расхождение между результатом процессора и графического процессора здесь возникает не из-за слияния-умножения, а из-за порядка операций. Как указано в другом ответе, первый элемент, который нужно «исправить» в сравнении, - это позволить фактическому вычислению результата ЦП с помощью кода времени выполнения, а не предварительно вычисляемого компилятором. Если затем мы просто переупорядочим операнды x, y и z, мы можем изменить результат на обратный, так что результат ЦП теперь тот, который находится дальше от правильно округленного результата, чем результат ГП: 

$ cat t1699.cu
#include <stdio.h>

using namespace std;

/// Kernel
__global__ void SolveRE(double x, double y, double z)
{
    printf("From GPU: %.15lf\n", (x*x + y*y + z*z));
}


double SolveRE_CPU(double x, double y, double z){

  return (x*x + y*y + z*z);
}

int main ()
{
    // Define values
    double z = 3051774.263732617720962;
    double x = 7168331.130877199582756;
    double y = 983.741348489015081;

    // Call kernel
    SolveRE<<<1, 1>>>(x, y, z);
    cudaDeviceSynchronize();

    // Print from host
    printf("From CPU: %.15lf\n", SolveRE_CPU(x, y, z));

    return 0;
}
$ nvcc -o t1699 t1699.cu
$ ./t1699
From GPU: 60698298326430.992187500000000
From CPU: 60698298326430.984375000000000
$

Теперь, в этой ситуации, компиляция с -fmad=false снова выстроит в соответствие эти два результата (результат GPU станет таким же, как результат CPU, указанный выше, что является «менее точным» результатом), но это не из-за какого-либо эффекта самой операции слияния-умножения, а потому, что компилятор (GPU) выбирает другой порядок оценки промежуточных результатов.

...