Помогите пожалуйста!

Решил попытаться ускорить одну свою программку при помощи CUDA, столкнулся с глюком. Процедура, приведенная ниже какбэ должна считать стоимость пути в графе:

  1. __global__ void gputest(int *x, float *y, int *z, float *c, float *d, float *p, int N)
  2.     {
  3.       int idx = blockIdx.x * blockDim.x + threadIdx.x;
  4.       int i,ths,next;
  5.       float t=0;  
  6.       if (idx<N)
  7.         {
  8.            ths=0;
  9.            for (i=1;i<=33;i++)
  10.              {
  11.                 next=x[idx*33+ths];//x - это замкнутые маршруты, которые состоят ровно из 33х шагов
  12.                 t=t+c[ths*33+next];//с - это стоимости переходов от одного узла графа к другому
  13.                 ths=next;
  14.              }
  15.           y[idx]=t;//y - это стоимости маршрутов
  16.        }
  17.   }

Такая же процедура на обычном CPU, вызываемая в цикле, прекрасно забивает массив y стоимостями, на GPU она забивает его значениями -431602080, причем, если сократить количество шагов цикла for до 20ти, то получаются в некоторых элементах правильные стоимости, в других значения -35659499650496332000, в третьих #QNAN0. Может кто-нибудь знает почему так происходит?
Мне 20 шагов мало, в реальной задаче их вообще сотни, а маршрутов тысячи. Помогите пожалуйста.

PS вверху в приведенном коде должен бы стоять знак меньше перед N, но в таком случае движок форума отсекает весь текст который после него (тоже своего рода глюк :))

Forums: 

1) код обрамлять тегами

1) код обрамлять тегами  <code>  
2) Вопрос: не могу понять а зачем тут цикл for если в его теле не используется переменная i ?

Затем чтобы повторить то что

Затем чтобы повторить то что в теле много раз. Как я уже написал выше, "много", в реальных задачах, измеряется сотнями. А если удастся увеличить производительность, с помощью куды, или еще как, то может это будут и тысячи.
По существу есть что сказать?

Скажите, а эти массивы x/y -

Скажите, а эти массивы x/y - это в глобальной памяти?

Я не очень знаю, что у вас в исходном x, но не возможна ли, случайно ситуация, когда несколько потоков пишут в одну и ту же память?

P.S. Позволил себе поставить тег code.

Спасибо за тег, я редко на

Спасибо за тег, я редко на форумах бываю поэтому не разбираюсь в этом. :)

С данными там происходит следующее, прямо по учебнику: есть переменные X_h/Y_h и X_d/Y_d для хоста и девайса соответственно. Память для X_d/Y_d выделяется CudaMalloc. Я проверял, X_d не меняется во время выполнения кернела.
Исходный x разбит на порции по 33, каждая такая порция - это маршрут, а каждое отдельное значение - шаг маршрута. Например если в нулевом элементе порции находится число 15, это означает что с 0-вого узла нужно переходить на 15-тый. Теоретически каждая порция обрабатывается своим потоком, потому там и написано x[idx*33+ths], т.е. обрабатывается в элемент ths idx-ной порции. В x ничего не пишется, только читается из него.
Вообще, эта процедура - часть алгоритма популяционного поиска для решения специфической разновидности задачи коммивояжера. На самом деле процедура сложнее, там есть еще куча условий и используются еще два массива с данными, но она не работает даже в представленном выше примитивном варианте и это вызывает у меня отчаяние.
Я попытаюсь трассировать что происходит со значением t у каждого потока на каждой итерации, хоть это и очень муторно. Может что пойму.

Мне уже стало интересно. Если

Мне уже стало интересно.
Если машина многоядерная - попробуйте emulation mode, она соберет мультитредную программу исполняемую на CPU и если есть проблемы с синхронизацией вы их увидите в обычном отладчике.

Извините за занудство - в x - все значения строго в диапазоне 0..32?

В-принципе, если есть желание, пришлите мне тестовую задачу и тестовые данные (лучше - с теми, на которые вы знаете правильные ответы) перед выходными (лучше всего в пятницу), я на выходных могу поковыряться с этим. До выходных - не могу. Писать на lexa@lexa.ru

Лично я бы начал с того, что завел бы еще один выходной массив (p[]) такого же размера что и x
и на каждом шаге печатал бы в него маршрут (в p[idx*33+i-1]) - и сравнил бы с предполагаемым маршрутом.

Т.е. что-то у вас не так, это понятно, но на расстоянии понять трудно.

Ура!!!!! Я разобрался. Ошибка

Ура!!!!! Я разобрался. Ошибка была настолько тупая, что и рассказывать стыдно, CUDA никак не касающаяся.
Не заметил ускорения выполнения алгоритма. Видеокарта у меня правда не супер - 9600GT, скорость получается примерно такая же как на CPU (на одном ядре C2D E7300). Однако тема меня заинтересовала и появилось много дополнительных вопросов к опытным.
Как правильнее, создать новую тему , а эту попросить модераторов грохнуть, или писать прямо здесь?

Меня интересует как реализовать генератор случайных чисел и по памяти есть вопросы. Сейчас у меня все данные просто в device memory, для тестовой задачи я мог бы их сделать shared, но для реальной они просто туда не поместятся (400*400*sizeof(float)+400*2*sizeof(float)=67200).

Тему грохать не надо, пусть

Тему грохать не надо, пусть нарастает содержание, просто заведите новую.

Давайте сначала к исходной задаче вернемся: у вас на каждом шаге цикла идет случайное обращение к массиву в глобальной памяти. Это - медленно, удивительно что у вас оно работает сравнимо с одним ядром CPU.
Варианты ускорения для задачи путей
- таблицу стоимостей (если она небольшого размера) - поместить в constant memory (быструю)
- таблицу путей перед началом работы читать или в shared memory или в регистры.

Кроме того, вы используете целую 32-битную арифметику. Чудо в том, что 24-битные целые - гораздо быстрее, а для ваших задач их достаточно. И еще одно возможное (тут надо проверять) ускорение - у вас все переходы в диапазоне 0-32, 5 бит для них недостаточно, а вот 6-ти - вполне. Т.е. массив x можно здорово упаковать, например хранить в одном элементе (32-битном) сразу 5 элементов пути. Я не гарантирую, что битовые операции помогут, но пробовать точно надо.

А по второму вопросу - если у вас какая-то часть из 67 килобайт константная - можно в constant memory класть.

Не знаю правильно ли я делаю,

Не знаю правильно ли я делаю, у меня каждый поток перемещает свою порцию таблицы путей в shared перед тем как начать чтото делать:

  1. __shared__ int sx[128][33];
  2. ...
  3. for (i=1;i<=33;i++)
  4.       {
  5.         sx[idx][i]=x[idx*33+i];
  6.       }

Ускорение после после того как я так сделал невооруженным глазом не видно.
Сейчас как раз читаю про constant memory, все данные (таблицу стоимостей и еще две таблички) можно грузить туда, они не меняются в течение всего выполнения.
И еще о задаче. Алгоритм состоит из четырех основных этапов:
- оценки экземпляров популяции (она уже на GPU)
- выбора лучших экземпляров (тоже уже на GPU)
- размножения этих экземпляров (пока на CPU)
- модификации размноженных копий (тоже на CPU)
Первые два этапа занимают порядка 50% времени выполнения (если на CPU, на GPU не мерил), последние два - соответственно тоже 50%. Сейчас получается так, что я засылаю на GPU массив с популяцией на каждой итерации, а назад на каждой же итерации получаю массив нулей и единиц, которые показывают, какие экземпляры размножать, а какие нет. Думаю это очень тормозит работу, мне бы хотелось и оставшиеся два этапа выполнять на GPU и не кидать данные туда-сюда, скорее всего только так я получу ощутимый прирост в скорости.
Для этого мне нужно придумать, каким образом экземпляры могли бы себя копировать (раз уж они в shared памяти) так чтобы не мешать друг другу (сейчас каждый просто записывает несколько своих копий в следующие за ним порции ), и реализовать генератор случайных чисел для модификации размноженных копий.
Кстати, я правильно понимаю, что после завершения кернела shared память сама собой очищается?
Спасибо за 24битные целые, я про них забыл.

33 - _очень_ плохое число, у

33 - _очень_ плохое число, у вас будут конфликты и в shared и в global. 32 - никак нельзя?

shared - очищается после выполнения блока. Но вы можете результаты сохранять в global (но надо аккуратно, чтобы тоже конфликтов не было) и запускать разные kernels с этой global

К сожалению это число может

К сожалению это число может быть самым разным, не обязательно кратным 16ти или 32ум.

Не могу понять как работать с константами. В учебнике пример есть как они объявляются и как заполняются CopyToSymbol. А обращаться то с ними как? Вечером в примерах SDK посмотрю.

Заметил такую вещь: бывает так что все компилируется в обычном режиме, запускается, но не работает нормально, а в режиме эмуляции выдает ошибки.

И еще вопрос. Как же все таки получать случайные числа? Сейчас я их получаю на CPU и передаю на GPU, это както криво..

Гонять данные с хоста на двс

Гонять данные с хоста на двс действительно не стоит, лучше генерировать сразу на гпу. Алгоритмов много, просто напишите подпрограмму под конкретный генератор, один из примеров реализации можно наити здесь: http://arxiv.org/abs/0903.3053, правда это один из "тяжелых" генераторов. Более простой и быстрый генератор исползуют, например, здесь:
Tobias Preis, Peter Virnau, Wolfgang Paul, Johannes J. Schneider, "GPU accelerated Monte Carlo simulation of the 2D and 3D Ising model", Journal of Computational Physics Volume 228, Issue 12, 1 July 2009, Pages 4468-4477,
(Если нет доступа к источнику, могу выслать електронную версию статьи)
Насколько конкретный генетор подходит для параллельных вычислений - это отдельная тема..

Я понял следующее: не так то

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

Быстрее будет работать при

Быстрее будет работать при соблюдении двух условий
- относительно много вычислений (в расчете на один байт данных)
- алгоритм параллелится без зависимостией между разными тредами.

И снова здравствуйте

Я переписал полную версию своего алгоритма на CUDA. Производительность - ну просто никакая. Не могу понять в чем дело, неужели 9600GT настолько медленный? Я уже все что можно загнал в constant, shared и просто регистры, убрал почти все ветвления, данные туда сюда почти не гоняю, все int'ы и местами даже float'ы поменял на short'ы, пробовал разные размеры блоков и все равно CPU делает GPU. Понятное дело, на тех данных, на которых я проверяю, все они помещаются в кэш второго уровня, но собственно так обстоят дела на всех практичных размерностях задачи, данные помещаются в пару мегабайт.
Почему я продолжаю пытаться добиться от GPU производительности хотя бы сравнимой с CPU: один китайский аспирант реализовал похожий алгоритм, и он работает по его словам шустро, намного быстрее чем на ЦПУ. Правда у него Tesla, но неужели Tesla настолько быстрее? К сожалению китайский товарищ отказался показать листинг, но на пальцах у него все то же самое что у меня.
А что у меня: GPU используется для вычисления стоимости пути в графе, для этого используется матрица стоимостей (в constant), и еще два вектора, значения из которых либо добавляются к стоимости, либо вычитаются из нее при прохождении каждого узла (оба тоже в constant). причем размер их таков, что они скорее всего полностью находятся в текстурном кэше, кроме того там есть штрафы за циклы, их стоимость вычисляется. На каждой итерации алгоритма на GPU передается пачка маршрутов, каждый из которых представляет собой одномерный массив (short), а обратно забирается пачка стоимостей (каждая float), перед вычислениями каждый поток читает свой маршрут в shared (не уверен что это имеет смысл кстати, потому что каждый элемент маршрута используется один раз, но я понадеялся что копирование в shared сразу кучки данных и последующее использование оттуда, происходит быстрее чем одиночные обращения к глобальной памяти). Около 80% времени на обычном процессоре (если не отображать результаты) как раз занимает вычисление стоимостей, т.е. теоретически, если GPU выполняет его вдвое быстрее ЦПУ, то и весь алгоритм должен выполняться почти вдвое быстрее, но этого не происходит, на самом деле он выполняется в два-три раза медленнее. В чем же дело?

У меня на ноутбуке 8600M Для

У меня на ноутбуке 8600M

Для подходящих ей задач - она примерно такая же быстрая, как CPU ноутбука (core2 duo 2.2 Ghz)

Как это понять? Т.е. не стоит

Как это понять? Т.е. не стоит ожидать от 9600 ураганных скоростей? Или наоборот у меня чтото криво, либо с алгоритмом, либо с самой задачей?

Кто нибудь делал чтонить связанное с обходом графов?

Десктопная 9600GT должна быть

Десктопная 9600GT должна быть в разы быстрее ноутбучной 8600M (плюс там память быстрее).

Т.е. если у вас на CPU используется SSE в полный рост, то 9600GT с хорошо приспособленной под нее задачей должна быть примерно как современный 4-ядерник с использованием SSE. Ну, плюс-минус.

Если CPU используется скалярно (без SSE), то оптимальная задача на 9600 должна выполняться в разы быстрее.

Спасибо! Сегодня я занимался

Спасибо!

Сегодня я занимался замерами времени выполнения. Оказалось что одна итерация у меня выполняется 15-16 миллисекунд (т.е. можно считать один кернел на GPU, потому что как выяснилось все memcpy выполняются меньше миллисекунды). Тогда я стал отключать участки кода в кернеле, чтобы узнать что же занимает все это время. Оказалось, что это цикл for, который выполняет большую часть действий. Как видно в примере цикл там только для того чтобы одно и то же действие выполнялось несколько раз. Я закомментил почти все тело, ничего не изменилось, тогда я его развернул, и (о чудо!) 15-16 миллисекунд стал выполняться только каждый 16тый кернел (при размере блока = 1), причем иногда, раз в две-три секунды, выпадают кернелы, которые выполняются дольше - около 100 миллисекунд (могу предположить что это случаи непопадания в текстурный кэш), остальные стали выполняться менее чем за миллисекунду (т.е. скорость выросла в разы). Чем может объясняться этот феномен? В кернеле есть другой цикл for, который выполняется такое же количество раз (загружает из глобальной памяти в shared значения) и он так не тормозит. Это какаято фишка компилятора может быть?

Если вас интересует

Если вас интересует производительность, то

1) Размер блока не может быть 1 (thread). Оптимально обычно что-то в районе 64-128-256, помните пожалуйста, что у вас выполняется целый warp (32 threads) на одном процессоре (коих на 9600 - 4 штуки, кажется).
Ну и блоков может быть много.

Не про GPU:

2) Чем вы меряете время? устойчивые 15-16 миллисекунд наводят меня на мысли о гранулярности таймера или планировщика.

3) У вас не Linux, случайно (на Linux оный эффект я наблюдал тут в одном проекте - мы там гарантировали некую скорость исполнения, она не соблюдалась на очень маленьких заданиях, а проблема была именно в точности измерения, а не в скорости исполнения). Правда там гранулярность таймера (точнее, планировщика) была 18 миллисекунд, а не 15-16.

Спасибо еще раз! 1) Т.е. вот

Спасибо еще раз!

1) Т.е. вот это неправильно?

  1. int block_size = 1;
  2. int n_blocks = 128/block_size + (128%block_size == 0 ? 0:1);
  3. gputest <<< n_blocks, block_size >>> (X_d, Y_d, 128);

Но я реально замечаю, что при уменьшении block_size ускоряется выполнение, изначально я его делал равным 64 (по числу потоковых процессоров в 9600, мне это показалось логичным). Может это признак чего то? У меня есть параноидальная мысль, что на самом деле прога не выполняется на GPU. Может такое быть?

2) Время меряю _ftime( ) (может это неправильно, но другого способа я не знаю), вызываю ее перед и после выполнения кернела и сохраняю значения миллисекунд с начала секунды. Кстати значения "около 100", о которых я говорил выше, на самом деле не около ста. Я просто на экран их выводил и не успевал рассмотреть. Сейчас вывел в файл и на самом деле они около -1000 :) Это происходит когда одна _ftime( ) вызывается в течение одной секунды, а вторая уже в другой. :)

3) Нет, у меня Windows XP. Я думаю что про гранулярность это верно. Что я не закоменчивал, время никогда не принимало промежуточных значений между 0 и 15-16.

Давайте начнем с конца. 3)

Давайте начнем с конца.
3) Как у XP устроен планировщик - не знаю. Но эффект с гранулярностью - есть.
2) В примерах CUDA есть использование таймера (cutCreateTimer), который в тулките свой, он хороший.
Ну или QueryPerformanceCounter (http://msdn.microsoft.com/en-us/library/ms644904%28VS.85%29.aspx) - у него просто разрешение - такт процессора.

1) Давайте я вам документацию не буду пересказывать. Возьмите Best Practices Guide, раздел 4.4, но один поток на блок - это ~3% от пиковой производительности (если все остальное идеально и не конфликтует)

Да, Tesla (новая) - это

Да, Tesla (новая) - это 280GTX примерно. Т.е. 224 процессора. На 9600 - их 64.