CUDA. Несколько вопросов и Решение СЛАУ методом Гаусса

Здравствуйте.
Начал недавно изучать CUDA. Возникло несколько вопросов (на мой взгляд нубских). И наверно в одну тему их помещать не надо было бы, но все же.
1) Про нити ( потоки). Количество их всегда должно равняться количеству обрабатываемых элементов? Т.е если у меня массив из 1000 элементов , то и потоков (нитей) я тоже создаю 1000. Независимо от задачи, количество нитей = количеству обрабатываемых элементов? 1 элемент обработки = 1 нить.?

2) Возникло желание попробовать написать на CUDA решение СЛАУ методом Гаусса, но потом подумал и решил, что наверно это сделать нельзя потому что под блочно - нитевидную архитектуру она не ложится.
Проблема в том, что когда мы избавляемся от первого столбца , по алгоритму, переход на следующий столбец будет произведен тогда когда закончится обработка первого столбца, вот как сделать эти задержки я не очень понимаю. (да и блок у нас какой то странный получится)

Если я ошибаюсь, подскажите возможно ли такое реализовать? есть ли уже готовые решения буду благодарен, может в той же CUBLAS есть.

Заранее спасибо.

Forums: 

>1) ... это уже целиком и

>1) ...
это уже целиком и полностью зависит не от задачи, а от конкретной реализации, т. .е. от вашей фантазии.
как сделаете, так и будет работать.
хотите 1 поток на элемент - делайте, хотите 1 поток на 10 элементов - да пожалуйста.

>2) ...
не знаю поможет ли вам это или нет, но в CUDA есть так называемая барьерная синхронизация (метод __syncthreads()) - почитайте, может натолкнет на новые мысли.

Про __syncthreads() , я знаю,

Про __syncthreads() , я знаю, но как она мне может помочь понятия не имею в этой задаче)

А вообще метод Гаусса для СЛАУ реализуем на CUDA? выполнимо такое?

Ребят, ну вы даете...Один не

Ребят, ну вы даете...
Один не помнит метод гаусса, второй не знает как его реализовывать...

Сразу на ум пришло такое:
1) создаем нитей сколько элеметнов в строке или в N раз меньше
2) блоков в M раз меньше чем строк

Синхронизируем нити.
А вот блоки либо через глобалку, либо через запуск функции с хоста.
Если их вообще надо синхронизировать...

Вот только один момент есть, погрешность флоата 7ой знак, считать матрицы меньше чем 1000*1000 смысла нет.
получим погрешность около 2-3 знака после запятой... тк там неск операций, оно тебе надо?
Особенно на фоне того что значения могут отличатся на два порядка)))

"Сразу на ум пришло

  1. "Сразу на ум пришло такое"

Возьми с полки пирожок.

  1. "1) создаем нитей сколько элеметнов в строке или в N раз меньше"

Чтобы отнять одну строку от другой, нужно знать коэффициент. Как этот коэффициент разбрасывать на потоки(особенно интересует как это делать между блоками)? Насколько такой разброс будет тормозить?
С каждой итерацией в прямом ходе, в строках появляются нолики. Что делать потокам у которых нолики? Молотить эти нули? Или тупо простаивать?
Как синхронизировать момент перехода к следующей итерации?

  1. "А вот блоки либо через глобалку, либо через запуск функции с хоста."

Насколько всё это будет тормозить?

  1. "2) блоков в M раз меньше чем строк"

То есть опять теже проблемы - синхронизация между итерациями,
количество строк над которыми нужно работать уменьшаются с каждой итерацией - простой?

Учитывая всё сказанное, зачем человеку который с GPGPU не имел дела, начинать с алгоритма в котором много параллельных граблей?

  1. "Особенно на фоне того что значения могут отличатся на два порядка)))"

То что метод гаусса можно реализовать на GPGPU это сомнений не вызывает.
Но, я не уверен насколько эффективной будет эта реализация.
Топикстартер так и не сказал какого типа у него задача. Разряженная матрица или нет. какой разброс по значениям и т.п.
А вы выдаёте какое-то решение, без чётко поставленной задачи.
Разные задачи требуют решения СЛАУ, но СЛАУ у каждого типа задач имеет свою специфику, в соответствии с этим необходимо выбирать более подходящий алгоритм.

Поэтому не стоит с голыми руками кидаться на мамонта

Но, я не уверен насколько

Но, я не уверен насколько эффективной будет эта реализация

Нормальной такой: (и кроме Волкова, повторяю, есть еще MAGMA)

Где здесь метод Гаусса(в

Где здесь метод Гаусса(в чистом виде, о котором говорил Opex)?

PS: Я не спорю о том, что есть прямые методы решения СЛАУ, которые могут быть эффективно реализованы на GPGPU.

Точнее, наоборот, конечно,

Точнее, наоборот, конечно, один из способов - по Гауссу.

Но судя по исходному вопросу /оговорке про CUBLAS/ - важен результат (прямое решение СЛАУ через треугольную матрицу), а не сам процесс.

"важен результат" Может быть.

"важен результат"
Может быть. В этом случае всё равно нужно указать тип задачи. Может результат это не "прямое решение СЛАУ через треугольную матрицу", а просто решение заданного СЛАУ, и может быть что матрица разряжена. Может важно именно найти обратную матрицу.. и т.п.

Вопрос об возможности

Вопрос об возможности эффективной реализации метода гаусса(в чистом виде) на GPGPU остаётся открытым.
Кэпство: не факт что для LU разложения там использовался метод Гаусса.

Может быть будет эффективно

Может быть будет эффективно рассматривать матрицу как блочную.
И соответственно вместо деления на число, нужно будет обращать под-матрицу, вместо умножения двух чисел нужно умножать две под-матрицы.
Вот только насколько я помню, например эффективное умножение матриц на GPU начинается где-то с N ~ 2k , соответственно исходную матрицу придётся делить на блоки по N~2k..

Но имхо, это уже точно не то, что имелось ввиду топикстартером.

Про пирожек, конечн,

Про пирожок, конечно, остроумно, но не надо показывать свое остроумие здесь...
Как я понял, топикстартеру эта задача интересна с точки зрения обучения.
Тк иначе рассматривал другие методы решения СЛАУ.
Учитывал бы погрешности, разряжение, размер матриц и тд.

Про коеффициент для вычитания строк:
посчитали коеф, далее его используем всеми нитями.
Когда появились нолики, нити простаивают.
Тормоза при вызове функции с хоста, простой порядка 100мк сек на вызов.

То что количество строк для обработки уменьшается и блоки будут простаивать - верно для ситуации когда блоков больше строк. Учитывая то, что матрица более 1000*1000, а блоков не более 240 ( для GT200, тк болье смысла делать нет ) это не сильно снизит производительность. Тк нитей 512 на блок , те параллельно решаем не более 120 строк, падение скорости только на 1/8 задачи.

Кроме вопроса - "насколько это будет тормозить", есть еще вопрос "надо ли ЭТО делать гауссом".

PS
"А вы выдаёте какое-то решение, без чётко поставленной задачи."
А Вы только умничали и не дали никакого совета.

Про пирожок, конечно,

  1. Про пирожок, конечно, остроумно, но не надо показывать свое остроумие здесь...
  1. Ребят, ну вы даете...
  2. Один не помнит метод гаусса, второй не знает как его реализовывать...

Прежде всего, не думать что лучше, умнее других. Даже если кто-то не знает/не помнит метод Гаусса, не нужно открыто показывать надменное удивление, тогда в ответ не будете получать отсылы к пирожкам.

  1. Про коеффициент для вычитания строк:
  2. посчитали коеф, далее его используем всеми нитями.
  3. Когда появились нолики, нити простаивают.

Как используем? Я не вижу ответа на вопрос "Как этот коэффициент разбрасывать на потоки(особенно интересует как это делать между блоками)? ", а потом уже на "Насколько такой разброс будет тормозить?"

  1. Тормоза при вызове функции с хоста, простой порядка 100мк сек на вызов.

"порядка" это понятие раз в 10 растяжимое.
Даже если брать 100мкс: допустим система из 10k уравнений (можно и больше ). Сколько по-вашему нужно сделать запусков ядра? 10000, 50000000 или сколько? И сколько времени это займёт суммарно?

  1. То что количество строк для обработки уменьшается и блоки будут простаивать - верно для ситуации когда блоков больше строк.
  2. Учитывая то, что матрица более 1000*1000, а блоков не более 240 ( для GT200, тк болье смысла делать нет ) это не сильно снизит производительность.
  3. Тк нитей 512 на блок , те параллельно решаем не более 120 строк, падение скорости только на 1/8 задачи.

Если вы действительно хотели вложить в этот абзац какие-то внятные мысли, просьба перефразировать. А так, что-то несвязное и с индуским акцентом.

  1. Кроме вопроса - "насколько это будет тормозить", есть еще вопрос "надо ли ЭТО делать гауссом".

Браво, вы всё-таки прочитали моё сообщение, даже что-то поняли.

  1. "А вы выдаёте какое-то решение, без чётко поставленной задачи."
  2. А Вы только умничали и не дали никакого совета.

В этой стране удивительные люди, решают задачи без условия - http://twitter.lexa.ru/2010/04/02/chem_seichas_zanimayutsya_perl_program... .

Неужели никакого?
http://www.gpgpu.ru/node/238#comment-1702
Также я пытался выяснить, какая всё-таки задача, если получу ответ, то дам ещё ссылок. У меня всё-таки есть небольшой опыт решения СЛАУ на GPU, также я искал информацию на эту тему.

Удивление насчет метода

Удивление насчет метода гаусса вполне оправдано, это самый очивидный метод решения СЛАУ, его не помнить в подробностях, сложно.

Если вам не ясны мои выкладки насчет того простоя блоков на < 1/8 задачи, могу посочувствовать...

Порядка 100мкс - да, не буду утверждать что не может быть 10мкс или 1 мс.
Но то что не более 1 мс уверен.

Ваше ехидство говорит только о вашей глупоти :)

Удивление насчет метода

Удивление насчет метода гаусса вполне оправдано, это самый очивидный метод решения СЛАУ, его не помнить в подробностях, сложно.
Да очевидный, но не каждый обязан его знать. Но тем не менее ваши реплики не тактичны
Ребят, ну вы даете... Один не помнит метод гаусса, второй не знает как его
Зачем переходить на личности?
Сразу на ум пришло такое
Причём предложены самые наивные решения, детали реализаций которых я от вас так и не получил.
Конкретно, вы не дали ответа о том как разбрасывать коэффициенты, и сколько по вашему необходимо вызовов ядра с хоста(для обеих реализаций). Вы эти вопросы проигнорировали.
Если вам не ясны мои выкладки насчет того простоя блоков на < 1/8 задачи, могу посочувствовать...
Так, по порядку:
То что количество строк для обработки уменьшается и блоки будут простаивать - верно для ситуации когда блоков больше строк.
Я так понял, по обеим предложенным схемам, количество блоков(вообще-то изначально разговор шёл про потоки) будет меньше либо равно числу строк. То есть по-вашему "количество строк для обработки уменьшается" неверно никогда? Вы будите всё время молотить все строки?
Учитывая то, что матрица более 1000*1000, а блоков не более 240 ( для GT200, тк болье смысла делать нет ) это не сильно снизит производительность
Кто сказал, что больше делать смысла не имеет? Пруфы, тесты? Что не сильно снизит производительность? То что часть потоков будет простаивать, или что?

Тк нитей 512 на блок , те параллельно решаем не более 120 строк, падение скорости только на 1/8 задачи.

Вот это уже отменная идусятина, то есть есть 512 потоков умножить на 240 блоков, а вы хотите "решать" параллельно 120 строк, пусть даже не умножая на 240.. Что это за метод? комбинация первых двух предложенных? То есть вы придумали какой-то метод и описываете какие-то его свойства, без объяснения метода?
Ваше ехидство говорит только о вашей глупоти :)
Я больше не желаю с вами общаться.

Тогда не имеет смысла

Тогда не имеет смысла объяснять суть предложенного мной метода.
Тк кроме вас, никто не говорил, что, не ясно спросят - объясню.

В краце идею озвучил.

1) создаем нитей сколько

1) создаем нитей сколько элеметнов в строке или в N раз меньше
2) блоков в M раз меньше чем строк

Насколько я понял это два варианта предложенных вами методов.

Теперь немного предыстории.
Около 5 месяцев назад, один человек попросил меня помочь распараллелить/отладить метод Гаусса(не GPGPU, соответственно параллельно не так много потоков). Первые две идеи распараллеливания были - разделить строки между потоками, либо разделить элементы в строке между потоками, то есть те же самые наивные реализации предложенные вами.
Был реализован метод, когда шло разделение по не обработанным строкам. Программа писалась на Delphi(так нужно было) с использованием SSE через asm вставки, во время разработки кстати был найдет косяк в каком-то модуле Delphi (то ли связанным с потоками, то ли с синхронизацией), что подтвердилось гуглением. Алгоритм работает, прирост скорости был получен.
Но здесь мы как-бы говорим об GPGPU, где одновременно работают тысячи потоков, причём дёшево синхронизировать ВСЕ потоки нельзя, поэтому предолгая какой-либо метод, тем более новичку, нужно продумывать все возможные аспекты реализации на GPGPU, что вы не сделали. Не достаточно выделить только параллельность данных/потоков вычислений, для хорошей GPGPU реализации.

Зачем я всё это пишу, чтобы у вас не было иллюзии о том, что предложенные методы оригинальны, и никому из участников дискуссии раньше не приходили в голову. Как сказал уважаемый lexa, чтобы померятся эго, тем более что вы изначально вели себя крайне некультурно по отношению к форумчанам - "Live by the sword, die by the sword".

  1. Тогда не имеет смысла объяснять суть предложенного мной метода.
  2. Тк кроме вас, никто не говорил, что, не ясно спросят - объясню.

Да, не надо, курить вашу траву я не собираюсь.

PS: Если кто-то найдёт примеры эффективных реализаций метода Гаусса на GPGPU, с интересом посмотрю.

Если кто-то найдёт примеры

Если кто-то найдёт примеры эффективных реализаций метода Гаусса на GPGPU, с интересом посмотрю.
Я вчера вечером погуглил (правда больше на ACM.org) на эту тему минут 10.
Причем, даже не столько про GPGPU, сколько про кластеры (там должно быть больше наработано, а суть примерно одна).

Реализаций не смотрел, смотрел всякие презентации про них. Не очень убедительные, как мне показалось (но я не очень специалист). Но по сути все сводится к множественным вызовам BLAS2 или аналогов - по строкам и все такое.

Предварительный вывод мой таков - действительно эффективных реализаций нет т.к. не нужно, на параллельных системах блочные методы LU гораздо интереснее.

Господа, постарайтесь меньше

Господа, постарайтесь меньше переходить на личности. Тут можно меряться эго, а вот ругаться - не стоит

Да реализуем, почему

Да реализуем, почему нет?
Запускаете в одном потоке весь алгоритм. Какой вопрос, такой и ответ.

Вообще, сначала определитесь, какого типа метод вам нужен прямой(типа Гаусса), итерационный, или конкретно обычный Гауссовский метод?
У прямых и итерационных методов есть свои плюсы и минусы, и соответственно разная область применения. Быть может для вашей задачи подойдут больше итерационные методы.
Или задача это реализовать метод Гаусса на Cuda?

Имхо итерационные методы легче реализуются на gpgpu, дают больший прирост по отношению к CPU(так как требуют меньше памяти в ходе решения, могут поместиться полностью в память GPU, за счёт этого меньше копий между хостом и видеокартой), легче распараллелить на кластер видеокарт - коммуникация минимальна(пару вещественных чисел на одну итерацию).

Вот статья про прямой решатель:
http://saahpc.ncsa.illinois.edu/09/papers/Krawezik_paper.pdf
Вот про умножение разряженной матрицы на вектор(самая ресурсоёмкая операция в итерационных решателях):
http://www.nvidia.com/docs/IO/66889/nvr-2008-004.pdf
http://www.nvidia.com/docs/IO/77944/sc09-spmv-throughput.pdf

J0hn спасибо, почитаю. Я имел

J0hn спасибо, почитаю. Я имел ввиду обычный метод Гаусса (самый стандартный). Просто я пока мало представляю как эту задачу можно разложить на блоки и потоки =)
Спасибо за статью еще раз.