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

Версия 06:53, 17 ноября 2018

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

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

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

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

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

Вер­сии 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 по­то­ков за один за­пуск яд­ра.

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

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

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

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