ПрограммированиеФорумОбщее

Порядок выполнения блоков в cuda

#0
11:04, 4 ноя 2010

Обычно считается, что этот порядок неопределен, но недавно в одной из статей (линка к сожалению дать пока не могу, забыл) было замечание, что вроде бы в Ферми, этот порядок можно определить, во всяком случае он предсказуем.
Поясню важность этой проблемы.
Если бы можно было быть всегда уверенным, что i-тый блок  выполнится раньше i + 1 -го, то можно было бы сильно упростить многие алгоритмы, например, редикс сортировки.
Какая инфа есть на этот счет?

#1
12:10, 4 ноя 2010

Нашел статью, о которой упомянул http://nvworld.ru/articles/cuda_parallel/page2/

Там между прочим сказано:

Но нити не обязаны исполняться совсем независимо друг от друга. Для нитей внутри одного блока (это примерно 256 нитей), есть стандартные инструкции синхронизации исполнения Wait — ждать, когда все нити синхронизируются. Напомню, что нити абсолютно произвольны в своем поведении и совсем не похожи на SSE
-инструкции. Есть так же инструкции атомарной записи данных в память блока, для обмена данными между нитями.
То есть, нити в одном блоке могут переплетаться между собой. А вот блоки, в рамках CUDA
-идеологии, предполагаются более независимыми. Они как будто бы исполняются на разных, далеко друг от друга расположенных, кластерах одного большого суперкомпьютера. Каждый из этих кластеров получил свое задание и, в каком порядке они исполняются, программисту неизвестно. Но некоторую синхронизацию, в частных случаях, можно организовать с помощью флагов в глобальной памяти. Например, организовать счетчик выполнения и старта блоков и, в зависимости от номера, блок будет выполнять свою часть задачи. Например, последний исполняющийся блок узнает, что он последний и может выполнить специальный завершающий этап работы. Или можно запустить новую программу, новый этап вычислений, который воспользуется сохраненными в памяти GPUрезультатами работы предыдущей программы.

Вопросы собственно такие:

- Как организовать синхронизацию с помощью флагов в глобальной памяти ?
- Как организовать счетчик выполнения и старта блоков ?

#2
15:09, 4 ноя 2010

О. Федор
О счетчике выполнения блоков я вот что подумал - это свой независимый счетчик вместо blockIdx. Блок стартует - и читает этот счетчик, который нужно атомарно увеличивать (скинуть перед выполнением кернела). А вот насчет синхронизации не помню, надо поглядеть что есть, забылось.

#3
15:12, 4 ноя 2010

_vasa_
И где менять этот blockIdx? Если из блока - то там это константа, ты даже можешь написать const ib = blockIdx.x например.

#4
15:23, 4 ноя 2010

Синхронизация на атомиках только приходит на ум.

int atomicExch(int* address, int val);
int atomicAdd(int* address, int val);
#5
15:25, 4 ноя 2010

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

#6
15:59, 4 ноя 2010

_vasa_
> Так в том и дело что менять blockIdx не нужно (и невозможно), вместо него все
> потоки читают значение счетчика (нашего собственного).
> Кроме этого нужно определить, какой из потоков блока меняет наш счетчик -
> например 0-й поток после syncthreads(), когда все потоки блока прочли текущий
> номер.

Как ты себе это представляешь?
Сооруди простенький, но рабочий код, там посмотрим.

#7
16:43, 4 ноя 2010

О. Федор
Ок, немного позже поставим эксперимент.

#8
18:30, 4 ноя 2010

О. Федор
>> Как организовать синхронизацию с помощью флагов в глобальной памяти ?
Ну как и сказал _vasa_ -  атомики.

По поводу порядка выполнения блоков не знаю ничего. Ты в документации смотрел?
Просто это нетривиально потому что разные блоки могут выполняться разное время.

Может быть тебе поможет persistent threads + атомики. М .б. ты знаешь, но я все-равно напишу.

Запускаешь ровно столько warp-ов сколько может вместить железо (надо посчитать).
Каждый warp выполняет работу не для одного а для большего числа warp-ов.

for(int warp=0;warp < N; warp++)
{
  uint tid = BLOCK_SIZE*blockIdx.x + threadIdx.x + warp*step;
  // code here
}

step собственно равен числу потоков которое железо может держать.
Соответственно, используя атомики можно сделать какую хочешь синхронизацию.
Например в конце цикла можно проверить нужное условие и как-то подождать если оно не выполнено.
Опиши чуть подробнее задачу и конкретнее - что именно за синхронизация требуется.

#9
11:02, 5 ноя 2010

FROL
> Ну как и сказал _vasa_ - атомики.

Что-то я в это не верю. Атомарная функция конечно хорошая штука, но здесь это работать не будет.

> По поводу порядка выполнения блоков не знаю ничего. Ты в документации смотрел?

Смотрел сегодня документацию. Безнадега. Вроде нашел одну, близкую по смыслу функцию

cudaDeviceBlockingSync

есть еще
cudaStreamSynchronize
но это для контроля выполнения в тредах cpu.

> Просто это нетривиально потому что разные блоки могут выполняться разное
> время.

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

> Может быть тебе поможет persistent threads + атомики. М .б. ты знаешь, но я
> все-равно напишу.

Не, не знаю, а что это такое?

>
> Запускаешь ровно столько warp-ов сколько может вместить железо (надо
> посчитать).
> Каждый warp выполняет работу не для одного а для большего числа warp-ов.
>
> for(int warp=0;warp < N; warp++)
> {
> uint tid = BLOCK_SIZE*blockIdx.x + threadIdx.x + warp*step;
> // code here
> }

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

> Опиши чуть подробнее задачу и конкретнее - что именно за синхронизация требуется.

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

#10
11:03, 5 ноя 2010

_vasa_
> Ок, немного позже поставим эксперимент.

Хорошо, посмотрим, может действительно что получится.

#11
15:59, 5 ноя 2010

Можно еще наверное попробовать разбить на разные kernel-ы и запускать в очереди. Fermi же умеет по-умному вродже манаджить, конкуретное исполнение там и проч.
Но очередь гарантирует что что-то выполнется раньше чего-то другого.

#12
15:50, 6 ноя 2010

FROL
> Можно еще наверное попробовать разбить на разные kernel-ы и запускать в
> очереди.

В принципе я так и делаю, см. например http://www.gamedev.ru/code/forum/?id=89113&page=31#m464,
там тоже первый раз кернель вызывается для всех блоков, а второй раз по сути только для одного мультипроцессора

    //вычисление частных результатов блоков
    accumulate <blockSize, Polyce, Tin,  Tout><<<blockSize, blockSize>>>(d_idata, d_odata, size);
    //вычисление общего результата всех блоков
    accumulate <blockSize, Polyce,  Tout,  Tout><<<1, blockSize>>>(d_odata, d_result, blockSize);

> Fermi же умеет по-умному вродже манаджить, конкуретное исполнение там и проч.

Это они хорошо придумали. Там даже есть возможность грузить в девайс и выгружать одновременно. Только мне кажется конкурентное использование кернелей лучше все же делать из разных тредов cpu

> Но очередь гарантирует что что-то выполнется раньше чего-то другого.

В принципе можно сделать полностью независимое выполнение разных программ на одном девайсе, хотя польза при нынешней доступной памяти весьма сомнительна. Тогда нужно хотя бы 4-6 гб.

#13
10:22, 12 ноя 2010

О. Федор
Совсем завал у меня, звиняй, найду время займусь этим обязательно :)

#14
11:23, 13 ноя 2010

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

ПрограммированиеФорумОбщее

Тема в архиве.