- О. Федор
- Постоялец
Обычно считается, что этот порядок неопределен, но недавно в одной из статей (линка к сожалению дать пока не могу, забыл) было замечание, что вроде бы в Ферми, этот порядок можно определить, во всяком случае он предсказуем.
Поясню важность этой проблемы.
Если бы можно было быть всегда уверенным, что i-тый блок выполнится раньше i + 1 -го, то можно было бы сильно упростить многие алгоритмы, например, редикс сортировки.
Какая инфа есть на этот счет?
Нашел статью, о которой упомянул http://nvworld.ru/articles/cuda_parallel/page2/
Там между прочим сказано:
Но нити не обязаны исполняться совсем независимо друг от друга. Для нитей внутри одного блока (это примерно 256 нитей), есть стандартные инструкции синхронизации исполнения Wait — ждать, когда все нити синхронизируются. Напомню, что нити абсолютно произвольны в своем поведении и совсем не похожи на SSE
-инструкции. Есть так же инструкции атомарной записи данных в память блока, для обмена данными между нитями.
То есть, нити в одном блоке могут переплетаться между собой. А вот блоки, в рамках CUDA
-идеологии, предполагаются более независимыми. Они как будто бы исполняются на разных, далеко друг от друга расположенных, кластерах одного большого суперкомпьютера. Каждый из этих кластеров получил свое задание и, в каком порядке они исполняются, программисту неизвестно. Но некоторую синхронизацию, в частных случаях, можно организовать с помощью флагов в глобальной памяти. Например, организовать счетчик выполнения и старта блоков и, в зависимости от номера, блок будет выполнять свою часть задачи. Например, последний исполняющийся блок узнает, что он последний и может выполнить специальный завершающий этап работы. Или можно запустить новую программу, новый этап вычислений, который воспользуется сохраненными в памяти GPUрезультатами работы предыдущей программы.
Вопросы собственно такие:
- Как организовать синхронизацию с помощью флагов в глобальной памяти ?
- Как организовать счетчик выполнения и старта блоков ?
О. Федор
О счетчике выполнения блоков я вот что подумал - это свой независимый счетчик вместо blockIdx. Блок стартует - и читает этот счетчик, который нужно атомарно увеличивать (скинуть перед выполнением кернела). А вот насчет синхронизации не помню, надо поглядеть что есть, забылось.
_vasa_
И где менять этот blockIdx? Если из блока - то там это константа, ты даже можешь написать const ib = blockIdx.x например.
Синхронизация на атомиках только приходит на ум.
int atomicExch(int* address, int val); int atomicAdd( int* address, int val);
Так в том и дело что менять blockIdx не нужно (и невозможно), вместо него все потоки читают значение счетчика (нашего собственного).
Кроме этого нужно определить, какой из потоков блока меняет наш счетчик - например 0-й поток после syncthreads(), когда все потоки блока прочли текущий номер.
_vasa_
> Так в том и дело что менять blockIdx не нужно (и невозможно), вместо него все
> потоки читают значение счетчика (нашего собственного).
> Кроме этого нужно определить, какой из потоков блока меняет наш счетчик -
> например 0-й поток после syncthreads(), когда все потоки блока прочли текущий
> номер.
Как ты себе это представляешь?
Сооруди простенький, но рабочий код, там посмотрим.
О. Федор
Ок, немного позже поставим эксперимент.
О. Федор
>> Как организовать синхронизацию с помощью флагов в глобальной памяти ?
Ну как и сказал _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 собственно равен числу потоков которое железо может держать.
Соответственно, используя атомики можно сделать какую хочешь синхронизацию.
Например в конце цикла можно проверить нужное условие и как-то подождать если оно не выполнено.
Опиши чуть подробнее задачу и конкретнее - что именно за синхронизация требуется.
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 двоичным разрядам. Так вот, чтоб предварительно отсортированный порядок не нарушился, желательно иметь предсказуемый порядок вызовов блоков.
Пока я просто все разбиваю на дипазоны, так чтоб каждый блок обрабатывал строго свой диапазон. Недостаток этого метода - нужно выделять больше памяти для вспомогательных массивов, указывающих на начала диапазонов.
_vasa_
> Ок, немного позже поставим эксперимент.
Хорошо, посмотрим, может действительно что получится.
Можно еще наверное попробовать разбить на разные kernel-ы и запускать в очереди. Fermi же умеет по-умному вродже манаджить, конкуретное исполнение там и проч.
Но очередь гарантирует что что-то выполнется раньше чего-то другого.
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 гб.
О. Федор
Совсем завал у меня, звиняй, найду время займусь этим обязательно :)
OK, но мне кажется я понял, что ты собираешься сделать.
Управлять порядком выполнения блоков по содержимому одной ячейки в глобальной памяти (т.е. запускать выполнение блока только при определенном значении счетчика).
Если так, то боюсь, что весь смысл теряется, поскольку это будут встроенные в код не хилые тормоза.
Тема в архиве.