LXF170:Паралельные вы­чис­ления: CUDA

Подписка на печатную версию Весь 2015 год (12 номеров) Первое полугодие (6 номеров) Второе полугодие (6 номеров)

Подписка на электронную версию Весь 2015 год (12 номеров) Первое полугодие (6 номеров) Второе полугодие (6 номеров)

Подшивки старых номеров журнала (печатные версии) Весь 2014 год (12 номеров) Первое полугодие (6 номеров) Второе полугодие (6 номеров)

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

Кон­стан­тин Кал­гин и Ев­ге­ний Бал­дин не мо­гут спо­кой­но ви­деть мат­ри­цу – им не­пре­мен­но на­до ее транс­по­ни­ро­вать.

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

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

Хо­чет­ся вол­шеб­ст­ва... вот бы ска­зать ком­пь­ю­те­ру: «По­счи­тай-ка, го­луб­чик, мою за­да­чу, да по­бы­ст­рее! А я ко­фе по­пью...» Ска­зать-то мож­но, да со­вре­мен­ный ком­пь­ю­тер сей­час на это не ­реа­ги­ру­ет. На­вер­ное, это непло­хо: про­снувшись, ма­шин­ный ра­зум нач­нет вся­че­­ски от­лынивать от ра­бо­ты. И че­ло­ве­ку все равно при­дет­ся раз­би­рать­ся с ап­па­ра­ту­рой.

Вер­сии CUDA

Сле­ду­ет раз­ли­чать вер­сии про­грамм­ных па­ке­тов CUDA Toolkit (ком­пи­ля­тор, про­фи­ли­ров­щик, биб­лио­те­ки), CUDA SDK (при­ме­ры с ис­ход­ны­ми ко­да­ми, вспо­мо­га­тель­ные биб­лио­те­ки), CUDA Driver (сис­тем­ный драй­вер) и вер­сии CUDA, ко­то­рую под­дер­жи­ва­ет гра­фи­че­­ский уско­ри­тель. В до­ку­мен­та­ции Nvidia C Programming Guide вер­сия, ко­то­рую под­дер­жи­ва­ет гра­фи­че­­ский уско­ри­тель, на­зы­ва­ет­ся Compute capability. Для про­сто­ты на­зо­вем это «вер­си­ей CUDA».

Вер­сия CUDA со­сто­ит из двух чи­сел – стар­ше­го и млад­ше­го, на­при­мер 1.3. Гра­фи­че­­ские уско­ри­те­ли с одинаковым стар­шим чис­лом име­ют од­ну ар­хи­тек­ту­ру яд­ра. Млад­шее чис­ло ука­зы­ва­ет на улуч­шения в ар­хи­тек­ту­ре яд­ра. Ка­кой вер­сии со­от­вет­ст­ву­ет ва­ше уст­рой­ст­во, лег­ко уз­нать из той же Ви­ки­пе­дии, в ста­тье CUDA.

Ар­хи­тек­ту­ра про­цес­со­ра

Про­цес­сор гра­фи­че­­ско­­го уско­ри­те­ля со­сто­ит из планиров­щи­ка бло­ков по­то­ков, на­бо­ра муль­ти­про­цес­со­ров и кэ­ша L2. На­ли­чие и объ­ем кэ­ша L2 за­ви­сит от вер­сии CUDA, а ко­ли­че­­ст­во муль­ти­про­цес­со­ров – еще и от мо­де­ли гра­фи­че­­ско­­го уско­ри­те­ля. Ме­няя ко­ли­че­­ст­во муль­ти­про­цес­со­ров, про­из­во­ди­те­ли про­пор­цио­наль­но ме­ня­ют по­треб­ляе­мую мощ­ность и про­из­во­ди­тель­ность гра­фи­че­­ско­­го уско­ри­те­ля, и непро­пор­цио­наль­но его це­ну.

» Планиров­щик бло­ков по­то­ков При за­пуске яд­ра мож­но от­пра­вить на ис­полнение до 65 535 бло­ков по­то­ков. Планиров­щик сле­дит за за­гру­жен­но­стью муль­ти­про­цес­со­ров, где ис­пол­ня­ют­ся бло­ки по­то­ков и по за­вер­шении ра­бо­ты одних бло­ков от­прав­ля­ет еще не от­ра­бо­тан­ные бло­ки на осво­бо­див­шие­ся муль­ти­про­цес­со­ры.

» Муль­ти­про­цес­сор со­сто­ит из на­бо­ра по­то­ко­вых про­цес­со­ров, планиров­щи­ка по­то­ков, раз­де­ляе­мой па­мя­ти, бан­ка ре­ги­ст­ров, а так­же тек­стур­но­го, кон­стант­но­го и L1 кэ­шей.

» Кэш L2 ав­то­ма­ти­че­­ски кэ­ши­ру­ет дан­ные при об­ра­щении к гло­баль­ной па­мя­ти, чем уско­ряет как по­сле­дую­щий по­втор­ный доступ, так и доступ к со­седним дан­ным. Кэш L2 поя­вил­ся в ар­хи­тек­ту­ре CUDA, на­чи­ная с вер­сии 2.x (Fermi).

Планиров­щик бло­ков по­то­ков

» CUDA 1.x По­ка не ис­полнились все бло­ки по­то­ков од­но­го яд­ра, бло­ки по­то­ков дру­го­го бу­дут планиро­вать­ся. По су­ти, это не планиро­вание, а про­стая раз­да­ча ра­бо­ты (бло­ков по­то­ков) в по­ряд­ке оче­ред­но­сти (ко­ор­ди­нат) и мо­но­поль­ное ис­поль­зо­вание ре­сур­сов гра­фи­че­­ско­­го уско­ри­те­ля одним ядром.

» CUDA 2.x (Fermi) Планиров­щик мо­жет планиро­вать бло­ки по­то­ков от раз­ных яд­ер од­но­го про­цес­са, по­вы­ша­я эф­фек­тив­ность уско­ри­те­ля на неболь­ших сет­ках в несколь­ко бло­ков по­то­ков.

» CUDA 3.5 (Kepler) Реа­ли­зо­ван ди­на­ми­че­­ский па­рал­ле­лизм, по­зво­ляю­щий за­пускать яд­ра с са­мо­го гра­фи­че­­ско­­го уско­ри­те­ля и син­хронизи­ро­вать­ся по ре­зуль­та­ту. Вы­зо­вы ядер мо­гут быть вло­жен­ны­ми, что­ от­кры­ва­ет боль­шие воз­мож­но­сти эф­фек­тив­ной реа­ли­за­ции за­дач с нере­гу­ляр­ным и ди­на­ми­че­­ским па­рал­ле­лиз­мом, а так­же по пе­реводу час­ти ко­да, управ­ляю­ще­го за­пуском ядер, на гра­фи­че­­ский уско­ри­тель.

Муль­ти­про­цес­сор

Ко­ли­че­­ст­во и объ­ем тех или иных эле­мен­тов мультипроцессора за­ви­сит от вер­сии CUDA. На од­ном муль­ти­про­цес­со­ре мо­гут планиро­вать­ся несколь­ко бло­ков по­то­ков. Минималь­ной единицей ис­полнения и планиро­вания на муль­ти­про­цес­со­ре яв­ля­ет­ся варп [англ. warp – зд. скрут­ка] – груп­па из 32 по­то­ков од­но­го бло­ка. На ка­ж­дом так­те планиров­щик вы­би­ра­ет груп­пу по­то­ков, и над ка­ж­дым по­то­ком из груп­пы ис­пол­ня­ет­ся од­на и та же ко­ман­да. Кор­рект­ная об­ра­бот­ка услов­ных пе­ре­хо­дов по­то­ка­ми од­но­го вар­па про­ис­хо­дит за счет то­го, что неко­то­рые по­то­ки/по­то­ко­вые про­цес­со­ры мо­гут про­стаи­вать, то есть не ис­пол­нять те­ку­щую ко­ман­ду. Ито­го­вое вре­мя об­ра­бот­ки вет­вей услов­но­го пе­ре­хо­да в слу­чае, когда про­изош­ло раз­де­ление по­то­ков вар­па, скла­ды­ва­ет­ся из вре­мен ис­полнения обе­их вет­вей.

Имею­щие­ся у муль­ти­про­цес­со­ра 4-бай­то­вые ре­ги­ст­ры де­лят­ся ме­ж­ду планируе­мы­ми по­то­ка­ми, а объ­ем раз­де­ляе­мой па­мя­ти – ме­ж­ду планируе­мы­ми бло­ка­ми по­то­ков.

» CUDA 1.x Со­дер­жит 8 по­то­ко­вых про­цес­со­ров, вы­пол­няю­щих ин­ст­рук­ции с це­лы­ми чис­ла­ми и чис­ла­ми с пла­ваю­щей точ­кой оди­нар­ной точ­но­сти [float]. Вы­чис­ления с пла­ваю­щей за­пя­той с двой­ной точ­но­стью [double] ста­ли доступ­ны в CUDA 1.3 – на од­ном муль­ти­про­цес­со­ре на­хо­ди­лось толь­ко од­но ис­полнитель­ное уст­рой­ст­во, т. е. ско­рость ра­бо­ты с двой­ной точ­но­стью бы­ло в 8 раз мень­ше ско­ро­сти ра­бо­ты с оди­нар­ной точ­но­стью.

» СUDA 2.x Со­дер­жит 32 (CUDA 2.0) или 48 (CUDA 2.1) по­то­ко­вых про­цес­со­ров, вы­пол­няю­щих ин­ст­рук­ции с це­лы­ми чис­ла­ми и чис­ла­ми с пла­ваю­щей точ­кой оди­нар­ной точ­но­сти. Уве­ли­чи­лось ко­ли­че­­ст­во ис­полнитель­ных уст­ройств для ра­бо­ты с чис­ла­ми двой­ной точ­но­сти – 16 (CUDA 2.0) и 24 (CUDA 2.1). Ко­ли­че­­ст­во планиров­щи­ков вар­пов – 2. В CUDA 2.1 на ка­ж­дом так­те ка­ж­дый планиров­щик вы­да­ет по две ин­фор­ма­ци­он­но неза­ви­си­мых ин­ст­рук­ции од­но­го из вар­пов, ес­ли толь­ко ни од­на из ин­ст­рук­ций не ра­бо­та­ет с чис­ла­ми двой­ной точ­но­сти.

В СUDA 2.x поя­вил­ся «пол­но­цен­ный» кэш L1 и L2 – ранее все кэ­ши бы­ли доступ­ны толь­ко на чтение, т. е. кэ­ши­ро­ва­ли кон­стант­ные дан­ные. У ка­ж­до­го муль­ти­про­цес­со­ра кэш пер­во­го уров­ня L1 свой, а кэш вто­ро­го уров­ня L2 об­щий для всех муль­ти­про­цес­со­ров гра­фи­че­­ско­­го уско­ри­те­ля. Объ­е­мы кэ­ша пер­во­го уров­ня и раз­де­ляе­мой па­мя­ти в сум­ме да­ют 64 KБ, и мо­гут быть скон­фи­гу­ри­ро­ва­ны с по­мо­щью cudaFuncSetCacheConfig() в 48 KБ/16 KБ или 16 KБ/48 KБ, со­от­вет­ст­вен­но. Объ­е­мы кэ­ша L1 и раз­де­ляе­мой па­мя­ти по умол­чанию – 16 KБ/48 KБ. Объ­ем кэ­ша L2 – 768 KБ.

» CUDA 3.x Со­дер­жит 192 по­то­ко­вых про­цес­со­ра, вы­пол­няю­щих ин­ст­рук­ции с це­лы­ми чис­ла­ми и чис­ла­ми с пла­ваю­щей точ­кой оди­нар­ной точ­но­сти. Ко­ли­че­­ст­во ис­полнитель­ных уст­ройств для ра­бо­ты с чис­ла­ми двой­ной точ­но­сти – 8 (CUDA 3.0) и 64 (CUDA 3.5). Планиров­щи­ки усо­вер­шен­ст­во­ва­ны та­ким об­ра­зом, что те­перь мо­гут вы­да­вать по две ин­ст­рук­ции за такт вне за­ви­си­мо­сти от ти­па ин­ст­рук­ций. Объ­е­мы кэ­ша L1 и раз­де­ляе­мой па­мя­ти те­перь мо­гут быть скон­фи­гу­ри­ро­ва­ны как 32 KБ/32 KБ. Объ­ем и про­пу­ск­ная спо­соб­ность кэ­ша L2 вы­росли в два раза.

На ар­хи­тек­ту­рах CUDA 1.x и 2.x, что­бы об­ме­нять­ся зна­чения­ми ре­ги­ст­ров ме­ж­ду по­то­ка­ми од­но­го бло­ка, необ­хо­ди­мо бы­ло ис­поль­зо­вать раз­де­ляе­мую па­мять. В ар­хи­тек­ту­ре CUDA 3.x поя­ви­лись ин­ст­рук­ции, по­зво­ляю­щие об­менивать­ся зна­чения­ми ре­ги­ст­ров ме­ж­ду по­то­ка­ми од­но­го вар­па без ис­поль­зо­вания раз­де­ляе­мой па­мя­ти. За счет это­го эко­но­мит­ся вре­мя об­ра­щения и объ­ем ис­поль­зуе­мой раз­де­ляе­мой па­мя­ти.

» Ор­ганиза­ция па­мя­ти В гра­фи­че­­ском уско­ри­те­ле ие­рар­хию па­мя­ти со­став­ля­ют сле­дую­щие эле­мен­ты: файл ре­ги­ст­ров, раз­де­ляе­мая па­мять, кэ­ши и гло­баль­ная па­мять. В офи­ци­аль­ной до­ку­мен­та­ции ко­ли­че­­ст­вен­ные дан­ные по ла­тент­но­сти при­во­дят­ся толь­ко для гло­баль­ной па­мя­ти. Ос­таль­ные дан­ные – ка­че­­ст­вен­но­го ха­рак­те­ра: ла­тент­ность ре­ги­ст­ров рав­на ла­тент­но­сти раз­де­ляе­мой па­мя­ти, ла­тент­ность кэ­шей мень­ше ла­тент­но­сти гло­баль­ной па­мя­ти, ла­тент­ность со­гла­со­ван­но­го досту­па в раз­де­ляе­мую/гло­баль­ную па­мять су­ще­ст­вен­но мень­ше ла­тент­но­сти несо­гла­со­ван­но­го досту­па.

» Локаль­ная па­мять В ар­хи­тек­ту­ре CUDA име­ет­ся ап­па­рат­ное ог­раничение на ко­ли­че­­ст­во ис­поль­зуе­мых ре­ги­ст­ров одним по­то­ков – 63 (CUDA < 3.0) и 255 (CUDA 3.5). Ес­ли во вре­мя ком­пи­ля­ции ком­пи­ля­то­ру не хва­тит доступ­ных ре­ги­ст­ров, он ото­бра­зит их на локаль­ную па­мять. Локаль­ная па­мять – это об­ласть в гло­баль­ной па­мя­ти, вы­де­лен­ная ком­пи­ля­то­ром для хранения локаль­ных зна­чений по­то­ков. Она ис­поль­зу­ет­ся для хранения локаль­ных дан­ных по­то­ков при нехват­ке ре­ги­ст­ров или объ­яв­ления локаль­ных мас­си­вов внут­ри яд­ра без клю­че­во­го сло­ва __shared__:

__global__ void kernel(int *a, int *s) {

int l[N], res;

int ind = (blockIdx.x*blockDim.x+threadIdx.x)*N;

for( int j=0; j<N; j++) l[ j ] = a[ ind + j ];

for( int j=0; j<N; j++) res += l[ j ] * j;

s[ ind ] = res;

}

В этом при­ме­ре при боль­ших N мас­сив l бу­дет рас­по­ла­гать­ся в локаль­ной па­мя­ти. У ка­ж­до­го по­то­ка бу­дет свой мас­сив. При ма­лых N ком­пи­ля­тор мо­жет раз­вер­нуть цик­лы, по­сле че­го ото­бра­зить эле­мен­ты мас­си­ва на ре­ги­ст­ры, по­сколь­ку от­па­дет ну­ж­да в об­ра­щении к эле­мен­там по ме­няю­ще­му­ся ин­дек­су.

» Доступ в раз­де­ляе­мую па­мять. Кон­флик­ты Вся раз­де­ляе­мая па­мять де­лит­ся на 16 (CUDA 1.x) или 32 (CUDA 2.x/3.x) бан­ков. По­сле­до­ва­тель­но рас­по­ло­жен­ные 32-бит­ные сло­ва по­ме­ща­ют­ся в по­сле­до­ва­тель­ных бан­ках [interleaved]. Про­пу­ск­ная спо­соб­ность ка­ж­до­го бан­ка – 32 би­та за два так­та. Кон­флик­том на­зы­ва­ет­ся од­но­вре­мен­ное об­ра­щение по­то­ков к раз­ным 32-бит­ным сло­вам од­но­го бан­ка. Кон­фликт­ные об­ра­щения в банк ис­пол­ня­ют­ся по­сле­до­ва­тель­но. Об­ра­щение к раз­де­ляе­мой па­мя­ти на­зы­ва­ет­ся со­гла­со­ван­ным ес­ли от­сут­ст­ву­ют кон­флик­ты.

> CUDA 1.x Ин­ст­рук­ция об­ра­щения к раз­де­ляе­мой па­мя­ти ис­полня­ет­ся за два ша­га – по по­ло­ви­нам вар­па. Кон­флик­ты мо­гут возник­нуть толь­ко внутри ка­ж­дой из по­ло­вин вар­па.

> CUDA 2.x/3.x Кон­флик­ты мо­гут возник­нуть в рам­ках вар­па в це­лом.

» Доступ в гло­баль­ную па­мять

> CUDA 1.0/1.1 Доступ в гло­баль­ную па­мять яв­ля­ет­ся со­гла­со­ван­ным, ес­ли для ка­ж­дой по­ло­ви­ны вар­па вы­пол­ня­ют­ся сле­дую­щие усло­вия:

1 Раз­мер слов, к ко­то­рым об­ра­ща­ет­ся ка­ж­дый по­ток, ра­вен 4, 8 или 16 бай­там.

2 Ес­ли раз­мер ра­вен N, то все 16 слов ле­жат в 16×N-байт­ном сег­мен­те.

3 По­то­ки об­ра­ща­ют­ся к сло­вам по­сле­до­ва­тель­но: k-й по­ток в по­ло­вине вар­па об­ра­ща­ет­ся к k-му сло­ву в сег­мен­те.

Вто­рое усло­вие для при­клад­но­го про­грам­ми­ста пе­ре­фор­му­ли­ру­ет­ся сле­дую­щим об­ра­зом: при об­ра­щении к эле­мен­там мас­си­ва пер­вый по­ток ка­ж­дой по­ло­ви­ны вар­па дол­жен об­ра­щать­ся к эле­мен­ту, но­мер ко­то­ро­го кра­тен 16. При со­гла­со­ван­ном об­ра­щении к 4/8/16-байт­ным сло­вам для ка­ж­дой по­ло­ви­ны вар­па вы­пол­ня­ет­ся од­на 64-байт­ная/од­на 128-байт­ная/две 128-байт­ных тран­зак­ции. Не­вы­полнение ин­ст­рук­ции об­ра­щения к гло­баль­ной па­мя­ти неко­то­ры­ми по­то­ка­ми за счет ранее ис­полнен­но­го услов­но­го ветв­ления не влия­ет на со­гла­со­ван­ность.

В слу­чае невы­полнения усло­вий со­гла­со­ван­но­го досту­па об­ра­щение раз­би­ва­ет­ся на 16 от­дель­ных 32-байт­ных тран­зак­ций.

> CUDA 1.2/1.3 Для этой и по­сле­дую­щих ар­хи­тек­тур не ис­поль­зу­ет­ся тер­мин со­гла­со­ван­но­го об­ра­щения, но опи­сы­ва­ет­ся ал­го­ритм оп­ре­де­ления ко­ли­че­­ст­ва и раз­мер тран­зак­ций с гло­баль­ной па­мя­тью. Для CUDA 1.2/1.3 ко­ли­че­­ст­во и раз­мер тран­зак­ций оп­ре­де­ля­ет­ся сле­дую­щим об­ра­зом.

1 Для ка­ж­дой по­ло­ви­ны вар­па бе­рет­ся минималь­ное ко­ли­че­­ст­во сег­мен­тов, ко­то­рые по­кры­ва­ют все за­пра­ши­вае­мые эле­мен­ты этой по­ло­ви­ной. Раз­мер ка­ж­до­го сег­мен­та ра­вен 32 бай­там для 1 байт­ных дан­ных, 64 бай­там для 2 байт­ных, 128 бай­там для 4/8/16 байт­ных дан­ных.

2 Ка­ж­дый за­гру­жае­мый сег­мент умень­ша­ет­ся по пра­ви­лам:

» Ес­ли за­пра­ши­вае­мые дан­ные ле­жат толь­ко в ле­вой или пра­вой по­ло­вине 64/128-байт­но­го сег­мен­та, то сег­мент умень­ша­ет­ся до со­от­вет­ст­вую­щей 32/64-байт­ной по­ло­ви­ны; » Ес­ли за­пра­ши­вае­мые дан­ные ле­жат толь­ко в од­ной из чет­вер­тей 128-байт­но­го сег­мен­та, то сег­мент умень­ша­ет­ся до со­от­вет­ст­вую­щей 32-байт­ной час­ти; > CUDA 2.x/3.x Все об­ра­щения в гло­баль­ную па­мять кэ­ши­ру­ют­ся; где именно, оп­ре­де­ля­ет про­грам­мист во вре­мя ком­пи­ля­ции че­рез фла­ги: в кэ­шах L1 и L2 (-Xptxas -dlcm=ca, по умол­чанию) или толь­ко в кэ­ше L2 (-Xptxas -dlcm=cg). Ва­ри­ант кэ­ши­ро­вания оп­ре­де­ля­ет раз­мер тран­зак­ций с па­мя­тью – 128-байт­ные для пер­во­го слу­чая и 32-байт­ные для вто­ро­го. Та­ким об­ра­зом, кэ­ши­ро­вание толь­ко в L2 мо­жет со­кра­тить вре­мя об­ра­щения в гло­баль­ную па­мять в слу­чае, когда по­то­ки од­но­го вар­па об­ра­ща­ют­ся к раз­бро­сан­ным [scattered] дан­ным.

» Еди­ное ад­рес­ное про­стран­ст­во На­чи­ная с CUDA 2.x, реа­ли­зо­ва­но еди­ное ад­рес­ное про­стран­ст­во. Т. е. мно­же­ст­во ад­ре­сов по­де­ле­но на уча­ст­ки, со­от­вет­ст­вую­щие локаль­ной, раз­де­ляе­мой и гло­баль­ной па­мя­тям. Это су­ще­ст­вен­но уп­ро­ща­ет про­грам­ми­ро­вание ал­го­рит­мов с ад­ре­са­ци­ей, за­ви­ся­щей от дан­ных.

Транс­пониро­вание мат­риц

Для де­мон­ст­ра­ции неко­то­рых из осо­бен­но­стей ар­хи­тек­ту­ры ре­шим про­стую мо­дель­ную за­да­чу: транс­пониру­ем мат­ри­цу. На вход про­грам­ме по­да­ет­ся мат­ри­ца A раз­мером N×N. На вы­хо­де необходимо по­лу­чить мат­ри­цу B, та­кую, что Bi,j = Aj,i.

Тес­ти­ро­вание про­из­во­ди­тель­но­сти пред­ло­жен­ных ал­го­рит­мов бу­дет про­во­дить­ся на гра­фи­че­­ских уско­ри­те­лях Nvidia GeForce GTS 8800 (CUDA 1.1, 128 ядер, 512 МБ) и Nvidia Quadro FX 480 (CUDA 1.3, 192 яд­ра, 1536 МБ).

Для сравнения при­ве­дем лис­тинг по­сле­до­ва­тель­ной реа­ли­за­ции транс­пониро­вания:

void transpose_host(float *a, float *b, int N) {

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

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

b[ j*N + i ] = a[ i*N + j ];

}

}

}

Ва­ри­ант 1

Чис­ло по­ро­ж­дае­мых вы­чис­ли­тель­ных по­то­ков рав­но чис­лу эле­мен­тов мат­ри­цы. По­ток с ко­ор­ди­на­та­ми (i,j):

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

j = threadIdx.y+blockIdx.y * blockDim.y

ко­пи­ру­ет зна­чение эле­мен­та (i,j) из мат­ри­цы A в эле­мент (j,i) мат­ри­цы B. При та­ком под­хо­де по­то­ки од­но­го вар­па чи­та­ют зна­чения из max(32/blockDim.x,1) строк мат­ри­цы, что при­во­дит к та­ко­му же чис­лу тран­зак­ций с па­мя­тью на чтение. При этом по­то­ки од­но­го вар­па за­пи­сы­ва­ют эти зна­чения в min(blockDim.x,32) строк, что при­во­дит к та­ко­му же чис­лу тран­зак­ций с па­мя­тью на запись. Та­ким об­ра­зом, при уве­ли­чении blockDim.x умень­ша­ет­ся чис­ло тран­зак­ций на чтение, но уве­ли­чи­ва­ет­ся чис­ло тран­зак­ций на запись. При умень­шении – на­обо­рот. Пол­ный лис­тинг про­грам­мы:

int N; // matrix size NxN

int BSX, BSY;

__global__ void transpose_1(float* A, float* B, int N) {

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

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

B[ j * N + i ] = A[ i * N + j ];

}

float *host_a, *host_b;

float *dev_a, *dev_b;

int main(){

N = 1024;

BSX = BSY = 16;

host_a = (float*)malloc(sizeof(float)*N*N);

host_b = (float*)malloc(sizeof(float)*N*N);

cudaMalloc(&dev_a, sizeof(float)*N*N);

cudaMalloc(&dev_b, sizeof(float)*N*N);

cudaMemcpy(dev_a, host_a, sizeof(float)*N*N, cudaMemcpyHostToDevice);

dim3 gdim = dim3(N/BSX,N/BSY,1);

dim3 bdim = dim3(BSX,BSY,1);

transpose_1 <<< gdim, bdim>>> (dev_a, dev_b, N);

cudaDeviceSynchronize();

cudaMemcpy(host_b, dev_b, sizeof(float)*N*N, cudaMemcpyDeviceToHost);

}

Здесь в стро­ках опи­са­но яд­ро transpose_1, транс­понирую­щее мат­ри­цу A, ре­зуль­тат за­пи­сы­ва­ет­ся в мат­ри­цу B.

Да­лее оп­ре­де­ля­ют­ся ука­за­те­ли на мат­ри­цы host_a и host_b, ко­то­рые рас­по­ла­га­ют­ся в опе­ра­тив­ной па­мя­ти ком­пь­ю­те­ра, и dev_a, dev_b – в гло­баль­ной па­мя­ти гра­фи­че­­ско­­го уско­ри­те­ля. Са­ми ука­за­те­ли бу­дут хранить­ся в опе­ра­тив­ной па­мя­ти, по­сколь­ку они инициа­ли­зи­ру­ют­ся в основ­ной про­грам­ме и в их оп­ре­де­лении от­сут­ст­ву­ют __device__ и __constant__. По­сле это­го вы­де­ля­ет­ся об­ласть в опе­ра­тив­ной (malloc) и гло­баль­ной (cudaMalloc) па­мя­ти.

Для вы­зо­ва яд­ра transpose_1 необ­хо­ди­мо оп­ре­де­лить две пе­ре­мен­ные струк­тур­но­го ти­па dim3, со­дер­жа­щие в се­бе раз­ме­ры бло­ка по­то­ков и мас­си­ва бло­ков по­то­ков. Ко­ли­че­­ст­во по­ро­ж­дае­мых по­то­ков рав­но N2 = (N/BSX)*(N/BSY)*BSX*BSY. Яд­ру в ка­че­­ст­ве ар­гу­мен­тов пе­ре­да­ют­ся ука­за­те­ли на мат­ри­цы и линей­ный раз­мер са­мих мат­риц. По­сле вы­зо­ва яд­ра основ­ная про­грам­ма до­жи­да­ет­ся за­вер­шения его ис­полнения (cudaDeviceSynchronize), что­бы убе­диться, что зна­чения всех по­то­ков за­пи­са­ны в па­мять, и ко­пи­ру­ет дан­ные из гло­баль­ной па­мя­ти (cudaMemcpyDeviceToHost).

Ва­ри­ант 2

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

Возь­мем то же чис­ло по­ро­ж­дае­мых по­то­ков и их ото­бра­жение на эле­мен­ты, как в Ва­ри­ан­те 1. По­то­ки од­но­го бло­ка ра­бо­та­ют со сле­дую­щи­ми эле­мен­та­ми мат­ри­цы A: {(x+tx,y+ty)| x = blockDim.x*blockIdx.x, y = blockDim.y*blockIdx.y, 0 ≤ tx < blockDim.x, 0 ≤ ty < blockDim.y}. Та­ким об­ра­зом, по­то­ки од­но­го бло­ка транс­пониру­ют под­мат­ри­цу раз­ме­ра blockDim.x*blockDim.y и ре­зуль­тат за­пи­сы­ва­ют в мат­ри­цу B, на­чи­ная с эле­мен­та (y,x). Под­мат­ри­ца доста­точ­но неболь­шая и мо­жет быть раз­ме­ще­на в раз­де­ляе­мой па­мя­ти, где ста­дия транс­пониро­вания бу­дет ис­пол­нять­ся бы­ст­рее.Там мы и бу­дем про­во­дить транс­пониро­вание под­мат­ри­цы:

1 За­гру­жае­мый из гло­баль­ной па­мя­ти эле­мент (i,j) = (x+threadIdx.x, y+threadIdx.y) мат­ри­цы A за­пи­сы­ва­ет­ся в эле­мент (threadIdx.x, threadIdx.y) до­полнитель­но­го мас­си­ва sh, рас­по­ло­жен­но­го в раз­де­ляе­мой па­мя­ти.

2 Про­ис­хо­дит барь­ер­ная син­хрониза­ция по­то­ков в бло­ке, что­бы быть уве­рен­ным, что все по­то­ки за­гру­зи­ли свое зна­чение в мас­сив sh.

3 По­ток (i,j) за­пи­сы­ва­ет зна­чение эле­мен­та (threadIdx.y,threadIdx.x) мас­си­ва sh в эле­мент (y+threadIdx.x,x+threadIdx.y) мас­си­ва B.

Дан­ный спо­соб бу­дет ра­бо­тать толь­ко с квад­рат­ны­ми бло­ка­ми по­то­ков, на­при­мер, 16 × 16 или 32 × 32. Вот но­вый лис­тинг яд­ра:

__global__ void transpose_2(float* a, float* b, int N) {

__shared__ float sh[ BSY ][ BSX ];

int x = blockIdx.x * blockDim.x;

int y = blockIdx.y * blockDim.y;

int i = x + threadIdx.x;

int j = y + threadIdx.y;

sh[ threadIdx.y ][ threadIdx.x ] = a[ (y+threadIdx.y) * N + (x+threadIdx.x) ];

__syncthreads();

b[ (x+threadIdx.y) * N + (y+threadIdx.x) ] = sh[ threadIdx.x ][ threadIdx.y ];

}

Ва­ри­ант 3

За счет ис­поль­зо­вания раз­де­ляе­мой па­мя­ти по­лу­чи­лось миними­зи­ро­вать чис­ло тран­зак­ций с гло­баль­ной па­мя­тью. Од­на­ко в дан­ной реа­ли­за­ции есть неболь­шой недоста­ток: в по­следней строч­ке во вре­мя чтения из мас­си­ва sh на всех вер­си­ях CUDA об­ра­зу­ют­ся кон­флик­ты бан­ков па­мя­ти. Рас­смот­рим ва­ри­ант с раз­ме­ром бло­ка по­то­ков 16 × 16 и 16 бан­ка­ми в раз­де­ляе­мой па­мя­ти (CUDA 1.x). Ад­рес эле­мен­та sh[threadIdx.x][threadIdx.y] ра­вен (void*) sh+4*(threadIdx.x*16 + threadIdx.y), сле­до­ва­тель­но, эле­мент sh[threadIdx.x][threadIdx.y] раз­ме­ща­ет­ся в бан­ке, но­мер ко­то­ро­го ра­вен остат­ку от де­ления (threadIdx.x*16+threadIdx.y) на 16. В дан­ном слу­чае но­мер бу­дет ра­вен threadIdx.y. По­это­му все по­то­ки ка­ж­дой из по­ло­вин вар­па бу­дут об­ра­щать­ся к од­но­му и то­му же бан­ку. Та­кое об­ра­щение при­ве­дет к уве­ли­чению вре­мени ис­полнения со­от­вет­ст­вую­щей ин­ст­рук­ции в 16 раз. Что­бы из­бе­жать кон­флик­тов по бан­кам, в этой реа­ли­за­ции доста­точ­но уве­ли­чить раз­мер строк до 17 эле­мен­тов:

__global__ void transpose_4(float* a, float* b, int N) {

__shared__ sh[ BSY ][ BSX+1 ];

…

Это обес­пе­чит от­сут­ст­вие кон­флик­тов, по­сколь­ку ад­ре­са эле­мен­тов sh[threadIdx.x][threadIdx.y] бу­дут рав­ны (threadIdx.x*17 + threadIdx.y)%16, и для раз­лич­ных threadIdx.x из (0,1, ..., 15) эти но­ме­ра бу­дут раз­ли­чать­ся.

Ре­зуль­та­ты

По ре­зуль­та­там тес­ти­ро­вания вид­но, что применение раз­де­ляе­мой па­мя­ти су­ще­ст­вен­но умень­шает вре­мя ра­бо­ты яд­ра, т. к. миними­зи­ру­ет­ся ко­ли­че­­ст­во тран­зак­ций c па­мя­тью – как на чтение, так и на запись. Ва­ри­ант 3, уст­ранив кон­флик­ты по бан­кам, да­ет еще несколь­ко про­цен­тов вы­иг­ры­ша по вре­мени ра­бо­ты.

Эф­фек­тив­ное про­грам­ми­ро­вание

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

» Ко­ли­че­­ст­во вы­чис­ли­тель­ных по­то­ков От раз­ме­ра сет­ки бло­ков по­то­ков и са­мо­го бло­ка по­то­ков за­ви­сит сте­пень за­гру­жен­но­сти планиров­щи­ков гра­фи­че­­ско­­го уско­ри­те­ля. Чем боль­ше вар­пов планиру­ет планиров­щик, тем боль­ше у него воз­мож­но­стей скрыть за­держ­ки, свя­зан­ные с об­ра­щением в гло­баль­ную па­мять. Кро­ме то­го, доста­точ­но боль­шие раз­ме­ры сет­ки и бло­ка обес­пе­чат эф­фек­тив­ное мас­шта­би­ро­вание на но­вых гра­фи­че­­ских уско­ри­те­лях без пе­ре­пи­сы­вания и пе­ре­ком­пи­ли­ро­вания про­грам­мы. Обыч­но ре­ко­мен­ду­ет­ся по­ро­ж­дать по­ряд­ка 105 по­то­ков за один за­пуск яд­ра.

» Рав­но­мер­ность за­груз­ки вы­чис­ли­тель­ных по­то­ков Не­рав­но­мер­ность таковой – очень час­тое яв­ление; на­при­мер, вы­чис­ления зна­чений во внут­ренних и гранич­ных уз­лах сет­ки яв­ной раз­но­ст­ной схе­мы обыч­но раз­ли­ча­ют­ся. Од­на­ко нерав­но­мер­ность за­груз­ки по­то­ков мо­жет при­вес­ти к су­ще­ст­вен­ной де­гра­да­ции про­из­во­ди­тель­но­сти. Во-пер­вых, это свя­за­но с тем, что блок по­то­ков за­вер­ша­ет­ся и осво­бо­ж­да­ет муль­ти­про­цес­сор, когда все его по­то­ки за­вер­ши­ли ис­полнение. Во-вто­рых, за­пуск яд­ра за­вер­ша­ет­ся, когда все бло­ки по­то­ков за­вер­ши­ли ис­полнение.

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

» Локаль­ность за­гру­жае­мых дан­ных Этот фак­тор на­пря­мую сле­ду­ет из пре­ды­ду­ще­го. Уве­ли­чение локаль­но­сти за­гру­жае­мых дан­ных на уровне вар­па по­зво­ля­ет умень­шить ко­ли­че­­ст­во и объ­ем тран­зак­ций с па­мя­тью и по­вы­сить эф­фек­тив­ность ра­бо­ты кэ­шей. Уве­ли­чение локаль­но­сти за­гру­жае­мых дан­ных на уровне бло­ка по­то­ков по­зво­ля­ет уве­ли­чить эф­фек­тив­ность ис­поль­зо­вания раз­де­ляе­мой па­мя­ти и ра­бо­ты кэ­шей.

» Де­ление вар­пов на услов­ных пе­ре­хо­дах Когда раз­ные по­то­ки од­но­го вар­па раз­би­ва­ют­ся по раз­ным вет­вям услов­но­го пе­ре­хо­да, вре­мя ис­полнения услов­но­го пе­ре­хо­да скла­ды­ва­ет­ся из вре­мен ис­полнения его вет­вей. Т. е. час­тое де­ление вар­пов по вет­вям при­во­дит к де­гра­да­ции про­из­во­ди­тель­но­сти; ее сте­пень за­ви­сит от числа де­лений вар­пов и от раз­ме­ра вет­вей пе­ре­хо­да. |

Об­рат­ная связь

При­гла­ша­ем вы­ска­зать­ся по­тен­ци­аль­ных ав­то­ров ста­тей по па­рал­лель­ным вы­чис­лениям – цен­ные пред­ло­жения, кри­ти­ку и со­ве­ты при­сы­лай­те по элек­трон­ной поч­те: kalgin@ssd.sscc.ru, E.M.Baldin@inp.nsk.su.