Посоветуйте что выбрать OpenCL/Cuda/Stream

Я собираюсь использовать gpgpu в своей дипломной работе.
В основном будут решаться большие СЛАУ, желательно с double точностью.. но СЛАУ я не собираюсь ограничиваться..

Я ещё не решил какое API я буду использовать OpenCL/Cuda/Stream.
Сейчас у меня есть ATI HD4850, но я планирую взять новую видеокарту, максимальная стоимость 15000 т.р.
Посоветуйте какая из видеокарт больше подходит для каждого из трёх API (конкретные модели), а также на каком API мне остановиться.
Или хотя бы выскажете своё мнение, буду очень признателен.
(Материнка asus P5Q3. если понадобиться купить другую - не вопрос)
Приоритетом выбора является кол-во GFlops(двойной точности) и обьём памяти.

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

О конечном выборе и о результатах своей работы я отпишусь.

Forums: 

Да, я видел - я посмотрел

Да, я видел - я посмотрел код. Я имел ввиду в общем.
Я когда писал, думал как раз про тот вариант, при котором для достижения большей скорости, не достаточно переустановить драйвер., а нужно ещё перекомпилировать код - то есть лишние апдейты для конечных пользователей cuda приложений.

Я вот с трудом могу себе

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

Я не разбирался с тем, что

Я не разбирался с тем, что именно из себя представляет PTX код, но тем не менее, почему его компилятор(в PTX) не может быть оптимизирован? Может конечно PTX код это просто просто разпарсеный код ядра - то в этом случае врядли что-то можно оптимизировать.
Можно сделать тест - использовать PTX компиляторы(в PTX) разных версий(целевую архитектуру поставить одну и ту же, если возможно) и сравнить полученный код.

PTX - это пседоассемблер,

PTX - это пседоассемблер, вроде OpenGL-ного. .cubin - оно же, но бинарное (судя по результатам работы decuda). Не оптимизированный по использованию регистров, насколько я вижу.

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

Я скорее про то, что "при прочих равных" (одинаковой оптимизации) перекомпилировать из .cu в .ptx/cubin вроде бы не надо (с поправкой на поддерживаемые форматы cubin)

Так, мы уже далеко ушли от

Так, мы уже далеко ушли от темы.
Но тем не менее, давайте чётко определимся, что мы обсуждаем.
Мой вариант: Компилятор, который транслирует код ядра в PTX, может быть более оптимизирован в новых версиях тулкита(так ли это?), по-этому для получения этой оптимизации, ядра нужно перекомпилировать. То есть с точки зрения распространения такой оптимизации, OpenCL лучше - всё необходимое для компиляции есть в драйвере, и нет необходимости пересобирать приложения или их части(конечно у этого подхода есть свои минусы.. или например в OpenCL можно использовать и прекомпилированные ядра(clCreateProgramWithBinary)).

"Да, естественно, за счет лучшей оптимизации в более других версиях может стать получше."
"Я скорее про то, что "при прочих равных" (одинаковой оптимизации) перекомпилировать из .cu в .ptx/cubin вроде бы не надо"

Я правильно понимаю, что последним предложением вы хотели сказать следующие: если разные версии компилятора генерируют одинаково оптимальный код, то перекомпилировать не надо. ?
Понятное дело, что при "при прочих(всех) равных" нет необходимости использовать новое.

Я к PTX (и cubin) относился

Я к PTX (и cubin) относился именно как к псевдокоду. Такому, вроде OpenGL-овских шейдеров.
А драйвер из него делает что-то для конкретного железа (в частности, экономит регистры - если это нужно)

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

Важно именно то, что PTX - это еще не окончательный код.

У меня есть такие ядра как

У меня есть такие ядра как копирование из одного вектора в другой, умножение вектора на скаляр.
Ядра простые - один поток на одно значение. Время выполнения для векторов из float(32 бита) размером 2000000 для обоих ядер одинаковое(конечно есть небольшая погрешность с каждым запуском) - 0.000131 секунд.
Получается 2000000*4*2/(0.000131*2^30)=113.75GB/sec

GTX480 новые результаты

не мог заснуть, по-этому решил сделать тест на bandwidth dtod:
код ядер:

  1. #pragma OPENCL EXTENSION cl_khr_fp64 : enable
  2. #define AEQB(FUNCNAME,VECTORTYPE) \
  3. __kernel void FUNCNAME(__global VECTORTYPE *vA,__global VECTORTYPE *vB,const unsigned int n) \
  4. { \
  5.     uint tid = get_global_id(0); \
  6.     if(tid<n) \
  7.     { \
  8.         vA[tid]=vB[tid]; \
  9.     } \
  10. }
  11. AEQB(DAEqB,double)
  12. AEQB(SAEqB,float)

Размер группы 512 (определилось недолгим подбором, много времени строить графики нет..)
С увеличением N, увеличивалась и скорость передачи. Но сильно большое N не получиться взять, так как OpenCL вроде ограничивает размер одного буфера.

Время выполнения замерялось OpenCL'скими средствами.

  1. clGetEventProfilingInfo(events[0],CL_PROFILING_COMMAND_START,sizeof(cl_ulong),&startTime,0);
  2. clGetEventProfilingInfo(events[0],CL_PROFILING_COMMAND_END,sizeof(cl_ulong),&endTime,0);

У меня есть предположение, что в это время входит и коммуникация с хостом - передача ядра, параметров запуска и т.п., по-этому с увеличением N, увеличивался показатель bandwidth измеряемый мной.
Вот результаты

SAEqB N=134217728 : 124.4 GB/s
DAEqB N=67108864 : 136.9 GB/s

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

Сделал грубый замер - запустил ядро с N=1 - время С=0.000015с (NVidia выдаёт в CL_DEVICE_PROFILING_TIMER_RESOLUTION разрешение 1000ns),
запускал ядра с меньшим N:
SAEqB N=2097152
DAEqB N=1048576
показатель в нескольких вариантах - с вычетом C и без.

Вычет C Raw
SAEqB 125.5GB/s 112.1GB/s
DAEqB 137.9GB/s 122.2GB/s

Итого, результат с вычетом C достаточно близкий к запуску при большом N(без вычета).

Для перевода из байт в гигабайты делил на 2^30

Проверил в профилеровщике -

Проверил в профилеровщике - global load hit=0%
Но вот, что интересно: NVidia OpenCL сама разбивает глобальное количество блоков так, чтобы за один раз запукалось 65535, например для N=67108864 ядро запускалось три раза с количеством блоков 65535, 65535, 2 (размер блока 512).
Причём видимо запуск новой группы блоков каждый раз инициируется с хоста, а не с устройства, так как если я запускаю с таким N, что всё умещается в одну группу блоков(N=33553920), показатель немного больше:
SAEqB: 124.7Gb/s
DAEqB: 137.3Gb/s

Я предполагаю, что на задачах с большим количеством потоков выгоднее использовать двухмерные и трёхмерные блоки, а не давать драйверу запускать несколько раз ядро.

Кстати, в OpenCL 1.1 что-то

Кстати, в OpenCL 1.1 что-то сделали на тему планировщика (его ручного управления) - судя по пресс-релизу.

В детали не вдавался пока.

Результат для Linux'а такой

Результат для Linux'а такой же(с погрешностью меньше процента). CentOS 5.4 x64, драйвера свежие.
На моём приложении в Linux скорость тоже такая же, как и в Windows.

"С другой стороны, я надеюсь

"С другой стороны, я надеюсь через несколько дней сам такой тест проделать, вместе с прочими."

По поводу установки, помимо питания, нужно учесть размеры - я взял корпус Thermaltake Aguila.
У него в том месте куда может упереться видеокарта стоит съёмный блок на 4 винчестера. Надеялся что снимать не прийдётся, но там слишком в притирку получается(не стал пытаться запихивать) + лучшая вентиляция видеокарты.
А для винчестеров там много другого места(несколько 5.25 отсеков + отсек где кнопка питания).
http://www.thermaltake.com/product_info.aspx?PARENT_CID=C_00000193&id=C_...
http://www.overclockercafe.com/Reviews/cases/Tt_Aguila/pg3.htm

Во время установки была ещё одна небольшая неприятность: в корпусе для установки PCI устройств предусмотрены "tool-free" задвижки, но они не подходят к корпусу видеокарты из-за пластика.. но они легко снимаются и видеокарта прикручивается обычным образом(хорошо что резьбу сделали).

По поводу шумности, во время простоя не выделяется из общего звукового фона компьютера - очень тихо.
Если нагрузка кратковременная - не успевает нагреться.
Если же длительная - проходит больше минуты пока начинает заметно раскручиваться.
Когда рускручивается, получается очень шумно, но замерить шум нечем(обычные микрофоны ведь не подойдут, хотя..)

Размеры

По поводу установки, помимо питания, нужно учесть размеры

Размеры такие же, как у GTX280 и у ATI5870, до копейки.

Шумность-перегретость тоже преувеличены, как мне кажется.

По поводу перегретости - мне

По поводу перегретости - мне всё равно, главное чтобы не сломалась и особо не грела всё вокруг.

По поводу шумности - >4700RPM (GPU-Z) для меня шумновато, но если использовать для расчётов, то это не проблема - загрузил и чай пить.
Но, я предполагаю если постоянно находится рядом, то будет не очень комфортно.
Как уже говорил, замерить шум нечем, но можно поискать инфу - например http://www.youtube.com/watch?v=ubNulVNCXT0 . Где-то также читал, что действительно, измеряемый шум намного меньше чем тот, который был указан в первых обзорах(дБ на 10!), но это не достоверная информация.
Для кого шумность и перегретость проблема, могут поставить водяное охлаждение

Ну я что вижу - то пою.

Ну я что вижу - то пою. Температура idle у меня 75, но охлаждение хорошее (корпус большой) и пока никакими тестами не удалось нагреть выше 86. При этом вентилятор не шумит так, как шумел на 285-й.

Ящик под столом, никакого значимого увеличения уровня шума в idle-режиме я не вижу.

А вот две карты (эта и 5870) действительно уживаются так себе, мощность у 480й изрядно побольше все-таки и все вокруг греется несколько сильнее. Радеон пока вытащил, дальше посмотрим.

Цифирки bandwidth у меня примерно такие же, MatrixMul OpenCL-ный тоже большого прироста не дал. SGEMM посмотрю на выходных.

idle не громкий.. если

idle не громкий.. если недолгая нагрузка тоже не громко..
Корпус нормально проветривается, даже есть дырочки на дне.
А вот попробуйте nbody запустить с -n=31000, и подождите минут десять. Кстате этот же тест у меня на Linux'е выдает примерно на 100 GFlops больше, может там конечно какой-нибудь режим другой.

А вот cuBlas'овский SGEMM показывал очень маленький результат, но про это написано в ReleaseNotes.

Ситуация, когда на Linux все

Ситуация, когда на Linux все быстрее (или медленнее) - нормальная в первые полгода. Да и потом бывает. С драйверами вечно всякие странности.

nbody подогрел карту до 92, шумит примерно как 280-я на полном ходу по громкости, но 280-я шумела противнее.

Вообще, на старых приложениях прирост не такой большой, как внутренне ожидалось: горшков ведь стало вдвое больше, частота примерно та же, а двукратного роста я нигде не вижу. Вот на nbody - меньше чем раза в полтора (было, если правильно помню, 350 гигафлопс, сейчас 527), на моих всяких задачах - даже меньше.
Явно надо несколько иначе писать (скажем вот текстурный кэш уже не нужен, раз есть свой, нормальный).

По поводу NBody - на линуксе

По поводу NBody - на линуксе что-то около 650 было.
По поводу увлечения производительности в двое - можно какое-нибудь bogoFlops запустить.
По поводу других приложений - многие зависят от скорости памяти. Также теперь запрос в память идёт от варпа, а не от половины, это надо учитывать.
По поводу кэша - он открывает новые возможности, например если для варпа нужно считать два числа из глобальной памяти, раньше считывали двумя потоками в локальную память, а оттуда читали остальные(сейчас ещё +broadcasting), сейчас можно всеми потоками считать сразу из глобальной памяти(на таком примере я получил прирост скорости, но могу ошибаться, может причина в другом). Это даёт меньший divergence branch.

Насчёт 92 градуса - сначало

Насчёт 92 градуса - сначало на 92 двух градусах турбина не сильно крутиться, но потом крутиться сильнее видно для поддержания этой температуры.

Bandwidth тесты делят байты

Bandwidth тесты делят байты на 2^20 и соответственно получают мегабайты. Чтобы получить Gb/sec вы на сколько результат делили(2.3% - тоже разница)?

Я в уме делил, отбрасывая три

Я в уме делил, отбрасывая три последние цифры. 2% - не разница, если мы ожидаем оной разницы десятки процентов.

Итого, Если брать результат

Итого,
Если брать результат по "неправильным" цифрам прирост скорости 13.18%
Если по правильным 15.91%
По спекам 25.19%

вот бы ещё с GTX285 сравнить

на Tesla S1070

на Tesla S1070 (CentOS 5.3, CUDA Toolkit 2.3, driver ver. 190.53):
device 3:Tesla T10 Processor
Quick Mode
Host to Device Bandwidth for Pageable memory

Transfer Size (Bytes) Bandwidth(MB/s)
33554432 3579.5

Quick Mode
Device to Host Bandwidth for Pageable memory

Transfer Size (Bytes) Bandwidth(MB/s)
33554432 2520.9

Quick Mode
Device to Device Bandwidth

Transfer Size (Bytes) Bandwidth(MB/s)
33554432 73896.2

Nexus

Ещё раз отправил запрос на аккаунт для Nexus, в этот раз они активировали его.
Если будет время то попробую запустить, но везде где написано про совместимость нету архитектуры Fermi(вот пример http://supportcenteronline.com/ics/support/KBAnswer.asp?questionID=66 ).

OpenCL Visual profiler (пока

OpenCL Visual profiler (пока проверил только тот, что под Windows), для GTX480 не показывает таких счётчиков как coalesced и т.д.
Или я не там смотрю?
http://img17.imageshack.us/img17/5445/counters.png
В документации к профилеровщику (openclprof.html) это тоже описано.
Засада...

CudaProfiler даёт на один счётчик больше - sm cta launched

Для себя полезными вижу только счётчики по кэшу и divergent branch

В чем подвох? Точно помню что

В чем подвох? Точно помню что где-то читал что GTX480 может одновременно выполнять >60k потоков.
Количество мультипроцессоров 15, пусть даже имелись в виду резидентные потоки..
CudaProgogramming Guide:
Compute Capability 2.0:
Maximum number of resident threads per multiprocessor:1536
в общем явно не дотягивает до 60k

Другой вопрос - сколько потоков одновременно может выполнять один SP на GTX480? 4 или 2(два планировщика варпов на SM, и в одном SM 32 SP)?
И где это подробно описано(железная архитектура на уровне SP,SM...), интересуют именно информация от NVidia, или хотя бы из очень надёжного источника.

Да нет никакого подвоха, это

Да нет никакого подвоха, это речь про одновременное исполнение на одном мультипроцессоре.
В 8800 было 768 тредов на SM, в 280-й - 1024, теперь 1536. Ограничение планировщика, даже если ваши блоки(треды) кушают мало регистров и shared memory, больше блоков на один SM не запихают.

Соответственно, на GTX480 одновременно будет 45k в пределе

О, точно, не 45k а 22.5к,

О, точно, не 45k а 22.5к, мультипроцессоры то потолстели вдвое. Кажется. Надо разбираться. И это, похоже, причина того, почему исполнение старого кода мне перестало нравится, он именно под старую набивку SM-ов оптимизировался.

То есть получается, что у

То есть получается, что у GTX280 могут быть резидентными в сумме 30k потоков - больше чем у GTX480?
А количество одновременно работающих такое же(так как у GTX480 на SP работают два, а не четыре)?
[Понятное дело, что архитектурные улучшения, увеличенные частоты, большая ПСП, может быть за счёт всего этого нет необходимости в большом количестве резидентных потоков]...

Итак, опишу некоторые

Итак, опишу некоторые результаты которые я получил.
Сразу говорю, что времени на всё не хватает, помимо диплома есть ещё и работа.
Задача которую я решал(точнее использованный метод) - МКЭ.
Метод решения СЛАУ - метод сопряжённых градиентов с preconditioner Якоби.(почему-то попадалась литература на английском языке, как точно переводится этот термин не знаю.. может предобуславливатель.. может всё целиком метод Якоби.. хотелось бы найти нормальное описание на русском)
Так как полученные результаты скорости нужно с чем-то сравнить, причём сравнивать с собственной программой для CPU не культурно(вот, смотрите, 1000 кратное увлечение скорости, по сравнению с наивной CPU реализацией), решил сравнить результат с AnSys'ом(в нём есть JCG, причём выводится подробная информация - сколько итераций, сколько MFlops на SPMV достигнуто). В связи с этим нужно использовать одинаковые конечные элементы. AnSys позволяет создавать новые типы элементов(либо через DLL, либо через статическую(!!! - exe'шник AnSys'а перелинковывается) линковку), но прежде чем кидаться в бой, решил проверить тот элемент, который у меня был - двухмерный треугольник с двумя степенями свободы в узлах, плоское-напряжение. Выяснилось, что в AnSys'е есть точно такой же элемент - PLANE42(я сравнивал матрицы жёсткости - нашёл около пяти способов достать матрицы жёсткости из AnSys).
Размерность задачи - 3000000 степеней свободы(включая граничные условия), причём в AnSys'е там где опора из глобальной матрицы жёсткости вычёркивается строка и столбец , а я, к сожалению, когда писал поставил в этом месте на диагональ большое число, переделывать нет времени, может после сдачи займусь(то есть у меня размерность СЛАУ больше на число опёртых степеней свободы - на 2000).
По поводу double, и single вычислений - single не на много быстрее чем double(за счёт того, что помимо самих матриц и векторов, есть ещё служебная информация о структуре матрицы), а сходится очень медленно, я считаю что нету смысла даже начинать single.
Оптимизацию ядра SPMV делал под double вычисления, причём матрица хранится во float, а вектора и промежуточные данные в double (назовём это SD - SingleDouble)- это связанно с тем, что обычно данные на основе которых формируется матрица жёсткости достаточно неточные, вряд ли есть смысл использовать double. Сделал это для экономии bandwidth, так как SPMV целиком зависит от неё.
Хотя ядра не оптимизировал под SS и DD, результаты приведу.
Тестовая платформа:
Asus P6T, 3xDDR3-1333 2GiB (SisSoftware Sandra показывает 22.8 Гб/c), Core I7 930(включены HT и TurboBoost), GTX480, Windows 2008R2
В Ansys'е включено 8 потоков(по сравнению с одним, очень маленький прирост, около 10%, наверно в связи с тем, что Bandwidth важнее для SPMV. А также с TurboBoost и SmartCache у Core I7).
Количество итераций в секунду(Iter/S): 11.5
GFlop/s(только SPMV): 1.31
Результаты моей программы:

Тип матрицы и векторов Iter/s GFlop/s Эффективное GiB/s Flops ускорение Iter/S ускорение
Single Double 118.18 17.22 130.97 13.15 10.28
Double Double 94.11 11.65 110.35 8.89 8.18
Single Single 141.58 19.22 110.36 14.67 12.31

Причём я знаю ряд улучшений которые можно сделать, для увеличения скорости процентов на 30% при текущем алгоритме. Либо получить где-то 3х(оптимистично, но пока кажется выполнимым, так как планируется в большей степени нагрузить GPU вычислениями) при смене общей концепции.
Важно отметить, что используя не оптимизированные ядра, можно получить раза в 2.5 меньшую скорость(как было у меня в начале).
Также, посмотрев существующие реализации CG на GPGPU, увидел много недостатков. Одна из причин в следующем: итерация CG состоит из серии шагов, причём в некоторых местах выход одного шага сразу используется на следующем, в итоге лучше реализовывать ядра, которые делают сразу несколько шагов, а не тасовать данные в global mem и обратно.

Главный вопрос на который я

Главный вопрос на который я пытался ответить: стоит ли вообще пытаться решать МКЭ на GPGPU.
Ответ: на данный момент времени, конечно же стоит - видеокарты превосходят CPU на порядки по основными скоростным показателям - ПСП и GFlop/s, при решении задач МКЭ эти показатели являются определяющими скорость(bottleneck либо в ПСП, либо в GFlop/s).

Кстате, если у меня будет время, желание и стимул, то интегрирую свой решатель в ANSYS(пока нашёл только несколько возможных вариантов реализации).

У меня тема диплома - решение

У меня тема диплома - решение СЛАУ методом Холецкого (еще называется квадратного корня) на OpenMP. То есть само преобразование, и затем решение двух треугольных систем. Изначально система положительно определенная, ленточная, симметричная, что значительно упрощает решение. Вот подумываю, а не портировать ли быстренько на CUDA, сравнить результаты. Если что - отпишусь.

Если время позволяет, то

Если время позволяет, то почему бы и нет. Всё же я советую написать сначала пояснительную записку, а потом уже экспериментировать и в случае получения какого-то результата вставить отдельную главу.

Имхо реализовывать на Cuda будет несколько быстрей чем на OpenCL - меньше служебных вызовов, соответственно меньше возможных багов.
Также если не возиться с интерфейсами(абстрактными классами) и не заботиться о гибкости, а непосредствено реализовывать решатель, можно управиться за день, тем более у вас есть опыт реализации этого алгоритма параллельно.

Интересно, почему было выбрано OpenMP, а не допустим MPI(конечно это разные вещи, но всё же). И дало-ли использование многопоточности на системе с разделяемой памятью ощутимый прирост скорости(тут конечно всё от системы зависит, может огромная ПСП и относительно слабые потоки как на видеокарте, то без параллельности никак).

Два вопроса для понимания цифр

1) у AnSys - все в double?

2) А почему так получается, что ускорение по итерациям в секунду и по флопсам - разное? Не сильно, но заметно....

1) Скорей всего при расчётах

1) Скорей всего при расчётах - всё double, по крайней мере матрицы в double формируется(около 17 значащих цифр) и вектор сил хранится в double.
2) По-этому поводу есть несколько предположений:
1]Это может быть связанно с немного разными размерностями системы, но вроде эта разница не такая большая и не должна сильно влиять. В связи с этим может быть, что ускорение времени выполнения SPMV меньше чем GFlop/S ускорение(Количество FLOP считал так: NonZeros*2).
2]Хоть матрица и симметричная, хранится она у меня не как симметричная. SPMV запускается как несколько потоков на строку матрицы, если использовать при этом симметричность - может возникать ситуация записи несколькими потоками в одно место global mem(есть несколько идей как это обойти - например, иметь несколько векторов результатов, хранить дополнительную инфу кто куда должен писать, а после суммировать. Но тут надо ещё продумать нормальное общение с памятью - если потоки будут писать в произвольные места памяти, будет очень uncoalesced). Кстате, многие реализации CG тоже забивают на симметричность.
3]Матрица хранится у меня в формате CSR, хотя для МКЭ выгодней использовать блочный CSR - BCSR(я не стал с этим заморачиваться, так как у меня каждый узел может иметь разное количество степеней свободы, в связи с чем блоки могут быть разного размера - в таком случае я думаю делать несколько раз SPMV для частей матриц с разными размерами блоков). Может Ansys использует BCSR (в одном из форматов записи матриц жёсткости, AnSys использует CSR)
4]Выгодно умножать главную диагональ отдельно, точнее даже блочную диагональ, в этом случае нету никакой служебной инфы, а только числа. Может AnSys использует это.
5]Может не все операции на итерации у меня имеют одинаковое ускорение
6]Во время одной итерации у меня есть копирование с устройства результата скалярного произведения, для вычисления коэффициента, на который должен умножаться вектор, в принципе это можно избежать - запустив ядро для вычисления коэффициента, которое будет просто делить два числа, но профилировщик говорит что это копирование занимает очень мало времени. Тем не менее полностью избежать копирования чего-то с устройства может не получиться - нужно ведь знать какая сейчас невязка, чтобы знать когда остановиться. Может можно попробовать как-то маякнуть с устройства вызвав что-то типа exception(если вообще такое возможно). Проверять невязку не на каждой итерации можно, но особенность CG, в том, что невязка уменьшается не на каждой итерации - бывает сильно скачет.

Конечно можно ещё повозиться с текущим алгоритмом, выжать из него ещё процентов 50 скорости, но я считаю для МКЭ лучше использовать принципиально другой подход.

Сделал сравнение

Сделал сравнение результатов.
DoubleDouble идентично(6 значащих цифр. дальше не смотрел, там скорей всего начнутся расхождения из-за разного порядка операций и немного разными СЛАУ) тому же, что выдаёт AnSys на такой же точности.
SingleDouble на небольших сетках даёт такой же результат, но с увеличением размерности результат начинает расходится, может можно снизить это расхождение некоторыми приёмами..

Также по-быстрому прогнал задачу на Asus P5Q3,Core2DuoE8500,2xDDR3-1333(реально из-за FSB всё равно, 1333 или 800), там в три раза медленней чем на CoreI7. ПСП по тестам тоже в 3 раза меньше. Соответственно на видеокарте примерно в 30 раз быстрее.

T-Edge

Сегодня получил возможность прогнать свою задачу на кластере из стоек T-Edge 8 (2.33GHz), - http://www.t-platforms.ru/ru/clusters/clusters/complete/tedge.html.
Пиковая производительность стойки 300GFlops.
Всего четыре таких стойки.
В одной стойке 4 узла+1 мастер узел(вроде такой-же).
В каждом узле 2 четырёх ядерных Xeon'а (архитектура типа Core2Duo), 16 ГБ памяти.
Установлен ANSYS (MPI) - лицензия на 32 процесса(можно ими забить полностью 4 узла, но можно и распределить).
Раньше ANSYS запускали только на одной из двух стоек (т.к. думали что быстрее), и соответственно под ansys настроены только две стойки.
Сегодня я экспериментировал и показал, что если распределять потоки по разным стойкам, то получается быстрее(примерно в два раза), за счёт того, что итерационные методы упираются в основном в ПСП.

Сравнение по скорости с GTX480:
(GTX480 оказалась быстрее)
ускорение по времени: 1.98 раз
ускорение GFlops(на SPMV): 1.33 раз.
Интересно, что на кластере SPMV занимает примерно третью часть времени от всего решения, когда на одном компьютере(Core I7 из предыдущего поста) две трети. Это может быть связанно с латентностью сети на операциях редукции(которые происходят на каждой итерации, как и SPMV(на SPMV, нет обмена между узлами)).

Защитил сегодня диплом на

Защитил сегодня диплом на отлично - комиссия внимательно слушала доклад и ответы на вопросы, работа им понравилась. Спасибо данному сайту и его автору за поддержку!
Если будет время, выложу более подробное описание результатов - с графиками, с картинками моделей и т.п.

Зашёл скачать новый cuda

Зашёл скачать новый cuda toolkit, и увидел много вкусного:

- СUSPARSE, supporting sparse matrix computations.
- CUSPARSE, a new library of GPU-accelerated sparse matrix routines for sparse/sparse and dense/sparse operations
- Conjugate Gradient Solver, demonstrating the use of CUBLAS and CUSPARSE together Function Pointers, a sample that shows how to use function pointers to implement the Sobel Edge Detection filter for 8-bit monochrome images

- CUBLAS performance tuned for Fermi architecture GPUs, especially for matrix multiplication of all datatypes and transpose variations.
-CUBLAS performance improved 50% to 300% on Fermi architecture GPUs, for matrix multiplication of all datatypes and transpose variations

- Support for malloc() and free() in CUDA C compute kernels

GTX 480: SGEMM -

GTX 480:
SGEMM - 797GFlops
DGEMM - 162GFlops

Хотя это поверхностные тесты

Pages