LXF169: Вы­чис­ления. CUDA

Текущая версия на 09:42, 14 ноября 2018

Па­рал­лель­ные тех­но­ло­гии

[править] CUDA: Ус­ко­ря­ем гра­фи­ку

Кон­стан­тин Кал­гин и Ев­ге­ний Бал­дин ух­ва­ти­лись за иг­ро­вые тех­но­ло­гии с це­лью про­дви­нуть ре­ше­ние на­уч­ных за­дач.

Наш эксперт Кон­стан­тин Кал­гин программист, который превращает персоналку в маленький супер-компьютер.

Ев­ге­ний Бал­дин Физик, который действительно знает, что такое нехватка вычислительных ресурсов.

Вам нуж­ны ги­гаф­лоп­сы здесь и сей­час? Вы не хо­ти­те ждать свет­ло­го бу­ду­ще­го OpenCL и со­глас­ны ра­ди дей­ст­ви­тель­но стоя­ще­го де­ла вля­пать­ся в «вен­дор-лок»? Тогда вы­бо­ра как бы и нет – CUDA ждет но­во­го адеп­та.

[править] Пре­дыс­то­рия

Как пра­ви­ло, лю­ди тра­тят день­ги ли­бо на хлеб, ли­бо на зре­ли­ща. На пер­вое тра­тят по­то­му, что хо­чет­ся до­жить до вто­ро­го. Иг­ро­вой плебс тре­бу­ет кра­си­вой кар­тин­ки – так и возник­ли гра­фи­че­­ские уско­ри­те­ли.

Из­на­чаль­но гра­фи­че­­ские уско­ри­те­ли (ви­део­кар­ты, GPU, Graphical Processing Units) пред­на­зна­ча­лись для вы­во­да дву­мер­ной или трех­мер­ной гра­фи­че­­ской ин­фор­ма­ции на эк­ран. Дол­гое вре­мя про­цесс ото­бра­жения гра­фи­че­­ской ин­фор­ма­ции на эк­ране управ­лял­ся лишь струк­ту­ра­ми дан­ных – мас­си­ва­ми при­ми­тив­ных фи­гур, тек­сту­ра­ми и про­стей­ши­ми цве­то­вы­ми фильт­ра­ми.

Толь­ко в XXI ве­ке, хоть и в са­мом его на­ча­ле, в гра­фи­че­­ских уско­ри­те­лях поя­ви­лась под­держ­ка шей­де­ров или мини-про­грамм об­ра­бот­ки дан­ных на раз­лич­ных ста­ди­ях гра­фи­че­­ско­­го кон­вей­е­ра. Про­цесс ото­бра­жения кар­тин­ки те­перь стал управ­лять­ся не толь­ко дан­ны­ми, но и мини-про­грам­ма­ми, ис­пол­няю­щи­ми­ся на са­мом гра­фи­че­­ском уско­ри­те­ле. Шей­де­ры да­ли воз­мож­ность раз­ра­бот­чи­кам соз­да­вать свои соб­ст­вен­ные спе­цэф­фек­ты, а не ог­раничи­вать­ся уже встро­ен­ны­ми в же­ле­зо. Это зна­чи­тель­но уве­ли­чи­ло слож­ность и реа­ли­стич­ность ком­пь­ю­тер­ной гра­фи­ки и, со­от­вет­ст­вен­но, при­вле­ка­тель­ность ито­го­вой кар­тин­ки для из­ба­ло­ван­но­го зри­те­ля.

Фак­ти­че­­ски сра­зу по­сле пер­вых шей­де­ров гра­фи­че­­ские уско­ри­те­ли на­ча­ли ис­поль­зо­вать­ся эн­ту­зиа­ста­ми для ре­шения негра­фи­че­­ских за­дач, то есть для за­дач об­ще­го на­зна­чения (GPGPU, General Purpose computations on GPU). На­ча­ли по­яв­лять­ся ста­тьи, в том чис­ле и в на­уч­ных жур­на­лах, о при­менении гра­фи­че­­ских уско­ри­те­лей для ре­шения до­воль­но важ­ных и час­то ис­поль­зуе­мых за­дач линей­ной ал­геб­ры, а так­же за­дач, ко­то­рые не име­ют ана­ли­ти­че­­ско­­го ре­шения в об­щем ви­де – вро­де мо­де­ли­ро­вания сис­те­мы N гра­ви­ти­рую­щих тел. За­чем серь­ез­ные, ка­за­лось бы, лю­ди тра­ти­ли си­лы и вре­мя на же­ле­зо, це­ли­ком и пол­но­стью ори­ен­ти­ро­ван­ное на иг­ро­вую ау­ди­то­рию? От­вет прост: це­на на про­из­во­ди­тель­ность. В си­лу боль­шо­го спро­са ка­зу­аль­ной ау­ди­то­рии на раз­вле­чения гра­фи­че­­ские уско­ри­те­ли ста­ли мощ­ны­ми и от­но­си­тель­но де­ше­вы­ми, хо­тя и уз­ко спе­циа­ли­зи­ро­ван­ны­ми вы­чис­ли­тель­ны­ми сис­те­ма­ми. На пу­ти эн­ту­зиа­стов бы­ла толь­ко од­на про­бле­ма: как бы сде­лать про­це­ду­ру за­груз­ки всей этой мо­щи по­про­ще?

С по­яв­лением в 2007 го­ду про­грамм­но-ап­па­рат­ной ар­хи­тек­ту­ры CUDA гра­фи­че­­ских уско­ри­те­лей ком­пании Nvidia си­туа­ция кар­ди­наль­но из­менилась. Про­грам­мы ста­ли со­став­лять­ся не на спе­ци­аль­ном язы­ке опи­сания шей­де­ров, а на зна­ко­мом C/C++. В опи­сании гра­фи­че­­ских уско­ри­те­лей CUDA прак­ти­че­­ски пе­ре­ста­ли ис­поль­зо­вать­ся гра­фи­че­­ские тер­ми­ны, та­кие как шей­де­ры, точ­ки, тек­сту­ры, фильт­ра­ция, Z-бу­фер и пр. По­тен­ци­аль­но серь­ез­ный кон­ку­рент в ли­це OpenCL поя­вил­ся толь­ко в 2010 го­ду, по­это­му на се­го­дня CUDA яв­ля­ет­ся фак­ти­че­­ски един­ст­вен­ной зре­лой тех­но­ло­ги­ей ис­поль­зо­вания гра­фи­че­­ских уско­ри­те­лей. Во мно­же­ст­ве учеб­ных цен­трах су­ще­ст­ву­ют обу­чаю­щие про­грам­мы по тех­но­ло­ги­ям CUDA, и мно­же­ст­во спе­циа­ли­стов уже вла­де­ют необ­хо­ди­мы­ми знания­ми на вполне при­ем­ле­мом уровне.

К недостат­кам CUDA сле­ду­ет отнести за­кры­тость драй­ве­ров и непе­ре­но­си­мость ко­да за пре­де­лы плат­форм от Nvidia. Пер­вое мо­жет при­во­дить к весь­ма стран­ным абер­ра­ци­ям в по­ве­дении поль­зо­ва­тель­ских про­грамм и необъ­яснимым па­дениям про­из­во­ди­тель­но­сти в за­ви­си­мо­сти от вер­сии ком­понент этой про­грамм­но-ап­па­рат­ной плат­фор­мы. Вто­рое не по­зво­ля­ет от­ка­зать­ся от нее в поль­зу ис­поль­зо­вания дру­гих плат­форм, ко­то­рые сей­час доста­точ­но ак­тив­но раз­ви­ва­ют­ся и в пер­спек­ти­ве мо­гут обо­гнать сис­те­мы Nvidia.

Са­мый про­из­во­ди­тель­ный су­пер­ком­пь­ю­тер в Top500 (http://www.top500.org) на но­ябрь 2012 го­да вклю­ча­ет в ка­че­­ст­ве од­ного из сво­их эле­мен­тов уско­ри­те­ли Nvidia. Поч­ти де­сять про­цен­тов сис­тем из это­го спи­ска так­же ука­за­ли уско­ри­те­ли Nvidia в опи­сании сво­ей ар­хи­тек­ту­ры, и эта до­ля в обо­зри­мом бу­ду­щем бу­дет рас­ти.

[править] Ус­та­нов­ка CUDA

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

Ес­те­ст­вен­но, необ­хо­ди­мо убе­дить­ся, что в ва­шем ком­пь­ю­те­ре есть гра­фи­че­­ский уско­ри­тель от ком­пании Nvidia и что он под­дер­жи­ва­ет CUDA. За ин­фор­ма­ци­ей мож­но об­ра­тить­ся к той же Ви­ки­пе­дии: http://ru.wikipedia.org/wiki/CUDA.

За­тем нуж­но уста­но­вить за­кры­тый дво­ич­ный драй­вер по­све­жее от ком­пании Nvidia. На­при­мер, в Ubuntu 12.04 это де­ла­ет­ся ко­ман­дой

sudo aptitude install nvidia-experimental-310

При этом уста­нав­ли­ва­ет­ся драй­вер вер­сии 310.14. Здесь и да­лее мы ори­ен­ти­ру­ем­ся имен­но на этот ди­ст­ри­бу­тив. В слу­чае несов­па­дении пред­поч­тений с на­ши­ми в се­ти лег­ко най­ти по­ша­го­вую ин­ст­рук­цию фак­ти­че­­ски для лю­бо­го дру­го­го ди­ст­ри­бу­ти­ва.

По­сле осоз­нания то­го, что ваш до­машний ком­пь­ю­тер те­перь го­ден для уста­нов­ки CUDA, нуж­но за­ка­чать дво­ич­ный уста­нов­щик с сай­та раз­ра­бот­чи­ка https://developer.nvidia.com/cuda-downloads. Для Ubuntu 12.04 го­дит­ся би­нарник, со­б­ран­ный для вер­сии 11.10. Здесь же нуж­но вы­брать ме­ж­ду уста­но­воч­ны­ми фай­ла­ми 64-bit и 32-bit. Да­лее бу­дем дей­ст­во­вать в пред­по­ло­жении, что у вас 64-бит­ная вер­сия ди­ст­ри­бу­ти­ва. Раз­мер уста­но­воч­но­го фай­ла по­ряд­ка 670 МБ, но ка­ча­ет­ся он доста­точ­но бод­ро.

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

sudo aptitide install freeglut3-dev build-essential libx11-dev libxmu-dev libxi-dev libgl1-mesa-glx libglu1-mesa libglu1-mesa-dev

Ес­ли у вас 64-бит­ная вер­сия ди­ст­ри­бу­ти­ва, то для уста­нов­ки при­ме­ров необ­хо­ди­мо до­ба­вить сим­во­ли­че­скую ссыл­ку на биб­лио­те­ку libglut.so.

sudo ln -s /usr/lib/x86_64-linux-gnu/libglut.so.3 /usr/lib/libglut.so

Ус­та­но­воч­ный скрипт ищет эту биб­лио­те­ку по­че­му-то в /usr/lib/.

Те­перь за­пуска­ем про­грам­му уста­нов­ки:

chmod +x cuda_5.0.35_linux_64_ubuntu11.10-1.run

sudo ./cuda_5.0.35_linux_64_ubuntu11.10-1.run

Ак­ку­рат­но про­лис­ты­ва­ем и со­гла­ша­ем­ся с EULA, от­ка­зы­ва­ем­ся от уста­нов­ки драй­ве­ра для ви­део­кар­ты, но со­гла­ша­ем­ся с уста­нов­кой CUDA и при­ме­ров. По умол­чанию уста­нов­ка идет в /usr/local/, и ес­ли умол­чание не ме­ня­лось, то для на­ча­ла ра­бо­ты доста­точ­но пе­ре­оп­ре­де­лить пе­ре­мен­ные ок­ру­жения:

export PATH=$PATH:/usr/local/cuda/bin

export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/usr/local/cuda/lib64

Пе­ред на­ча­лом экс­пе­ри­мен­тов по­лез­но по­бро­дить по ди­рек­то­рии /usr/local/cuda, по­смот­рев на 100 МБ до­ку­мен­та­ции в ди­рек­то­рии doc и по­ко­вы­ряв­шись в при­ме­рах из ди­рек­то­рии samples. При­ме­ры со­би­ра­ют­ся с по­мо­щью ко­ман­ды make. В ди­рек­то­рии samples/0_Simple/template рас­по­ла­га­ет­ся за­го­тов­ка для стан­дарт­но­го CUDA-про­ек­та.

[править] Ис­ходники и их ком­пи­ля­ция

Про­грамм­ная часть ар­хи­тек­ту­ры CUDA опи­сы­ва­ет рас­ши­рение язы­ка C/C++, функ­цио­наль­ность и клю­чи управ­ления ра­бо­той ком­пи­ля­то­ра nvcc, ин­тер­фейс обо­лоч­ки CUDA Runtime сис­тем­но­го драй­ве­ра гра­фи­че­­ско­­го уско­ри­те­ля, про­фи­ли­ров­щик и от­лад­чик.

Во вре­мя ис­полнения на цен­траль­ном про­цес­со­ре ком­пь­ю­те­ра про­грам­ма за­пуска­ет функ­ции на гра­фи­че­­ском уско­ри­те­ле. Та­ких функ­ций в про­грам­ме мо­жет быть несколь­ко. Функ­ция, ис­пол­няе­мая на гра­фи­че­­ском уско­ри­те­ле, на­зы­ва­ет­ся ядром [kernel]. Во вре­мя за­пуска яд­ра по­ро­ж­да­ет­ся мно­же­ст­во по­то­ков, ко­то­рые бу­дут ис­пол­нять од­ну и ту же функ­цию на про­цес­со­рах гра­фи­че­­ско­­го уско­ри­те­ля. По­ве­дение вы­чис­ли­тель­ных по­то­ков за­ви­сит как от ко­да функ­ции, так и от их ко­ор­ди­нат.

Фай­лы ис­ходников име­ют рас­ши­рение .cu:

// example1.cu

1. include <stdio.h>

1. include <cuda_runtime_api.h>

int main() {

int N;

cudaDeviceProp prop;

// Под­счи­ты­ва­ем чис­ло уст­ройств CUDA

cudaGetDeviceCount(&N);

for (int i = 0; i < N; i++){

// По­лу­ча­ем ин­фор­ма­цию об уст­рой­ст­ве

cudaGetDeviceProperties(&prop, i);

// Вы­во­дим ифор­ма­цию об уст­рой­ст­ве

printf(“Уст­рой­ст­во N %d: %s\n”,i+1,prop.name);

}

}

Сбор­ка про­из­во­дит­ся с по­мо­щью ком­пи­ля­то­ра nvcc:

nvcc example1.cu -o example1

./example1

Уст­рой­ст­во N 1: GeForce GTS 450

В при­ве­ден­ном вы­ше при­ме­ре по­ка нет опи­сания яд­ра, и все ис­полнение идет на цен­траль­ном про­цес­со­ре без при­вле­чения уско­ри­те­ля.

Рас­ши­рение язы­ка C/C++ по­зво­ля­ет опи­сы­вать в од­ном ис­ход­ном фай­ле как основ­ную про­грам­му, так и яд­ра с пе­ре­мен­ны­ми и мас­си­ва­ми, ко­то­рые бу­дут рас­по­ла­гать­ся в па­мя­ти гра­фи­че­­ско­­го уско­ри­те­ля. Кро­ме то­го, в рас­ши­рении язы­ка име­ет­ся ком­пакт­ная кон­ст­рук­ция для за­пуска ядер, скры­ваю­щая вы­зов функ­ции драй­ве­ра и упа­ков­ку ар­гу­мен­тов.

Взаи­мо­дей­ст­вие с гра­фи­че­­ским уско­ри­те­лем осу­ще­ст­в­ля­ет­ся че­рез ин­тер­фейс CUDA Runtime сис­тем­но­го драй­ве­ра: за­пуск яд­ра, ди­на­ми­че­­ское вы­де­ление па­мя­ти в гра­фи­че­­ском уско­ри­те­ле, ко­пи­ро­вание дан­ных из/в па­мять гра­фи­че­­ско­­го уско­ри­те­ля.

Ком­пи­ля­тор nvcc раз­де­ля­ет вход­ной файл на две час­ти: од­на бу­дет ис­пол­нять­ся на цен­траль­ном про­цес­со­ре и ком­пи­ли­ро­вать­ся стан­дарт­ным ком­пи­ля­то­ром gcc, а вто­рая – на гра­фи­че­­ском уско­ри­те­ле, и ком­пи­ли­ро­вать­ся са­мим nvcc. Рас­ши­рения­ми CUDA язы­ка C/C++ мож­но поль­зо­вать­ся толь­ко в фай­лах с рас­ши­рением .cu.

[править] Ие­рар­хи­че­­ская ор­ганиза­ция по­то­ков

При про­грам­ми­ро­вании гра­фи­че­­ско­­го уско­ри­те­ля необ­хо­ди­мо быть в кур­се дик­туе­мой ар­хи­тек­ту­рой ие­рар­хии ис­пол­няе­мых по­то­ков и па­мя­ти. Это ключ к уско­рению па­рал­лель­ных вы­чис­лений. К со­жа­лению, в этом мес­те очень лег­ко оши­бить­ся и убить весь эф­фект боль­шо­го чис­ла про­цес­со­ров неэф­фек­тив­ным ко­пи­ро­ванием дан­ных.

При за­пуске про­грам­мы на гра­фи­че­­­ском уско­ри­те­ле по­ро­ж­да­ет­ся мно­же­ст­­во по­то­ков [thread]. Все по­то­ки по­де­ле­ны на груп­пы од­но­го раз­ме­ра – бло­ки по­то­ков [block]. Мак­си­маль­ный раз­мер бло­ка по­то­ков на со­вре­мен­ных гра­фи­че­­­ских уско­ри­те­лях ра­вен 1024. У ка­ж­до­го по­то­ка и бло­ка по­то­ков име­ют­ся свои уникаль­ные иден­ти­фи­ка­то­ры, на­зы­вае­мые ко­ор­ди­на­та­ми. Тем са­мым, для раз­ных по­то­ков ар­гу­мен­ты ис­пол­няе­мых ин­ст­рук­ций и их по­сле­до­ва­тель­ность мо­гут раз­ли­чать­ся, по­сколь­ку мо­гут за­ви­сеть от ко­ор­ди­нат по­то­ка и бло­ка по­то­ков. Мно­же­ст­­ва ко­ор­ди­нат по­то­ков и бло­ков по­то­ков об­ра­зу­ют од­но-, дву- или трех­мер­ные мас­си­вы-сет­ки [grid]. Раз­ме­ры бло­ка по­то­ков и мас­си­ва бло­ков по­то­ков за­да­ют­ся при за­пуске яд­ра.

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

Внут­ри ка­ж­до­го по­то­ка доступ­ны струк­ту­ры, ко­то­рые по­зво­ля­ют его иден­ти­фи­ци­ро­вать:

» threadIdx – ко­ор­ди­на­ты по­то­ка в бло­ке по­то­ков;

» blockIdx – ко­ор­ди­на­ты бло­ка по­то­ков в сет­ке;

» blockDim – раз­ме­ры бло­ка по­то­ков;

» gridDim – раз­ме­ры сет­ки бло­ков по­то­ков.

[править] Аз­бу­ка вы­зо­вов

В этом раз­де­ле пе­ре­чис­ле­ны про­стей­шие язы­ко­вые струк­ту­ры, ко­то­рые смо­гут при­го­дить­ся в про­цес­се вве­дения в сре­ду CUDA.

Функ­ции

Опи­сание про­из­воль­ной функ­ции мо­жет пред­ва­рять­ся сле­дую­щи­ми клю­че­вы­ми сло­ва­ми:

» __global__ – это функ­ция-яд­ро, ко­то­рая за­пуска­ет­ся на гра­фи­че­­ском уско­ри­те­ле. Тип воз­вра­щае­мо­го ре­зуль­та­та дол­жен быть void.

» __device__ – функ­ция, ко­то­рая вы­зы­ва­ет­ся из яд­ра.

» __host__ – функ­ция, ко­то­рая вы­зы­ва­ет­ся с цен­траль­но­го про­цес­со­ра.

Воз­мож­но од­но­вре­мен­ное ис­поль­зо­вание __device__ и __host__ – это оз­на­ча­ет, что функ­ция мо­жет быть вы­зва­на как из яд­ра, так и из про­грам­мы на цен­траль­ном про­цес­со­ре. По умол­чанию, при от­сут­ст­вии вы­ше­пе­ре­чис­лен­ных клю­че­вых слов, счи­та­ет­ся, что функ­ция бу­дет вы­зы­вать­ся толь­ко с цен­траль­но­го про­цес­со­ра (__host__).

Функ­ция-яд­ро дек­ла­ри­ру­ет­ся при­мер­но так:

__global__ void MyKernel(int *a,int *b,int *c) {

int i = threadIdx.x + blockDim.x * blockIdx.x;

c[i] = a[i] + b[i];

}

Здесь в ка­че­­ст­ве ар­гу­мен­тов пе­ре­да­ют­ся три ука­за­те­ля в об­ласть па­мя­ти ви­део­кар­ты. Ка­ж­дый по­ро­ж­ден­ный по­ток скла­ды­ва­ет i-е эле­мен­ты мас­си­вов a и b и за­пи­сы­ва­ет ре­зуль­тат в мас­сив c; но­мер эле­мен­та в мас­си­ве i вы­чис­ля­ет­ся в за­ви­си­мо­сти от ко­ор­ди­на­ты по­то­ка threadIdx.x, бло­ка по­то­ков blockIdx.x и раз­ме­ра бло­ка по­то­ков blockDim.x по оси Ox.

То же са­мое мож­но изо­бра­зить и с по­мо­щью вспо­мо­га­тель­ной функ­ции __device__:

__device__ int MyDev(int a,int b) {

return a + b;

}

__global__ void MyKernel(int *a,int *b,int *c) {

int i = threadIdx.x + blockDim.x * blockIdx.x;

c[i] = MyDev( a[i], b[i] );

}

Что­бы вы­звать функ­цию-яд­ро, необ­хо­ди­мо ука­зать раз­мер бло­ка по­то­ков и раз­мер сет­ки бло­ков. Для это­го ис­поль­зу­ют­ся трой­ные уг­ло­вые скоб­ки:

MyKernel<<< 256, N/256 >>>( a, b, c );

В слу­чае по­ро­ж­дения дву- или трех­мер­ных се­ток, что удоб­но при рас­че­те дву­мер­ных и трех­мер­ных за­дач со­от­вет­ст­вен­но, для пе­ре­да­чи ин­фор­ма­ции о раз­ме­ре раз­ме­рах бло­ка и сет­ки ис­поль­зу­ет­ся спе­ци­аль­ный тип дан­ных – dim3:

__global__ void MyKernel2D(int *a, int *b, int *c) {

int x = threadIdx.x + blockIdx.x * blockDim.x;

int y = threadIdx.y + blockIdx.y * blockDim.y;

int i = y * Nx + x; c[i] = b[i] + a[i];

}

dim3 blockSize( 16, 16, 1 );

dim3 gridSize( Nx/16, Ny/16, 1 );

MyKernel2D<<< blockSize, gridSize >>>( a, b, c );

Здесь про­ис­хо­дит сло­жение двух мат­риц раз­ме­ра Nx × Ny. Ко­ли­че­­ст­во по­ро­ж­дае­мых по­то­ков в точ­но­сти со­от­вет­ст­ву­ет ко­ли­че­­ст­ву эле­мен­тов в мат­ри­це.

[править] Ра­бо­та с па­мя­тью

При объ­яв­лении пе­ре­мен­ной для ее раз­ме­щения в гло­баль­ной па­мя­ти ис­поль­зу­ет­ся клю­че­вое сло­во __device__, для раз­ме­щения в раз­де­ляе­мой па­мя­ти – __shared__, а для раз­ме­щения в кон­стант­ной – __constant__.

Объ­яв­лен­ные внут­ри яд­ра пе­ре­мен­ные без ис­поль­зо­вания этих трех клю­че­вых слов ав­то­ма­ти­че­­ски ото­бра­жа­ют­ся на ре­ги­ст­ры или локаль­ную па­мять (об­ласть из гло­баль­ной па­мя­ти, об­ласть ви­ди­мо­сти – по­ток).

Тек­сту­ры объ­яв­ля­ют­ся с по­мо­щью шаб­ло­на texture< >. Раз­мер и об­ласть гло­баль­ной па­мя­ти, к ко­то­рой бу­дет при­вя­за­на тек­сту­ра, оп­ре­де­ля­ют­ся с по­мо­щью вы­зо­ва cudaBindTexture().

Функ­ция cudaGetSymbolAddress() че­рез пер­вый ар­гу­мент воз­вра­ща­ет ад­рес пе­ре­мен­ной или мас­си­ва, объ­яв­лен­но­го с по­мо­щью __device__ или __constant__:

texture< float, 1, cudaReadModeElementType > tex;

__device__ int dev_c;

__device__ float dev_a[ 1024*256 ];

__constant__ float dev_g[ 1024*256 ];

__global__ void MyKernel( float * c ){

int i = threadIdx.x + blockDim.x * blockIdx.x;

__shared__ int sb[ 256 ];

sb[ threadIdx.x ] = dev_g[ i ] + dev_c;

__syncthreads();

c[i] = sb[threadIdx.x] + tex1Dfetch(tex, i) + dev_g[i];

}

int main(){

float *p_dev_a;

cudaGetSymbolAddress( &p_dev_a, dev_a );

cudaBindTexture( 0, tex, p_dev_a, tex.channelDesc, 1024*256*sizeof(float) );

MyKernel <<< 1024, 256 >>>( dev_c );

}

Функ­ция cudaMalloc по­зво­ля­ет вы­де­лить об­ласть в гло­баль­ной па­мя­ти гра­фи­че­­ско­­го уско­ри­те­ля ука­зан­но­го раз­ме­ра и воз­вра­ща­ет ука­за­тель на эту об­ласть че­рез пер­вый ар­гу­мент:

cudaError_t cudaMalloc( void **, size_t );

При этом зна­че­ние са­мо­го ука­за­те­ля бу­дет хра­нить­ся в опе­ра­тив­ной па­мя­ти ком­пь­ю­те­ра:

int main(){

float *dev_a;

cudaMalloc( &dev_a, sizeof(int)*1024*256 );

cudaBindTexture( 0, tex, dev_a, tex.channelDesc, 1024*256*sizeof(float) );

MyKernel <<< 1024, 256 >>>( dev_c );

cudaFree( dev_a );

}

Функ­ция cudaMallocPitch вы­де­ля­ет об­ласть па­мя­ти для рас­по­ло­жения в ней дву­мер­ных мас­си­вов. При этом уве­ли­чи­ва­ет­ся раз­мер стро­ки в бай­тах до бли­жай­ше­го чис­ла, крат­но­го 128, что­бы ка­ж­дая стро­ка дву­мер­но­го мас­си­ва на­чи­на­лась со 128-байт­но­го сег­мен­та па­мя­ти. Это свя­за­но с оп­ти­ми­за­ци­ей досту­па в па­мять:

cudaError_t cudaMallocPitch( void** p_dev, size_t* pitch size_t width, size_t height );

где p_dev – воз­вра­щае­мый ад­рес вы­де­лен­ной па­мя­ти, pitch – воз­вра­щае­мый но­вый раз­мер стро­ки в бай­тах, width – раз­мер ис­ход­ной стро­ки в бай­тах, а height – ко­ли­че­­ст­во строк.

Ко­пи­ро­вание па­мя­ти осу­ще­ст­в­ля­ет­ся с по­мо­щью се­мей­ст­ва функ­ций cudaMemcpy:

cudaError_t cudaMemcpy(void* dst, const void* src, size_t count,enum cudaMemcpyKind kind);

cudaError_t cudaMemcpyToSymbol(const char* dst, const void* src, size_t count, size_t offset=0);

cudaError_t cudaMemcpyFromSymbol(void* dst, const char* src, size_t count, size_t offset=0);

cudaError_t cudaMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind);

где dst оп­ре­де­ля­ет об­ласть па­мя­ти, ку­да ко­пи­ру­ет­ся сег­мент па­мя­ти, оп­ре­де­ляе­мый пе­ре­мен­ной src. В функ­ции cudaMemcpy сег­мен­ты па­мя­ти оп­ре­де­ля­ет­ся че­рез ука­за­те­ли. В функ­ци­ях же cudaMemcpyToSymbol и cudaMemcpyFromSymbol од­на из об­лас­тей яв­ля­ет­ся пе­ре­мен­ной или мас­си­вом, оп­ре­де­лен­ным с по­мо­щью клю­че­вых слов __device__ или __constant__. Функ­ция cudaMemcpy2D ко­пи­ру­ет дву­мер­ный мас­сив, па­мять под ко­то­рый бы­ла вы­де­ле­на с по­мо­щью вы­зо­ва cudaMallocPitch.

Че­рез ар­гу­мент count ука­зы­ва­ет­ся ко­ли­че­­ст­во ко­пи­руе­мых байт, а че­рез offset – сме­щение от­но­си­тель­но на­ча­ла. По­следний ар­гу­мент kind ука­зы­ва­ет на­прав­ление ко­пи­ро­вания, воз­мож­ные зна­чения ко­то­ро­го:

» cudaMemcpyDeviceToHost

» cudaMemcpyHostToDevice

» cudaMemcpyDeviceToDevice

где Device – па­мять гра­фи­че­­ско­­го уско­ри­те­ля, а Host – па­мять ком­пь­ю­те­ра.

[править] Ожи­дание за­вер­шения

За­пуск ядер осу­ще­ст­в­ля­ет­ся в асин­хрон­ном ре­жи­ме. Это оз­на­ча­ет, что по­сле по­ста­нов­ки яд­ра драй­ве­ром в оче­редь на за­пуск управ­ление тут же воз­вра­ща­ет­ся в про­грам­му. То есть функ­ция cudaMalloc, вы­зван­ная сра­зу по­сле вы­зо­ва яд­ра, ско­пи­ру­ет не те дан­ные, ко­то­рые ожи­да­ют­ся, по­сколь­ку яд­ро не за­вер­ши­ло ра­бо­ту. Что­бы до­ж­дать­ся ис­полнения всех за­пу­щен­ных ядер, необ­хо­ди­мо восполь­зо­вать­ся функ­ци­ей cudaDeviceSynchronize:

int main(){

float *dev_a, *dev_b;

float *host_a;

host_a = (float*)malloc(sizeof(float) * 1024*256);

cudaMalloc( &dev_a, sizeof(float)*1024*256 );

cudaMalloc( &dev_b, sizeof(float)*1024*256 );

cudaMemcpy( dev_a, host_a, sizeof(float)*1024*256, cudaMemcpyHostToDevice );

MyKernel <<< 1024, 256 >>>( dev_a, dev_b );

cudaDeviceSynchronize();

cudaMemcpy( host_a, dev_a, sizeof(float)*1024*256, cudaMemcpyDeviceToHost );

cudaFree( dev_a );

cudaFree( dev_b );

free( host_a );

}

[править] Про­вер­ка оши­бок

Лю­бая функ­ция CUDA Runtime воз­вра­ща­ет зна­чение ти­па cudaError_t, сиг­на­ли­зи­рую­щее об успеш­но­сти вы­полнения. Функ­ция cudaGetLastError воз­вра­ща­ет ре­зуль­тат по­следней вы­зван­ной функ­ции CUDA Runtime или яд­ра. Функ­ция cudaGetErrorString воз­вра­ща­ет стро­ку, рас­шиф­ро­вы­ваю­щую ошиб­ку.

Для удоб­ст­ва мож­но оп­ре­де­лить сле­дую­щий мак­рос про­вер­ки оши­бок:

1. define cudaCheck {

cudaError_t err = cudaGetLastError();

if ( err != cudaSuccess ){

printf(« cudaError = '%s' \n in '%s' %d\n»,\ cudaGetErrorString( err ),\ __FILE__, __LINE__ );

exit(0);\

}

}

Этот мак­рос про­ве­ря­ет, бы­ла ли ошиб­ка, а ес­ли бы­ла, то вы­во­дит на пе­чать опи­сание, имя фай­ла и но­мер стро­ки, где возник­ла ошиб­ка. За счет про­вер­ки ка­ж­до­го вы­зо­ва функ­ций CUDA Runtime и ядер уда­ет­ся най­ти боль­шую часть оши­бок на ранних ста­ди­ях раз­ра­бот­ки про­грамм.

[править] Вре­мя ра­бо­ты?

Из­ме­рение вре­мени ра­бо­ты яд­ра с по­мо­щью сис­тем­ных функ­ций не точ­но, по­сколь­ку оно вклю­ча­ет вре­мя, за­тра­чи­вае­мое на взаи­мо­дей­ст­вие про­грам­мы с драй­ве­ром. Для точ­но­го из­ме­рения вре­мени ра­бо­ты ядер оп­ре­де­ле­ны спе­ци­аль­ный тип со­бы­тия cudaEvent_t и функ­ции ра­бо­ты с ним:

cudaError_t cudaEventCreate (cudaEvent_t *event)

cudaError_t cudaEventRecord (cudaEvent_t event)

cudaError_t cudaEventSynchronize (cudaEvent_t event)

cudaError_t cudaEventElapsedTime (float *ms, cudaEvent_t start, cudaEvent_t end)

Здесь cudaEventCreate инициа­ли­зи­ру­ет пе­ре­мен­ную ти­па cudaEvent_t, cudaEventRecord уста­нав­ли­ва­ет со­бы­тие, cudaEventSynchronize до­жи­да­ет­ся за­вер­шения со­бы­тия, а cudaEventElapsedTime из­ме­ря­ет вре­мя ме­ж­ду дву­мя со­бы­тия­ми в мил­ли­се­кун­дах. Функ­ция cudaEventElapsedTime вы­да­ет ошиб­ку, ес­ли со­бы­тия start или end не за­вер­ши­лись.

При­мер про­грам­мы с из­ме­ре­ни­ем вре­ме­ни ра­бо­ты яд­ра:

cudaEvent_t start, end;

float time;

cudaEventCreate( &start );

cudaEventCreate( &end );

cudaEventRecord( start );

MyKernel<<< 1024, 256 >>>();

• Метамодернизм в позднем творчестве В.Г. Сорокина
• ЛитРПГ - последняя отрыжка постмодерна
• "Ричард III и семиотика"
• 3D-визуализация обложки Ridero создаем обложку книги при работе над самиздатом.
• Архитектура метамодерна - говоря о современном искусстве, невозможно не поговорить об архитектуре. В данной статье будет отмечено несколько интересных принципов, характерных для построек "новой волны", столь притягательных и скандальных.
• Литература
• Метамодерн
• Рокер-Прометей против изначального зла в «Песне про советскую милицию» Вени Дркина, Автор: Нина Ищенко, к.ф.н, член Союза Писателей ЛНР - перепубликация из журнала "Топос".
• Как избавиться от комаров? Лучшие типы ловушек.
• Что делать если роблокс вылетает на windows
• Что делать, если ребенок смотрит порно?
• Почему собака прыгает на людей при встрече?
• Какое масло лить в Задний дифференциал (мост) Visco diff 38434AA050
• О чем может рассказать хвост вашей кошки?
• Верветки
• Отчетность бюджетных учреждений при закупках по Закону № 223-ФЗ
• Срок исковой давности как правильно рассчитать
• Дмитрий Патрушев минсельхоз будет ли преемником Путина
• Кто такой Владислав Поздняков? Что такое "Мужское Государство" и почему его признали экстремистским в России?
• Как правильно выбрать машинное масло в Димитровграде?
• Как стать богатым и знаменитым в России?
• Почему фильм "Пипец" (Kick-Ass) стал популярен по всему миру?
• Как стать мудрецом?
• Как правильно установить FreeBSD
• Как стать таким как Путин?
• Где лучше жить - в Димитровграде или в Ульяновске?
• Почему город Димитровград так называется?
• Что такое метамодерн?
• ВАЖНО! Временное ограничение движения автотранспортных средств в Димитровграде
• Тарифы на электроэнергию для майнеров предложено повысить

cudaEventRecord( end );

cudaEventSynchronize( end );

cudaEventElapsedTime( &time, start, end );

printf(«Time: %.2f\n», time/1000.0 );

[править] За­клю­че­ние

Это лишь ма­лая часть техниче­­ской ин­фор­ма­ции, ко­то­рая по­тре­бу­ет­ся на пу­ти изу­чения CUDA. С дру­гой сто­ро­ны, поль­зо­ва­тель­ский ин­тер­фейс – это не ра­кет­ные тех­но­ло­гии, и его вполне мож­но осво­ить за ра­зум­ный про­ме­жу­ток вре­мени. |