Вопрос про параллельное программирование

[dil]

Вдруг кто-нибудь из читателей этот вопрос уже проходил?

В CUDA параллельная программа запускается в виде множества тредов. Треды логически группируются в блоки, а блоки – в гриды. Треды внутри одного блока имеют доступ к общей быстрой памяти и могут синхронизироваться, а треды из разных блоков могут общаться только через медленную память и не имеют встроенных средств синхронизации.

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

Физически треды выполняются на thread processor’ах, блоки – на мультипроцессорах (каждый мультипроцессор обычно содержит 8 thread processor’ов, один блок вычислений с плавающей точкой с двойной точностью и быструю общую память, доступную только изнутри этого мультипроцессора).

А грид выполняется на всём GPU, содержащем N мультипроцессоров (например, 10).

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

Или где про это можно почитать?

Upd: на первый взгляд логично тредов делать ровно столько, сколько может физически одновременно выполнить GPU: размер блока сделать равным количеству thread processor’ов в мультипроцессоре, а размер грида равным количеству мультипроцессоров.

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

Оригинал этой записи. Комментировать можно тут или там.

Любые материалы из этого блога запрещается использовать на сайте livejournal.ru в любой форме и любом объёме

Comments

[dmarck.livejournal.com]

У Тутубалина б тебе спросить, вот что я думаю...

[dil.livejournal.com]

это на крайний случай, не могу ж я у него всё подряд спрашивать :)

[alextutubalin.livejournal.com]

Есть дурацкий способ - надо помянуть меня в ЖЖ и я сам приду.

С конца
1) Перенос большими порциями интегрально наверное быстрее (ну там ~4Gb/sec на хорошей машине, но время на сетап ненулевое). Но ты же можешь таскать данные асинхронно, пока кернел считает одну порцию, подтаскивать-утаскивать другие.
Если время трансфера существенно (относительно счета) - это окупится.

2) Оптимальный размер задачи при прочих равных.
Я даже начал писать, а потом понял, что излагаю 4-й раздел из best practices guide, проще там прочесть

[dil.livejournal.com]

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

2) кажется, я ещё рановато приступил к практическим задачам. таки пошёл дальше читать TFM.

[alextutubalin.livejournal.com]

Кстати, если дело происходит на ноутбуке, то есть еще zero copy.

[dil.livejournal.com]

не, это десктоп с внешней видеокартой

[alextutubalin.livejournal.com]

Ну тогда нет.

На Макбуке (9400M) просто волшебный эффект.

[dil.livejournal.com]

в смысле, которая не на матери

[dil.livejournal.com]

zero copy попробовал, получилось медленнее, чем с копированием из хоста в девайс

больше всего меня удивило, что на общую скорость приложения существенно влияет способ выделения памяти в хосте, которая потом копируется в девайс.
с обычным malloc'ом получилось медленнее всего, cudaHostAlloc(cudaHostAllocDefault) быстрее, cudaHostAlloc(cudaHostAllocWriteCombined) быстрее всего.

[alextutubalin.livejournal.com]

Ну собственно про это и в мане написано :)

[dil.livejournal.com]

да, но я думал, что cudaHostAllocDefault и malloc - это одно и то же, ан нет. все четыре варианта cudaHostAlloc работают быстрее, чем malloc.
а zeroCopy через cudaHostGetDevicePointer() медленнее. хотя логично, что доступ из видеокарты к памяти хоста медленнее, чем к своей собственной.
хотя это должно существенно зависеть от объёмов. у меня данных относительно много, а расчёты относительно простые - двойной md5.

[alextutubalin.livejournal.com]

malloc() и родные функции могут отличаться выравниванием, например.

А система какая, кстати? Винды?

[dil.livejournal.com]

да не, откуда у меня винды :) линукс
с выравниванием вряд ли, память выделяется большими пачками, кратными 2 килобайтам (32*64).

да, и что интересно, при разных методах выделения памяти в хосте максимальное быстродействие достигается для разных сочетаний blocks per grid и threads per block.
для malloc: 160x128,
для cudaHostAllocDefault и cudaHostAllocPortable.txt: 480x320
для out-cudaHostAllocMapped.txt: 512x480, и возможно, она не максимальная, я только до 512 блоков на грид проверял.
для cudaHostAllocWriteCombined: 480x288

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

[alextutubalin.livejournal.com]

Ну а под linux тебе strace или как его там зовут небось расскажет, как именно там маллок зовется и зачем. Собственно, это должен быть бы не malloc, а mmap на /dev/zero с разными флажками.

А время трансфера померять и асинком его, асинком!

[dil.livejournal.com]

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

[dil.livejournal.com]

до асинхронной передачи данных я ещё не дошёл, но даже просто подготовка следующей порции данных на CPU параллельно с обсчётом предыдущей на GPU дала прирост производительности с 12 до 20 миллионов ключей/с. какие горизонты для оптимизации..

[alextutubalin.livejournal.com]

Вообще, программирование этого дела - чистое удовольствие.

[dil.livejournal.com]

натурально :)
а то я за повседневным скриптингом уже почти разучился получать удовольствие от красивого кода :)