CUDA Dynami c Индексирование

Question

Я заметил, что использование динамических c индексов снижает скорость кода CUDA в 12 раз - см. Следующий пример:

__global__ void static3Ops(int start, int end, const float* p, const int* prog_dont_use, float* c)
{
    int i = threadIdx.x;
    float buf[5];
    buf[0] = 1.0e7;
    buf[1] = c[i];
    const int prog[] = { 0,1,2,3,4,5 };

    for (long j = start; j < end; j++) {
        buf[2] = p[j];
        buf[3] = buf[prog[0]] + buf[prog[1]];
        buf[4] = buf[prog[2]] - buf[prog[3]];
        buf[1] = buf[prog[4]] * buf[prog[5]];
    }
    c[i] = buf[1];
}

в 12 раз быстрее, чем

__global__ void static3Ops(int start, int end, const float* p, const int* prog, float* c)
{
    int i = threadIdx.x;
    float buf[5];
    buf[0] = 1.0e7;
    buf[1] = c[i];

    for (long j = start; j < end; j++) {
        buf[2] = p[j];
        buf[3] = buf[prog[0]] + buf[prog[1]];
        buf[4] = buf[prog[2]] - buf[prog[3]];
        buf[1] = buf[prog[4]] * buf[prog[5]];
    }
    c[i] = buf[1];
}

Любой намек, как минимизировать эти накладные расходы? Динамическая c природа - ключевая особенность моего кода ... поэтому я вряд ли обойдусь без нее ...

Обратите внимание, что загрузка ЦП составляет всего около 20%.

Answer 1

Две возможности, о которых я могу думать:

Если прога - это небольшой массив : используйте свое собственное решение! т.е. используйте prog так же, как это определено в верхнем примере, если prog действительно массив с небольшим количеством элементов (как в вашем примере). Но ваш комментарий о "динамической природе c - это ключевая особенность моего кода" - звучит так, будто это не вариант для вас. Когда я изменяю const int prog[] = { 0,1,2,3,4,5 } на int prog_0 = 0, prog_1 = 1, ... и использую prog_0, prog_1, ... вместо prog[], я получаю ту же производительность. Это указывает на то, что значения prog[] хранятся непосредственно в регистрах без использования глобальной памяти. Если prog не является небольшим массивом или не известен во время компиляции, этот метод может привести к интенсивному использованию локальной памяти и значительно ухудшить производительность.

Если prog - большой массив : Потоки загружают prog параллельно в разделяемую память, а затем соответственно получают доступ к разделяемой памяти в остальной части вашего ядра (уровень блока tiling ).

__shared__  int prog_sh[6]; // or dynamically allocate if size is not known
int i = threadIdx.x;
if (i < 6)
    prog_sh[i] = prog[i];
__syncthreads();

// and then use prog_sh instead of prog....

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

Answer 2

Избегайте псевдонимов указателей

Первый бизнес: Используйте __restrict на всех ваших указателях! Это супер важно! Прочитайте об этом здесь:

Совет CUDA Pro: Оптимизируйте, чтобы избежать наложения указателей

Теперь, помимо этого ...

Играйте с вашими шаблонами доступа и попробуйте улучшить локальность памяти

Если:

• Размер prog ограничен небольшим значением, а
• Доступ к prog индексы, известные во время компиляции (т. е. не значение, а индекс)

Тогда:

• Использовать локальный ядро, простой массив C в стиле, или std::array -подобный класс, например kat::array из ветки разработки библиотеки cuda-kat (из-за раскрытия информации: это библиотека, над которой я работаю, поэтому я пристрастен здесь. реализация массива довольно стабильна). Загрузите их значения из указателя prog, полученного в качестве параметра.
• Тщательно расположите данные в памяти, чтобы можно было объединить загрузку в массив prog. Так, например, первый элемент prog для первого потока, затем первый элемент для второго потока et c, вплоть до первого элемента prog 31-го потока.
• Выполните все загрузки в прогу перед использованием любого из значений.

Если:

• Размер prog не ограничен небольшим значением, но
• Вы можете организовать использование prog таким образом, чтобы для каждого небольшого отрезка его ограниченной длины доступ к нему был с фиксированным смещением относительно некоторой базовой линии (например: l oop через i, на итерации i мы получаем доступ prog[k*i + 1], prog[k*i + 3], prog[k*i + 4] only),

Затем:

• Сделайте то же самое, что и в предыдущем случае, но для каждого отрезка фиксированной длины prog.

Если:

• prog не такой маленький, но не такой большой (от десятков элементов до тысяч элементов на поток), и
• шаблон доступа в него произвольный, произвольный или зависит от данных

Затем:

• Загрузка prog в общую память rst.
• Убедитесь, что загружаете его, чтобы не возникали конфликты банков, т. е. эквивалент каждого потока блока prog в общей памяти должен целиком содержаться в одном банке.

Если ничего из вышеперечисленного не выполняется, то:

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

Методическое примечание

Всегда помните, что все, что вы делаете - профиль и проанализируйте это , не соглашайтесь только на итоговый номер. И попробуйте разбить изменения и профилировать их отдельно. Например - сначала добавьте __restrict и посмотрите, что это даст вам. CUDA "nSight compute" также должен сообщать вам, где находятся ваши узкие места (хотя и не то, что с ними делать ...)

Answer 3

Идеи для улучшения скорости:

Если возможно, рассчитайте вариации прог в коде CUDA, используя threadIdx и blockIdx. Вычисления выполняются быстрее, чем доступ к памяти.

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

Так что, если у вас есть блоки размером 128 потоков, а я содержит номер потока:

__shared__ float buf[128 * 6];
buf[0] -> buf[0*128 + i];
buf[1] -> buf[1*128 + i];
buf[prog[0]] -> buf[prog[0]*128 + i];
...

Поскольку размер блока (128) делится на 32, каждый поток внутри деформации получает доступ к другому банку совместно используемой памяти, даже если индекс программы отличается. поток 0 всегда обращается к bank0 и т. д.

Альтернатива

Попытка хранить буферы непосредственно в регистрах вместо общей памяти: buf0, buf1, buf2, ...

Как получить к ним доступ по индексу? Просто напишите встроенную функцию или макрос с переключателем.

Есть 6 * 6 * 6 * 6 * 6 * 6 возможностей. Вы можете попробовать оптимизировать, сгенерировав код для 36 или 216 возможностей, а затем просто вызвать подходящий вариант. Например,

switch(prog01) {
case 0: buf3 = buf0 + buf0; break;
case 1: buf3 = buf0 + buf1; break;
...
case 6: buf3 = buf1 + buf0; break;
...
}

Но, возможно, это быстрее, если вы сделаете 6 переключателей по 6 случаев в каждом, тогда у вас будет меньше случаев / сравнений / переходов.

Лучше всего будет: сделать половину переключатели (например, 216) вне l oop, половина переключателей внутри одного из циклов 216.

Возможно, было бы еще лучше создать указатели функций устройства вне l oop и вызвать соответствующая функция, которая выбирает буф. Но тогда эти функции должны будут выбирать из буфера в качестве параметров функции вместо локальных переменных. Надеемся, что они по-прежнему эффективно хранятся в регистрах.

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

Пожалуйста, сравните и протестируйте с полными деформациями (не только с одним потоком), чтобы получить более реалистичные результаты c, включая банковские коллизии в общей памяти .

Answer 4

Спасибо всем за ваши подсказки!

, пока самый быстрый код, который я нашел, выглядит следующим образом:

_global__ void static3OpsShared(int start, int end, const float* prices, const int* __restrict__ prog, float* c)
{
    int i = threadIdx.x;
    __shared__ float buf[5];
    buf[0] = 1.0e7;
    buf[1] = c[i];
    // I never use more than 6 values of prog in a single thread - but each thread has its own set
    // values of prog are ranging from 0...5 
    // Performance needs to focus on what happens within the following loop typically having over 10000 iterations
    for (long j = start; j < end; j++) { 
        buf[2] = prices[j];
        buf[3] = buf[prog[0]] + buf[prog[1]];
        buf[4] = buf[prog[2]] - buf[prog[3]];
        buf[1] = buf[prog[4]] * buf[prog[5]];
    }
    c[i] = buf[1];
}

(Пожалуйста, на мгновение проигнорируйте индексирование общей памяти - я запустил это пока что с одним потоком)

с использованием регистров для prog [0] ... prog [5] в виде

r0 = prog[0];

и использование buf[r0] instead of buf[prog[0]], кажется, выполняется оптимизатор.

Наибольшее улучшение я получил благодаря использованию общей памяти для buf []. ограничение не помогло как-то. Особое ограничение не относится к buf, так как значения должны использоваться повторно.

Мой вывод таков: - Если вместо buf [] можно использовать регистры, код будет примерно в 5 раз быстрее.