Каков наилучший способ реализовать небольшую таблицу поиска в ядре OpenCL?

В моем ядре необходимо сделать большое количество случайных обращений к небольшой таблице поиска (всего 8 32-битных целых чисел). Каждое ядро ​​имеет уникальную таблицу поиска. Ниже приведена упрощенная версия ядра, иллюстрирующая использование таблицы поиска.

__kernel void some_kernel(  
    __global uint* global_table,
    __global uint* X,
    __global uint* Y) {

    size_t gsi = get_global_size(0);
    size_t gid = get_global_id(0);

    __private uint LUT[8]; // 8 words of of global_table is copied to LUT

    // Y is assigned a value from the lookup table based on the current value of X
    for (size_t i = 0; i < n; i++) {
        Y[i*gsi+gid] = LUT[X[i*gsi+gid]];
    }   
}

Из-за небольшого размера я получаю лучшую производительность, сохраняя таблицу в пространстве памяти __private. Однако из-за случайного характера обращения к таблице поиска производительность по-прежнему сильно снижается. Если код таблицы поиска удален (например, заменен простой арифметической операцией), хотя ядро ​​и выдаст неверный ответ, производительность возрастет более чем в 3 раза.

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

[править] Обратите внимание, что максимальное значение X равно 7, но максимальное значение Y равно 2^32-1. Другими словами, используются все биты таблицы поиска, поэтому ее нельзя упаковать в меньшее представление.


opencl
person Dustin    schedule 20.10.2011    source источник
comment
Просто чтобы убедиться, что я правильно понимаю, являются ли LUT [] и X [] уникальными для каждого отдельного рабочего элемента?   -  person Adam S.    schedule 20.10.2011
comment
Пробовали ли вы использовать __constant памяти для своих справочных таблиц? Графические процессоры часто реализуют отдельные кэши и пути доступа к памяти для постоянной памяти, чтобы ускорить такие вещи, как общие справочные таблицы.   -  person user57368    schedule 21.10.2011
comment
Я пытался передать global_table как __constant, но по какой-то причине это не помогло производительности. Каждое ядро ​​работает с независимым разделом global_table, X и Y (в зависимости от идентификатора потока).   -  person Dustin    schedule 21.10.2011
comment
Если это так, __constant, скорее всего, не поможет, потому что, по крайней мере, на большинстве графических процессоров, которые я видел, это в основном небольшой кеш для глобальной памяти. Если каждый поток считывает разные значения, кеш ничем не поможет.   -  person Adam S.    schedule 21.10.2011
comment
В руководстве Intel OpenCL рассказывается об использовании общей локальной памяти для LUT. software.intel. com/en-us/articles/ В целом может быть полезно.   -  person Peter Cordes    schedule 18.06.2015

Ответы (2)


arrow_upward
4
arrow_downward

Самое быстрое решение, которое я могу придумать, - это вообще не использовать массивы: вместо этого использовать отдельные переменные и использовать какую-то функцию доступа для доступа к ним, как если бы они были массивом. IIRC (по крайней мере, для компилятора AMD, но я уверен, что это верно и для NVidia): обычно массивы всегда хранятся в памяти, а скаляры могут храниться в регистрах. (Но мой разум немного затуманен в этом вопросе, я могу ошибаться!)

Даже если вам нужен гигантский оператор switch:

uint4 arr0123, arr4567;
uint getLUT(int x) {
    switch (x) {
    case 0: return arr0123.r0;
    case 1: return arr0123.r1;
    case 2: return arr0123.r2;
    case 3: return arr0123.r3;
    case 4: return arr4567.r0;
    case 5: return arr4567.r1;
    case 6: return arr4567.r2;
    case 7: default: return arr4567.r3;
    }
}

... вы все равно можете выйти вперед по производительности по сравнению с массивом __private, поскольку, если предположить, что все переменные arr помещаются в регистры, это связано исключительно с ALU. (Конечно, при условии, что у вас достаточно свободных регистров для переменных arr.)

Обратите внимание, что некоторые цели OpenCL даже не имеют имеют личную память, и все, что вы там объявляете, просто попадает в __global. Использование регистрового хранилища — еще большая победа.

Конечно, этот подход LUT, вероятно, будет медленнее инициализировать, поскольку вам потребуется как минимум два отдельных чтения памяти для копирования данных LUT из глобальной памяти.

person rtollert    schedule 20.10.2011
comment
Раньше я исключал такое решение, потому что боялся вызвать расхождение деформации, но я попробую, просто чтобы посмотреть. Почему вы использовали 2 вектора uint4 вместо одного uint8 или всего 8 переменных? - person Dustin; 20.10.2011
comment
Это повысило производительность до 15%... не от 200% до 300%, на которые я надеялся, но каждый бит помогает. - person Dustin; 21.10.2011
comment
Я использовал uint4s, потому что я идиот. :) Думаю, uint8s тоже подойдет. Сбор значений в одну логическую переменную означает, что вы можете инициализировать таблицу с помощью vload8() вместо выполнения (возможно) восьми отдельных обращений к памяти, по одному для каждой переменной. Но вы это уже знали. - person rtollert; 21.10.2011
comment
На самом деле расхождений почти не было бы. Операторы переключения обычно реализуются как арифметические jmps (вычисляемый переход). Поскольку вы не выполняете никакой работы в случаях, вы будете расходиться максимум в одной инструкции, что совсем неплохо. - person imallett; 22.06.2012

arrow_upward
1
arrow_downward

Как заявил rtollert, реализация должна решить, помещать ли LUT[] в регистры или в глобальную память. Обычно массивы в ядре - нет-нет, но поскольку он маленький, трудно сказать, где он будет размещен. Предполагая, что LUT [] помещается в регистры, я бы сказал, что причина, по которой это занимает много времени по сравнению с простой арифметической операцией, заключается не в том, что к ней обращаются случайным образом, а в том, что каждый рабочий элемент создает дополнительные 8 (Редактировать: очевидно, намного больше) глобальных считывает X для расчета индекса LUT. В зависимости от того, что пропущено, вы могли бы сделать что-то вроде Y[i*gsi+gid] = global_table[someIndex + X[i*gsi+gid]]];?

person Adam S.    schedule 20.10.2011
comment
8 чтений из глобальной памяти объединяются, а также вне цикла. Поскольку n велико (обычно около 1024), накладные расходы на чтение для копирования LUT эффективно амортизируются. - person Dustin; 21.10.2011