Оптимизация ядра OpenCL

Question

Я работаю в ядре с большим массивом беззнаковых символов, создаю объект памяти с помощью clCreateBuffer. Затем я скопировал через clEnqueueWriteBuffer кусок неподписанных символов в этот объект памяти. И затем я вызываю в цикле ядро, которое читает из этого объекта памяти, выполняет некоторую логику и записывает новые данные в одно и то же место (я не вызываю clEnqueueWriteBuffer или clEnqueueReadBuffer в этом цикле). Вот код ядра:

__kernel void test(__global unsigned char *in, unsigned int offset) {
    int grId = get_group_id(0);
    unsigned char msg[1024];
    offset *= grId;

    // Copy from global to private memory
    size_t i;
    for (i = 0; i < 1024; i++)
        msg[i] = in[ offset + i ];

    // Make some computation here, not complicated logic    

    // Copy from private to global memory
    for (i = 0; i < 1024; i++)
        in[ offset + i ] = msg[i];
}

Когда цикл завершен (цикл выполняется приблизительно 1000 раз), я считываю результат из объекта памяти через clEnqueueReadBuffer.

Можно ли оптимизировать этот код?

Answer 1

Некоторые предложения:

• сделать один in += get_group_id(0) * offset в начале ядра.
• читать 4 символа за раз (работа с uchar4 или uint).
• если возможно, обрабатывайте также 4 символа одновременно.
• с частным массивом 1 КБ в каждом потоке, размер рабочей группы и занятость будут сильно ограничены, может быть более эффективно запускать меньше потоков, обрабатывающих меньшеchars.
• кажется, что все потоки в каждой группе будут обрабатывать одни и те же данные;это может быть не то, что вы имели в виду.

Answer 2

Вы можете попробовать векторную версию (uchar8 вместо uchar), но компилятор в любом случае может оптимизировать ее.Самый важный профиль ваш код все время и эксперимент.

edit

Кажется, даже uchar16 теперь поддерживается: http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/vectorDataTypes.html

Answer 3

Для оптимизации вам нужно объяснить, какие расчеты вы делаете.Наибольшую выгоду для производительности можно получить, сгруппировав свои расчеты в рабочие группы и предоставив им возможность работать с локальной памятью.Вам нужно уделять много внимания размеру вашей личной памяти (наименьшей) и локальной памяти (маленькой).

Как часто называется ваше ядро?Все ли ядра используют одни и те же данные?Можно подумать о буфере локальной памяти, где все потоки в рабочей группе считывают часть данных в локальную память и впоследствии делятся данными.Вам нужно уделить немного внимания синхронизации.

Я предлагаю взглянуть на примеры поставщиков SDK.Я знаю только nVidia SDK.Примеры там довольно сложные, но очень интересные для чтения.

Изменение на векторные типы, такие как float4, должно подходить для плат ATI.Говорят, что nVidia лучше всего работает со скалярами и внутренней оптимизацией компилятора.Это что-то для тонкой настройки позже с помощью профилировщика.Оптимизация памяти позволяет повысить производительность.

Answer 4

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

Также может помочь использование общей памяти на чипах Nvidia (если ваша текущая локальная память не использует общую память по умолчанию)