Использование векторных типов для улучшения производительности ядра OpenCL - PullRequest
Новая
       294

Использование векторных типов для улучшения производительности ядра OpenCL

0 голосов
/ 23 октября 2019

У меня есть следующее ядро ​​OpenCL, которое копирует значения из одного буфера в другой, опционально инвертируя значение (аргумент 'инвертировать' может быть 1 или -1): - 

__kernel void extraction(__global const short* src_buff, __global short* dest_buff, const int record_len, const int invert)
{                                                                                                           
    int i = get_global_id(0); // Index of record in buffer                                              
    int j = get_global_id(1); // Index of value in record                                                       

    dest_buff[(i* record_len) + j] = src_buff[(i * record_len) + j] * invert;
}

Исходный буферсодержит одну или несколько «записей», каждая из которых содержит N (record_len) короткие значения. Все записи в буфере имеют одинаковую длину, а record_len всегда кратен 32.

Глобальный размер - 2D (количество записей в буфере, длина записи), и я выбрал его, как мне показалось,лучше всего использовать параллельную обработку GPU, причем каждый поток отвечает за копирование только одного значения в одну запись в буфере.

(Кстати, для локального рабочего размера установлено значение NULL, что позволяет OpenCL определять само значение).

После недавнего чтения векторов мне стало интересно, смогу ли я использовать их для улучшенияна спектакле? Я понимаю концепцию векторов, но я не уверен, как использовать их на практике, отчасти из-за отсутствия хороших примеров.

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

Нариск быть немного наивным здесь, могу ли я просто изменить два типа arg буфера на short16 и изменить второе значение в глобальном 2-D размере с «длина записи» на «длина записи / 16»? Приведет ли это к тому, что каждый поток ядра скопирует блок из 16 коротких значений между буферами?

1 Ответ

2 голосов
/ 23 октября 2019

Ваше наивное предположение в основном верно, хотя вы можете добавить подсказку компилятору, что это ядро ​​оптимизировано для векторного типа ( Раздел 6.7.2 спецификации ), в вашем случае выдобавит

атрибут ((vec_type_hint (short16)))

над вашей функцией ядра. Таким образом, в вашем примере у вас будет 

__attribute__((vec_type_hint(short16)))
__kernel void extraction(__global const short16* src_buff, __global short16* dest_buff, const int record_len, const int invert)
    {                                                                                                           
        int i = get_global_id(0); // Index of record in buffer                                              
        int j = get_global_id(1); // Index of value in record                                                       

        dest_buff[(i* record_len) + j] = src_buff[(i * record_len) + j] * invert;
    }

Вы правы в том, что ваше второе глобальное измерение должно быть разделено на 16, а ваш record_len также должен быть разделен на 16. Кроме того, если бы вы указалилокальный размер, вместо того, чтобы задавать его как NULL, вы также хотели бы разделить его на 16.

Однако следует учитывать и другие моменты.
Вы можете подумать, что выбор наибольшего размера вектора должен обеспечить наилучшую производительность,особенно с таким простым ядром. Но по моему опыту, это редко самый оптимальный размер. Вы можете попробовать запросить clGetDeviceInfo для CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT, но для меня это редко является точным (также, это может дать вам 1, что означает, что компилятор попытается выполнить автоматическую векторизацию или устройство не имеет векторного оборудования). Лучше всего попробовать разные размеры вектора и посмотреть, какой из них самый быстрый.

Если ваше устройство поддерживает автоматическую векторизацию, и вы хотите попробовать, это может помочь удалить ваш параметр record_len и заменить его на get_global_size (1) , поэтому компилятор / драйвер может позаботиться о делении record_len на любой размер вектора, который он выберет. В любом случае, я бы порекомендовал сделать это, предполагая, что record_len равен глобальному размеру, который вы дали этому измерению.

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

Наконец, для общих оптимизаций OpenCL, вы можете взглянуть на Рекомендации NVIDIA OpenCL. Руководство для оборудования NVidia или Руководство пользователя AMD APP SDK OpenCL для оборудования AMD GPU. NVidia один с 2009 года, и я не уверен, насколько сильно изменилось их оборудование с тех пор. Обратите внимание, что на самом деле говорится:

Архитектура CUDA является скалярной архитектурой. Следовательно, при использовании векторных типов и инструкций выигрыш в производительности отсутствует. Они должны использоваться только для удобства.

Более старое оборудование AMD (до GCN) получило выгоду от использования векторных типов, но AMD предлагает не использовать их на устройствах GCN (см. комментарий mogu ). Также, если вы ориентируетесь на процессор, будет использовать аппаратное обеспечение AVX, если доступно .

...