threads, warps && blocks

Здравствуйте, раньше когда работал с CUDA не задавался вопросами распределения, простой указывал разделение, как в примере с возведением в квадрат элементов массива, но возникла проблема в более серьезной программе, где это надо учитывать и в связи с этим несколько вопросов:
0. когда я загружаю функцию на обработку я так понимаю что она разбивается на множество копий каждая копия загружается в нить и у каждой идет свой набор параметров для обработки, при этом при загрузке я указываю количество необходимых блоков и размер каждого блока, если функция делает не сложный расчет то я сразу же могу рассчитывать на результаты, если нет то как это можно проверить, чтобы не обрабатывать данных которых нету, или управление передастся только когда дынные получаться?
1. Могут ли блоки иметь разный размер, если нет и допустим блок заполнен не полностью что будет с частью потоков которые не заполнены они не будут просто принимать участие в обработке или будут забиваться другими данными?
2. я начал загрузку нитей в блок и у меня там есть условие что допустим если результат выполнения какой-то функции вернет отрицательное значение то мне надо остановить загрузку нитей, их результаты мне уже не надо, сейчас я это делаю так:

  1. <br />
  2. __shared__ int  LastCount;//позиция символа<br />
  3. __device__ void CountLast(char  *main, char  *sub, int length, int numb, int start )<br />
  4. {<br />
  5.         int idx=length-threadIdx.x;<br />
  6.         if((idx>=start)  && (idx==LastCount))<br />
  7.         {<br />
  8.              if(main[numb+(length-idx)] ==sub[length-idx] )<br />
  9.                  LastCount--;<br />
  10.         }<br />
  11. }<br />

Перед вызовом этой функции пишу LastCount=length;
как себя поведет гпу при первом невыполнении условия он прервет загрузку нитей на обработку или их загрузит?
3. При запуске многих блоков на обработку они ставятся в очередь и GPU их потом берет из очереди и обрабатывает, как можно узнать размер этой очереди?

Forums: 

Возможно, я не все ваши

Возможно, я не все ваши вопросы понял правильно, если что не так - давайте уточнять.

0) Да, у вас есть множество копий, собранных в блоки (у блока - общая память). Количество потоков в блоке и самих блоков вы задаете на старте. При этом данные должны быть скопированы на GPU до старта kernel.

1) Блоки не могут иметь разный размер.

2) Остановить работу всех потоков из какого-то - можно через внешнюю синхронизацию (если нужно остановить, какой-то поток делает atomicAdd, все потоки проверяют значение этой переменной в глобальной памяти регулярно), но атомарные операции - довольно медленные. Остановить потоки блока - да, можно через shared memory, но остальные блоки об этом не узнают.

3) Да, грубо говоря ставятся в очередь. Количество одновременно исполняемых блоков на одном SP ограничено количеством регистров, потоков и shared memory (планировщик запустит столько блоков, чтобы им всего хватило), дальше умножаем на число SP, все что с первого раза не влезло - в очередь.

у меня возник ещё один вопрос

у меня возник ещё один вопрос аналогичный номеру 2, у меня есть функция:

  1. __shared__ int  LastCount;//позиция символа
  2. __device__ void CountLast(char  *main, char  *sub, int length, int numb, int start )
  3. {
  4.         int idx=length-threadIdx.x;
  5.         if((idx>=start)  && (idx==LastCount))
  6.         {
  7.              if(main[numb+(length-idx)] ==sub[length-idx] )
  8.                  LastCount--;
  9.         }
  10. }

эта функция для проверки элементов массивов с обратной стороны, вопрос следующий:
я запустил эту функцию и она выполняется в 50 потоках, но существует один нюанс мне надо чтобы с начал выполнилась первая часть проверки, тоесть проверились последние элементы, потом предпоследние элементы. мне подойдет любая задержка или же данные загружаются и выполняются абсолютно параллельно, тоесть одновременно выполняться будут первый, второй, ...последний потоки?

Потоки внутри warp

Потоки внутри warp выполняются строго одновременно, с точностью до инструкции. warp-ы внутри блока - перемешаны, часть блоков исполняется тоже одновременно на разных SP

Если вам нужно последовательное исполнение, то вам не на GPU, видеокарты быстры именно за счет распараллеливания (причем независимого, синхронизация будет снижать скорость очень сильно)