OpenCl: Как я могу использовать INT4 / 8?

Question

Я ищу способы повышения эффективности алгоритма на основе OpenCl.

В настоящее время я использую типы данных float и int на карте Radeon VII. Однако, тип данных, охватывающий числа от -8 до +7, будет достаточным.

Согласно следующей статье Radeon VII достигает максимальной производительности 53/110 TFlops при ограничении на INT8 / INT4, что намного выше, чем при float, что составляет 14 TFlops.

https://www.pcgameshardware.de/Radeon-VII-Grafikkarte-268194/Tests/Benchmark-Review-1274185/2/

Итак, мой вопрос, как я могу использовать операции INT8 / 4? Просто используйте тип данных char вместо int в OpenCl? Поскольку char является самым маленьким типом данных buildin, как я могу даже использовать INT4?

Answer 1

Для "int8", то есть 8-разрядных целых чисел, тип OpenCL действительно char (со знаком, от -128 до +127) или uchar (без знака, от 0 до 255). Не следует путать с типом OpenCL int8, который представляет собой вектор из 8 32-разрядных целых чисел.

Для достойной производительности вы можете использовать вектор их версии, такие как char4 или char16, хотя это должно основываться на ваших измерениях производительности, а не на предположениях.

Обратите внимание, что вам необходимо знать о поведении переполнения, особенно при умножении Вам может потребоваться выполнить промежуточные операции над 16-битными значениями. (short / ushort / short4 / ushort16 / et c.) OpenCL также обеспечивает "насыщающее" сложение и вычитание и несколько других полезных целочисленных встроенных функций .

Мне неизвестна какая-либо «нативная» поддержка упакованных 4-битных целочисленных математических выражений в OpenCL или любой другой платформе GPGPU, или даже каких-либо расширений. Может быть, кто-то, имеющий опыт в этом, может внести свой вклад, но я предполагаю, что вам действительно нужно распаковать значения uchar, используя битовые сдвиги и маскировку, выполнить свои операции со значениями uchar, а затем упаковать обратно в 4-разрядные кусочки для хранения. Повышение скорости, скорее всего, будет связано с тем, что вы можете безопасно умножать, используя 8-битную логику c, а не 16-битные, чтобы уловить переполнение.

Answer 2

Я провел тест с некоторыми ядрами, чтобы увидеть, есть ли разница в производительности между int8 и char8:

    typedef int8 type_msg;
    //typedef char8 type_msg;
    #define convert_type_msg(x) convert_int8(x)

    __kernel void some_operation(__global  type_msg *in_buff,
                                 __global  type_msg *out_buff)
    {
        out_buff[get_global_id(0)] = in_buff[get_global_id(0)] +(type_msg)(2);
    }

Сначала, чтобы посмотреть, что происходит на GPU, я использовал CodeXL, чтобы получить код ассемблера.

Вот часть кода ассемблера, где используется int8:

    global_load_dwordx4  v[4:7], v[2:3], off
    global_load_dwordx4  v[8:11], v[2:3], off inst_offset:16
    v_add_co_u32  v0, vcc, s6, v0
    v_mov_b32     v2, s7
    v_addc_co_u32  v1, vcc, v2, v1, vcc
    s_waitcnt     vmcnt(0)
    v_add_u32     v8, 2, v8
    v_add_u32     v9, 2, v9
    v_add_u32     v10, 2, v10
    v_add_u32     v11, 2, v11
    global_store_dwordx4  v[0:1], v[8:11], off inst_offset
    v_add_u32     v2, 2, v4
    v_add_u32     v3, 2, v5
    v_add_u32     v4, 2, v6
    v_add_u32     v5, 2, v7
    global_store_dwordx4  v[0:1], v[2:5], off

А вот часть кода ассемблера, где используется char8:

    global_load_dwordx2  v[2:3], v[2:3], off
    s_waitcnt     vmcnt(0)
    v_lshlrev_b32  v4, 8, v3 src1_sel:BYTE_3
    v_lshrrev_b32  v5, 8, v3
    v_add_u32     v6, 2, v3 src1_sel:WORD_1
    v_add_u32     v4, 0x00000200, v4
    s_movk_i32    s0, 0x00ff
    v_lshlrev_b32  v7, 8, v2 src1_sel:BYTE_3
    v_add_u32     v5, 2, v5
    v_bfi_b32     v4, s0, v6, v4
    s_mov_b32     s1, 0x02010004
    v_lshrrev_b32  v6, 8, v2
    v_add_u32     v8, 2, v2 src1_sel:WORD_1
    v_add_u32     v7, 0x00000200, v7
    v_add_u32     v3, 2, v3
    v_perm_b32    v4, v5, v4, s1
    v_add_u32     v5, 2, v6
    v_bfi_b32     v6, s0, v8, v7
    v_add_co_u32  v0, vcc, s6, v0
    v_mov_b32     v7, s7
    v_addc_co_u32  v1, vcc, v7, v1, vcc
    v_perm_b32    v3, v3, v4, s1
    v_add_u32     v2, 2, v2
    v_perm_b32    v4, v5, v6, s1
    v_perm_b32    v2, v2, v4, s1
    global_store_dword  v[0:1], v3, off inst_offset:4
    global_store_dword  v[0:1], v2, off

Я не специалист по ассемблеру, но, насколько я могу судить, в обоих случаях было выполнено 8 дополнений с использованием операции v_add_u32. Также char8, кажется, требует больше операций, таких как v_perm_b32 и v_bfi_b32. Возможно, кто-то может объяснить, что они делают.

Единственное преимущество использования char8 заключается в том, что требуется меньший доступ к глобальной памяти. Например, существует только один global_load_dwordx2 доступ для char8, но 2 global_load_dwordx4 доступа для int8.

Таким образом, с точки зрения производительности, возможно, char8 немного медленнее для вычислительных ограниченных алгоритмов, но быстрее для алгоритмов, ограниченных памятью.

Для проверки анализа я построил небольшой эксперимент, где арифметика c является узким местом. Чтобы компилятор не слишком упрощал for-l oop, я добавил несколько ветвлений внутри него.

    typedef int8 type_msg;
    #define convert_type_msg(x) convert_int8(x)

    //typedef char8 type_msg;
    //#define convert_type_msg(x) convert_char8(x)

    __kernel void some_complex_operation(__global  char8 *in_buff,
                                         __global  char8 *out_buff)
    {
        type_msg res = in_buff[get_global_id(0)];
        for(int i=0; i<1000000; i++)
        {
            res += select((type_msg)(-1), (type_msg)(4), res<(type_msg)100);
        }
        out_buff[get_global_id(0)] =(type_msg) res;
    }

В моей системе среднее время (100 раз) для

• int8 равно 0,0558 с c
• char8 равно 0,0754 с c
• short8 равно 0,0738 с c
• long8 равен 0,1105 se c

Так что char8 потребляет примерно на 35% больше времени. Это подтверждает наблюдение, что на языке ассемблера генерируется больше инструкций для char8. Тем не менее, некоторые профессиональные объяснения для дополнительных утверждений ассамблеи было бы неплохо.