Ответ 1
Вы можете получить приблизительную оценку производительности, подсчитав количество операций float, необходимых для реализации каждой операции с двойным поплавком. Вы хотите проверить двоичный код с помощью cuobjdump --dump-sass, чтобы получить точный счет. Я показываю двукратное умножение ниже, что в полной мере использует поддержку FMA (плавное многократное добавление) на графическом процессоре. Для кода с двойным поплавком я бы указал вам на статью Эндрю Талла, так как у меня нет времени, чтобы закодировать это прямо сейчас. Из предыдущего анализа я считаю, что код добавления, приведенный в документе, верен и что он избегает распространенных ошибок в более быстрых, но менее точных реализациях (которые теряют точность, когда величина операндов находится в два раза).
Если вы являетесь зарегистрированным разработчиком CUDA, вы можете скачать двухдверный код с сайта разработчика NVIDIA (войдите в https://developer.nvidia.com), который находится под лицензией BSD, и переделайте его относительно быстро в двойной плавающий код. Двухдиапазонный код NVIDIA поддерживает операции сложения, вычитания, деления, квадратного корня и обратного квадратного корня.
Как вы можете видеть, для умножения ниже требуется 8 float инструкций; унарное отрицание поглощается в FMA. Для добавления требуется около 20 float инструкций. Однако последовательности команд для операций с двойным поплавком также требуют временных переменных, что увеличивает давление в регистре и может уменьшить заполняемость. Поэтому разумно консервативная оценка может заключаться в том, что арифметика с двойным поплавком выполняет на 1/20 пропускную способность нативной float арифметики. Вы можете легко измерить это самостоятельно, в контексте, относящемся к вам, то есть в вашем случае (-ях) использования.
typedef float2 dblfloat; // .y = head, .x = tail
__host__ __device__ __forceinline__
dblfloat mul_dblfloat (dblfloat x, dblfloat y)
{
dblfloat t, z;
float sum;
t.y = x.y * y.y;
t.x = fmaf (x.y, y.y, -t.y);
t.x = fmaf (x.x, y.x, t.x);
t.x = fmaf (x.y, y.x, t.x);
t.x = fmaf (x.x, y.y, t.x);
/* normalize result */
sum = t.y + t.x;
z.x = (t.y - sum) + t.x;
z.y = sum;
return z;
}
Обратите внимание, что в разных приложениях полная двунаправленная арифметика может не понадобиться. Вместо этого можно использовать вычисление float, дополненное методами компенсации ошибок, одним из старейших из которых является суммирование Кахана. Я дал краткий обзор легкодоступной литературы по таким методам в недавнем размещении на форумах разработчиков NVIDIA. В комментариях выше Роберт Кровелла также указал на разговор GTC 2015 Скотта Легранда, который я еще не успел проверить.
Что касается точности, double-float имеет репрезентативную точность 49 (24 + 24 + 1) бит по сравнению с IEEE-755 double, которая обеспечивает 53 бит. Однако double-float не может поддерживать эту точность для операндов малой по величине, так как хвостовая часть может стать денормальной или нулевой. Когда поддержка denormal включена, 49 бит точности гарантируется для 2 -101 <= | x | < 2 128. Денормальная поддержка float включена по умолчанию в цепочке инструментов CUDA для архитектуp >= sm_20, что означает все архитектуры, поддерживаемые текущей продаваемой версией CUDA 7.0.
В отличие от операций с данными IEEE-754 double операции с двойным поплавком не корректно округлены. Для умножения с двойным поплавком выше, используя 2 миллиарда случайных тестовых случаев (со всеми исходными операндами и результатами в границах, указанных выше), я наблюдал верхнюю границу 1.42е-14 для относительной ошибки. У меня нет данных для добавления с двойным поплавком, но его ошибка должна быть схожей.