Что и почему

FP16 - это формат IEEE, в котором количество битов уменьшено по сравнению с традиционным форматом с плавающей запятой (например, 32 бита = ключевое слово «float», которое мы используем в C / C ++). Основная причина использования этого FP16 с пониженной точностью заключается в том, что при использовании FP16 доступно аппаратное ускорение (если вас устраивает потеря точности) и есть двукратная экономия места.

В частности, некоторые графические процессоры предлагают ускорение от 2X до 8X на FP16 по сравнению с FP32. Несмотря на это, мы часто придерживаемся FP32 (как новички), потому что начало работы с FP16 может быть немного сложным, в основном из-за:

  1. В языке C / C ++ нет встроенной поддержки FP16 - очевидно, версия C / C ++ для ARM имеет довольно хорошую поддержку, но на x86 нам нужно использовать либо специальные инструкции, либо библиотеки для преобразования из FP32- ›FP16 и обратно (что мы освещаем в этой статье)
  2. Как новичок - легко испугаться всех этих преобразований типов, особенно при написании кода хоста / процессора и передаче данных / указателей на сторону устройства / графического процессора.

Как

Поэтому в этой статье я постараюсь дать краткое введение о том, как писать код, использующий FP16. Первая проблема при написании ядер FP16 связана с написанием кода хоста, и для этого у нас есть 2 варианта создания массивов FP16 на ЦП.

Вариант №1:
uint16_t: Если подумать, FP16 на самом деле может быть сохранен на стороне хоста / процессора как unsigned short int (который также использует 16 бит), но нам нужно какой-то механизм для преобразования из float в этот 16-битный формат. На машинах X86 это может быть достигнуто с помощью встроенных функций преобразования Intel FP16, доступных как часть emmintrin.h - если вас это особенно интересует, вы можете прочитать больше здесь, но обратите внимание что для этого требуется icc (или компилятор Intel C ++ для использования библиотеки, которую nvcc может использовать в качестве основного компилятора - посмотрите здесь и здесь).

Вариант №2:
__half: это тип данных, который доступен как часть библиотеки NVIDIA FP16 “cuda_fp16.h”. На мой взгляд, это был самый простой способ заставить FP16 работать на стороне процессора, поскольку он легко предоставляет функции для преобразования в формат FP16 и из него в другие. Подробнее об этой опции я расскажу здесь, но для получения более подробной информации о самой библиотеке и предоставляемых ею функциях загляните здесь.

Например - взгляните на базовый пример кода ниже, он показывает, насколько легко использовать тип данных __half без проблем как для CPU, так и для GPU.

#include <iostream>
#include <stdint.h> 
#include <cuda_fp16.h>
using namespace std;
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
    if (code != cudaSuccess) 
    {   
        fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
        if (abort) exit(code);
    }   
}
__global__ void half_plus1(__half *in_array)
{
    const int idx = threadIdx.x + blockDim.x*blockIdx.x;
    in_array[idx] = __float2half(__half2float(in_array[idx]) + 1.0);
}
int main(void)
{
    const int n = 64;
    __half *h_in, *d_in;
    h_in = (__half*) malloc(n*sizeof(__half));
    gpuErrchk( cudaMalloc(&d_in, n*sizeof(__half)) );
    for (int i=0; i<n; i++)
        h_in[i] = __float2half(1.5);
    gpuErrchk( cudaMemcpy(d_in, h_in, n*sizeof(__half), cudaMemcpyHostToDevice) );
    dim3 block_dims(2,1,1);
    dim3 thread_dims(32,1,1);
    half_plus1<<<block_dims, thread_dims>>>(d_in);
    gpuErrchk( cudaPeekAtLastError() );
    gpuErrchk( cudaDeviceSynchronize() );
    gpuErrchk( cudaMemcpy(h_in, d_in, n*sizeof(__half), cudaMemcpyDeviceToHost) );
    for (int i=0; i<n; i++)
    {   
        if(__half2float(h_in[i]) != 2.5)
        {
            cout<< "Mismatch at " << i << " Expected = 2.5 " << "Actual = " << __half2float(h_in[i]) << endl;
            exit(1);
        }
    }   
    cout << "TEST PASSES" << endl;
    cudaFree(d_in);
    free(h_in);
    return 0;
}

Представление

Тип данных Half Float или FP16 дает нам два основных преимущества, а именно:
1. Экономия места
2. Ускорение арифметических операций (поскольку FP16 работает в 2 раза быстрее, чем FP32)

Но один из важных моментов, который следует отметить, заключается в том, что в архитектуре графического процессора NVIDIA (и, скорее всего, всех других графических процессорах) для достижения хорошей пропускной способности памяти критически важно использовать тип данных __half2, а не __half на стороне графического процессора. Это потому что :

  1. С __half и 32 потоками / warp - мы достигаем только ~ 64B / Load, тогда как для достижения хорошей пропускной способности нам нужно стремиться достичь как минимум 128B / Load (это выгодно как с точки зрения L2, так и с точки зрения шаблона доступа DRAM). В идеале мы хотели бы использовать еще большие нагрузки векторов, чего можно добиться, поместив несколько __half или __half2 в структуру, а затем преобразовав их как uint4 или float4s.
  2. Кроме того, чтобы получить вдвое большую математическую производительность, нам нужно упаковать 2 значения __half в регистры вместе, чтобы использовать 2X арифметические инструкции FP16 (подробнее здесь здесь и здесь). Упаковка необходима, поскольку это формат, в котором FPU ожидает присутствия данных.

Итак, по причинам, указанным выше, код меняется на:

__global__ void half2_plus1(__half2 *in_array)
{
    const int idx = threadIdx.x + blockDim.x*blockIdx.x;
    in_array[idx] = __hadd2(in_array[idx],  __float2half2_rn(1.0));
}
int main(void)
{
    const int n = 64;
    __half2 *h_in, *d_in;
    h_in = (__half2*) malloc(n*sizeof(__half2));
    gpuErrchk( cudaMalloc(&d_in, n*sizeof(__half2)) );
    for (int i=0; i<n; i++)
        h_in[i] = __float2half2_rn(1.5);
    gpuErrchk( cudaMemcpy(d_in, h_in, n*sizeof(__half2), cudaMemcpyHostToDevice) );
    dim3 block_dims(2,1,1);
    dim3 thread_dims(32,1,1);
    half2_plus1<<<block_dims, thread_dims>>>(d_in);
    gpuErrchk( cudaPeekAtLastError() );
    gpuErrchk( cudaDeviceSynchronize() );
    gpuErrchk( cudaMemcpy(h_in, d_in, n*sizeof(__half2), cudaMemcpyDeviceToHost) );
    for (int i=0; i<n; i++)
    {   
        if((__high2float(h_in[i]) != 2.5) || (__low2float(h_in[i]) !=2.5))
        {
            cout<< "Mismatch at " << i << " Expected = 2.5 " << "Actual = " << __half2float(h_in[i].x) << " " << __half2float(h_in[i].y) << endl;
            exit(1);
        }
    }   
    cout << "TEST PASSES" << endl;
    cudaFree(d_in);
    free(h_in);
    return 0;
}

Также обратите внимание, что некоторые функции, такие как __hadd2, поддерживаются только начиная с определенных версий архитектуры GPU, поэтому не забудьте добавить флаг «-arch» во время компиляции nvcc.