Решил попытаться ускорить одну свою программку при помощи CUDA, столкнулся с глюком. Процедура, приведенная ниже какбэ должна считать стоимость пути в графе:
- __global__ void gputest(int *x, float *y, int *z, float *c, float *d, float *p, int N)
- {
- int idx = blockIdx.x * blockDim.x + threadIdx.x;
- int i,ths,next;
- float t=0;
- if (idx<N)
- {
- ths=0;
- for (i=1;i<=33;i++)
- {
- next=x[idx*33+ths];//x - это замкнутые маршруты, которые состоят ровно из 33х шагов
- t=t+c[ths*33+next];//с - это стоимости переходов от одного узла графа к другому
- ths=next;
- }
- y[idx]=t;//y - это стоимости маршрутов
- }
- }
Такая же процедура на обычном CPU, вызываемая в цикле, прекрасно забивает массив y стоимостями, на GPU она забивает его значениями -431602080, причем, если сократить количество шагов цикла for до 20ти, то получаются в некоторых элементах правильные стоимости, в других значения -35659499650496332000, в третьих #QNAN0. Может кто-нибудь знает почему так происходит?
Мне 20 шагов мало, в реальной задаче их вообще сотни, а маршрутов тысячи. Помогите пожалуйста.
PS вверху в приведенном коде должен бы стоять знак меньше перед N, но в таком случае движок форума отсекает весь текст который после него (тоже своего рода глюк :))
Comments
Я переписал полную версию своего алгоритма на CUDA. Производительность - ну просто никакая. Не могу понять в чем дело, неужели 9600GT настолько медленный? Я уже все что можно загнал в constant, shared и просто регистры, убрал почти все ветвления, данные туда сюда почти не гоняю, все int'ы и местами даже float'ы поменял на short'ы, пробовал разные размеры блоков и все равно CPU делает GPU. Понятное дело, на тех данных, на которых я проверяю, все они помещаются в кэш второго уровня, но собственно так обстоят дела на всех практичных размерностях задачи, данные помещаются в пару мегабайт.
Почему я продолжаю пытаться добиться от GPU производительности хотя бы сравнимой с CPU: один китайский аспирант реализовал похожий алгоритм, и он работает по его словам шустро, намного быстрее чем на ЦПУ. Правда у него Tesla, но неужели Tesla настолько быстрее? К сожалению китайский товарищ отказался показать листинг, но на пальцах у него все то же самое что у меня.
А что у меня: GPU используется для вычисления стоимости пути в графе, для этого используется матрица стоимостей (в constant), и еще два вектора, значения из которых либо добавляются к стоимости, либо вычитаются из нее при прохождении каждого узла (оба тоже в constant). причем размер их таков, что они скорее всего полностью находятся в текстурном кэше, кроме того там есть штрафы за циклы, их стоимость вычисляется. На каждой итерации алгоритма на GPU передается пачка маршрутов, каждый из которых представляет собой одномерный массив (short), а обратно забирается пачка стоимостей (каждая float), перед вычислениями каждый поток читает свой маршрут в shared (не уверен что это имеет смысл кстати, потому что каждый элемент маршрута используется один раз, но я понадеялся что копирование в shared сразу кучки данных и последующее использование оттуда, происходит быстрее чем одиночные обращения к глобальной памяти). Около 80% времени на обычном процессоре (если не отображать результаты) как раз занимает вычисление стоимостей, т.е. теоретически, если GPU выполняет его вдвое быстрее ЦПУ, то и весь алгоритм должен выполняться почти вдвое быстрее, но этого не происходит, на самом деле он выполняется в два-три раза медленнее. В чем же дело?
Да, Tesla (новая) - это 280GTX примерно. Т.е. 224 процессора. На 9600 - их 64.
У меня на ноутбуке 8600M
Для подходящих ей задач - она примерно такая же быстрая, как CPU ноутбука (core2 duo 2.2 Ghz)
Как это понять? Т.е. не стоит ожидать от 9600 ураганных скоростей? Или наоборот у меня чтото криво, либо с алгоритмом, либо с самой задачей?
Кто нибудь делал чтонить связанное с обходом графов?
Десктопная 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) Т.е. вот это неправильно?
Но я реально замечаю, что при уменьшении block_size ускоряется выполнение, изначально я его делал равным 64 (по числу потоковых процессоров в 9600, мне это показалось логичным). Может это признак чего то? У меня есть параноидальная мысль, что на самом деле прога не выполняется на GPU. Может такое быть?
2) Время меряю _ftime( ) (может это неправильно, но другого способа я не знаю), вызываю ее перед и после выполнения кернела и сохраняю значения миллисекунд с начала секунды. Кстати значения "около 100", о которых я говорил выше, на самом деле не около ста. Я просто на экран их выводил и не успевал рассмотреть. Сейчас вывел в файл и на самом деле они около -1000 :) Это происходит когда одна _ftime( ) вызывается в течение одной секунды, а вторая уже в другой. :)
3) Нет, у меня Windows XP. Я думаю что про гранулярность это верно. Что я не закоменчивал, время никогда не принимало промежуточных значений между 0 и 15-16.
Давайте начнем с конца.
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% от пиковой производительности (если все остальное идеально и не конфликтует)
Я понял следующее: не так то и просто перенести на GPU алгоритм локального поиска. Точнее перенести то можно, но: много проблем, связанных с организацией памяти, не факт что он будет работать быстрее, результаты несколько отличаются от получаемых на CPU (наверное это связано с тем как хранятся данные на GPU).
Я две недели парился с этим, результат меня не удовлетворяет, но я рад тому, что научился программировать GPU, возможно пригодится в будущем.
Быстрее будет работать при соблюдении двух условий
- относительно много вычислений (в расчете на один байт данных)
- алгоритм параллелится без зависимостией между разными тредами.
Ура!!!!! Я разобрался. Ошибка была настолько тупая, что и рассказывать стыдно, 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 перед тем как начать чтото делать:
Ускорение после после того как я так сделал невооруженным глазом не видно.
Сейчас как раз читаю про constant memory, все данные (таблицу стоимостей и еще две таблички) можно грузить туда, они не меняются в течение всего выполнения.
И еще о задаче. Алгоритм состоит из четырех основных этапов:
- оценки экземпляров популяции (она уже на GPU)
- выбора лучших экземпляров (тоже уже на GPU)
- размножения этих экземпляров (пока на CPU)
- модификации размноженных копий (тоже на CPU)
Первые два этапа занимают порядка 50% времени выполнения (если на CPU, на GPU не мерил), последние два - соответственно тоже 50%. Сейчас получается так, что я засылаю на GPU массив с популяцией на каждой итерации, а назад на каждой же итерации получаю массив нулей и единиц, которые показывают, какие экземпляры размножать, а какие нет. Думаю это очень тормозит работу, мне бы хотелось и оставшиеся два этапа выполнять на GPU и не кидать данные туда-сюда, скорее всего только так я получу ощутимый прирост в скорости.
Для этого мне нужно придумать, каким образом экземпляры могли бы себя копировать (раз уж они в shared памяти) так чтобы не мешать друг другу (сейчас каждый просто записывает несколько своих копий в следующие за ним порции ), и реализовать генератор случайных чисел для модификации размноженных копий.
Кстати, я правильно понимаю, что после завершения кернела shared память сама собой очищается?
Спасибо за 24битные целые, я про них забыл.
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,
(Если нет доступа к источнику, могу выслать електронную версию статьи)
Насколько конкретный генетор подходит для параллельных вычислений - это отдельная тема..
Скажите, а эти массивы 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]) - и сравнил бы с предполагаемым маршрутом.
Т.е. что-то у вас не так, это понятно, но на расстоянии понять трудно.
1) код обрамлять тегами
<code>2) Вопрос: не могу понять а зачем тут цикл for если в его теле не используется переменная i ?
Затем чтобы повторить то что в теле много раз. Как я уже написал выше, "много", в реальных задачах, измеряется сотнями. А если удастся увеличить производительность, с помощью куды, или еще как, то может это будут и тысячи.
По существу есть что сказать?