Постановка неуклюжих массивов длины в разделяемой памяти - PullRequest
Новая
       3

Постановка неуклюжих массивов длины в разделяемой памяти

1 голос
/ 06 ноября 2011

x - массив длины N в глобальной памяти, управляемый ядром cuda / opencl из k блоков каждого из w потоков (так что k = ceil (N / w)).Каждый блок в ядре имеет локальный общий массив xlocal длины w.Задача состоит в том, чтобы каждый блок загружал свой кусок x в xlocal.

Если w точно делит N, то мы можем сделать это: 

int lid = threadIdx.x;
int gid = threadIdx.x + (blockIdx.x * blockDim.x);
xlocal[lid] = x[gid];

Если нет, то имеем (N% w) избыточные потоки в последнем блоке.Как мы должны иметь с ними дело?Я могу думать о следующих опциях:

  1. Malloc большей длины для х.то есть, выделите k * w элементов вместо N. Это полезно, потому что приведенный выше код будет просто работать.К сожалению, я не думаю, что в cuda или opencl есть эквивалент realloc.

  2. Перед загрузкой проверьте диапазон.Это хорошо, потому что нам не нужно возиться с распределением x.Но досадно добавлять работу в большинство потоков только из-за краевого условия. 

    if (gid < N) xlocal[lid] = x[gid];

  3. Загрузка из x по модулю N, так что избыточные потоки обертываются вокруг: 

    xlocal[lid] = x[gid%N];

Есть еще какие-нибудь мысли по решению этой проблемы?

Некоторые тесты

Вот некоторые результаты, сравнивающие вариант (2) Rangecheck (синий)против варианта (3) загрузка по модулю N (красным).

Мы фиксируем размер блока из 32 потоков и варьируем N от 45,6 до 45,6 К + 32, чтобы получить от 0 до 32 избыточных потоков в последнем блоке соответственно,Тест запускает простое ядро, которое предварительно загружает общий массив из глобальной памяти.График слева (/ справа) загружает один (/ три) элемент (ы) на поток.Я скомпилировал с флагами cuda 3.2.16 -O2 и запустил на карте Tesla M2070.

runtime

1 Ответ

2 голосов
/ 06 ноября 2011

  1. Вы можете выделить больше х с хоста.Затем вы должны подумать о дополнительном времени копирования, которое может быть введено без использования, а также о площади памяти.Кроме того, ваш код потеряет смысл и структуру.

  2. :

  3. с помощью этой опции вы добавляете дополнительную работу по вычислению gid%N к каждому потоку, что именно то, что вы пытаетесь избежать, плюс еще одна дополнительная копия из глобальной памяти (это может не повредить, потому что копия может быть объединена, но все же).На мой взгляд, 2 (или, возможно, 3) - ваш лучший вариант, вы бы добавили только пару инструкций в каждую ветку.Не о чем беспокоиться, учитывая, что ваш код останется ясным и понятным.

    Вам следует избегать варианта 1.

...