Индексы CUDA для петель со счетчиками
У меня есть вложенный цикл со счетчиком между ними.
Мне удалось использовать индексы CUDA для внешнего цикла, но я не могу придумать способа использовать больше параллелизма в циклах такого типа.
Есть ли у вас опыт работы с чем-то похожим на это?
int i = threadIdx.x + blockIdx.x * blockDim.x;
if (i < Nx) {
counter = 0;
for (k = 0; k < Ny; k++) {
d_V[i*Ny + k] = 0;
if ( d_X[i*Ny + k] >= 2e2 ) {
/* do stuff with i and k and counter i.e.*/
d_example[i*length + counter] = k;
...
/* increment counter */
counter++;
}
}
}
Проблема, которую я вижу, состоит в том, как бороться со счетчиком, как k может также быть проиндексирован в CUDA с threadIdx.y + blockIdx.y * blockDim.y
Решение
Наличие переменной счетчика / цикла, которая используется между итерациями цикла, является естественной противоположностью распараллеливанию. Идеальные параллельные циклы имеют итерации, которые могут выполняться в любом порядке, не зная друг друга. К сожалению, общая переменная делает ее зависимой от порядка и взаимно осведомленной.
Похоже, вы используете счетчик, чтобы упаковать d_example массив без пробелов. Подобные вещи могут быть более эффективными во время вычислений, тратя немного памяти; если вы позволите элементам d_example, которые не будут установлены, оставаться равными нулю из-за неэффективной упаковки d_exampleВы можете выполнить фильтр для d_example позже, после любых дорогостоящих вычислительных шагов.
Фактически, вы можете даже оставить фильтрацию модифицированному итератору при чтении массива, который просто пропускает любые нулевые значения. Если ноль является допустимым значением в массиве, просто используйте конкретное значение NaN или отдельный массив масок.
int i = threadIdx.x + blockIdx.x * blockDim.x;
if (i < Nx) {
for (k = 0; k < Ny; k++) {
d_V[i*Ny + k] = 0;
if ( d_X[i*Ny + k] >= 2e2 ) {
/* do stuff with i and k and counter i.e.*/
d_example[i*length + i*k] = k;
d_examask[i*length + i*k] = 1;
...
/* increment counter */
} else {
d_examask[i*length+i*k] = 0;
}
}
}
Другие решения
Обратите внимание, что вы можете использовать threadIDx.y в качестве второго индекса в вашем массиве. Для получения дополнительной информации смотрите здесь: http://www.cs.sunysb.edu/~mueller/teaching/cse591_GPU/threads.pdf
Например, если у вас есть блоки в двух измерениях, вы можете использовать threadix.x и threadix.y в качестве индикаторов и добавить смещение рабочей группы (blockidx.x * blockDim.x) в качестве вашего смещения.
Поскольку на графических процессорах ветвление очень дорого, и все потоки в данной рабочей группе всегда будут ожидать продолжения всех задач в группе, лучше просто вычислить все элементы и отбросить те, которые вам не нужны, если это возможно, это потенциально может избежать использования счетчика полностью. Если нет, лучшее решение — использовать функции атомарного приращения API CUDA на глобальном счетчике, как указано в его комментарии phoad.
Если это возможно, вы можете использовать cudpp или thrust (библиотеки, которые реализуют параллельные функции, такие как remove_if или compact — что-то, что у вас есть в примере).
Вы можете найти на этих страницах простые примеры, как их использовать. Я предпочитаю cudpp, потому что ИМХО быстрее, чем тяга.