CUDA веднага

NB: Статията е кратко въведение, едва ли е възможно да се обхванат всички нюанси на програмирането под CUDA в една статия :-)

Откъде да започна?

NB: Ако добавите опцията -keep, можете да намерите много интересни междинни файлове след компилация.

Дефиниция на функцията

Дефиниране на данни

__constant__ - задава променлива в постоянна памет. Моля, обърнете внимание, че стойностите за константите трябва да бъдат заредени с функцията cudaMemcpyToSymbol. Константите са достъпни от всички нишки, скоростта е сравнима с регистрите (когато попадне в кеша).__shared__ - задава променлива в споделената памет на блока на нишката (т.е. стойността ще бъде споделена от всички). Тук трябва да подходите с повишено внимание - компилаторът агресивно оптимизира достъпа тук (можете да го удушите с модификатора volatile), можете да получите състояние на състезание, трябва да използвате __syncthreads(); за да се гарантира, че данните са записани. Споделената памет е разделена на банки и когато 2 нишки едновременно се опитат да получат достъп до една и съща банка, възниква конфликт на банка и скоростта пада.

Всичкилокални променливи, които сте дефинирали в ядрото (__устройство__), са в регистри, най-бързата скорост на достъп.

Как една нишка знае върху какво да работи?

Основната идея на CUDA е, че стартирате хиляди и хиляди нишки, за да разрешите проблема си, така че не се страхувайте от това, което ще бъде написано тук :-) Да кажем, че трябва да извършите някаква операция върху изображение 200x200. Картината е разделена на парчета 10x10 и за всеки пиксел от такова парче минаваме по протежение на потока. Ще изглежда така: dim3 threads(10, 10);//размер на квадрата, 10*10 dim3 grid(20, 20);//колко квадрата са ви необходими, за да покриете цялото изображение

вашето_ядро >>(изображение,200,200);//Този ред ще стартира 40'000 нишки (не едновременно, приблизително 200-2000 нишки ще работят едновременно).

За разлика от Brook + от AMD, където ние веднага определяме коя нишка да работи върху какви данни, в CUDA всичко е различно: параметрите, предадени на ядрото, са еднакви за всички нишки и нишката трябва да получи данните за себе си, за да направи това, нишката трябва да изчисли къде в изображението се намира. Магическите променливи blockDim, blockIdx помагат за това. const int ix = blockDim.x * blockIdx.x + threadIdx.x; const int iy = blockDim.y * blockIdx.y + threadIdx.y;

В ix и iy - координати, с които можете да получите първоначалните данни от масива на изображението и да напишете резултата от работата.

Оптимизация

Няколко думи за това как да не направите програмата си много бавна (да напишете програма, която работи по-бавно от процесора, е много по-лесно, отколкото да напишете програма, която работи 10 пъти по-бързо :-) )

  • Използвайте __global__ памет възможно най-малко.
  • Когато работите с __споделена__ памет, избягвайте банкови конфликти (обаче много задачи могат да бъдат решени без споделена памет).
  • Възможно най-малко разклонения в кода, където различните нишки поемат по различни пътища. Такъв код не се изпълнява паралелно.
  • Използвайте възможно най-малко памет. Колкото по-малко памет използвате, толкова по-агресивно компилаторът и хардуерът ще могат да изпълняват вашето ядро ​​(например, може да отнеме 100 нишки и с помощта на още 100 регистъра да се изпълняват едновременно на един MP, драстично намалявайки латентността)

Не работи?

Първо, трябва да прочетете документацията заедно с SDK (NVIDIA_CUDA_Programming_Guide, CudaReferenceManual, ptx_isa), след това можете да попитате в официалния форум - дори разработчиците на nVidia често се отписват там и наистина многоумни хора. На български можете да ме попитате във форума например къде ще отговоря :-) Пък ​​и много хора живеят в gpgpu.ru.

Hardcore conf в C++. Каним само професионалисти.