физическа памет на AMD устройства: локална срещу частна

Пиша алгоритъм в OpenCL, в който ще имам нужда всяка работна единица да запомни справедлива част от данните, да кажем нещо между long[70] и long[200] или така на ядро.

Последните устройства на AMD имат 32 KiB __local памет, което е (за дадено количество данни на ядро) достатъчно за съхраняване на информация за 20-58 работни единици. Въпреки това, от това, което разбирам от архитектурата (и особено от този чертеж), всяко шейдърно ядро ​​също има специално количество частна памет. Но не успявам да намеря размера му.

Може ли някой да ми каже как да разбера колко частна памет има всяко ядро?

Особено съм любопитен за HD7970, тъй като смятам да купя някои от тях скоро.

Редактиране: Проблемът е решен, отговорът е тук в приложение D.


amd-processor architecture gpu opencl gpgpu
person user1111929    schedule 17.02.2012    source източник
comment
Не вярвам, че частната памет е предназначена за ядро ​​- тя се съпоставя с регистърния файл, който е ресурс на изчислителна единица. Всеки работен елемент получава регистри, разпределени от регистърния файл на изчислителната единица, колко са необходими определя броя на вълновите фронтове в полет във всеки даден момент.   -  person talonmies    schedule 17.02.2012
comment
От известния навсякъде виждан чертеж codeproject.com/KB/showcase/Memory -Spaces/image001.jpg Заключих, че личната памет е физически различна от __локалната памет, нали?   -  person user1111929    schedule 18.02.2012
comment
Да, физически са различни. Частната памет се преобразува към регистрационния файл на изчислителната единица, локалната памет към споделената памет на ниво изчислителна единица в повечето съвременни AMD устройства. Няколко ранни графични процесора, съвместими с OpenCL, нямаха споделена памет и локалната памет беше просто SDRAM. Нито едно от тях не е на ядро ​​и колко използвате на работен елемент за частни и на работна група за локални ефекти, броят на едновременните вълнови фронтове, изпълнявани на изчислителна единица.   -  person talonmies    schedule 18.02.2012
comment
Добре. Тогава трябва да преформулирам въпроса си: колко голям е този регистрационен файл? Как да разберете размера му, като цяло или конкретно за HD7970.   -  person user1111929    schedule 18.02.2012
comment
Мисля, че не сте разбрали - частната памет е (както казва името) частна за всеки работен елемент. Но той се разпределя за всеки работен елемент от файла(ите) на регистъра на изчислителната единица, който действа като общ пул от ресурси за всички работни елементи, изпълнявани на дадена изчислителна единица. И съм почти сигурен, че компилаторът на AMD поставя твърдо ограничение от 256 регистъра на работна единица, независимо от размера на регистрационните файлове на GPU.   -  person talonmies    schedule 18.02.2012
comment
И какъв е размерът на 1 регистър? 64 бита? Ако е така, това е твърдо ограничение от 2 KB на работен елемент, което е доста голямо, нали? Предполагам, че трябва да е много по-малък (в противен случай проблемът ми е тривиално решен, тъй като може да съдържа long[200] чисто в регистъра).   -  person user1111929    schedule 18.02.2012
comment
Мисля, че всеки регистър е 32-битова дума. Но не забравяйте, че всички други променливи във вашия код също консумират регистри. Мисля, че си спомням, че типичните графични процесори на AMD имат 64kb регистрационен файл на изчислителна единица, който трябва да бъде споделен от минимум 4 или 8 вълнови фронта от 64 работни елемента всеки. Но аз не използвам техния хардуер много, така че това може да не е правилно. Проверете текущите бележки за изданието в техния OpenCL SDK.   -  person talonmies    schedule 18.02.2012
comment
Наистина, ето го, благодаря! Намира се в Приложение D на ръководството за програмиране на AMD APP OpenCL developer.amd.com/sdks /amdappsdk/assets/. Очевидно един регистър е 128 бита (4x32) и има 16384 за всички съвременни устройства от висок клас, така че това са забележителните 256KB на изчислителна единица. хубаво! Ако можете да поставите това в нов отговор, мога да го приема и да затворя темата.   -  person user1111929    schedule 19.02.2012

Отговори (3)


arrow_upward
4
arrow_downward

Отговорът беше даден от потребител talonmies в коментарите, така че ще го напиша в нов отговор тук, за да затворя въпроса.

Тези стойности могат да бъдат намерени в Приложение D на ръководството за програмиране на AMD APP OpenCL http://developer.amd.com/sdks/amdappsdk/assets/amd_accelerated_parallel_processing_opencl_programming_guide.pdf (съществува подобен документ за nVidia). Очевидно регистърът е 128 бита (4x32) за AMD устройства и има 16384 регистъра за всички съвременни устройства от висок клас, така че това са забележителните 256KB на изчислителна единица.

person user1111929    schedule 01.03.2012

arrow_upward
0
arrow_downward

Мисля, че търсите __локална памет. Това е, за което се отнасят 32KB локално съхранение на данни. Не мисля, че можете да анкетирате устройството, за да получите обема на личната памет.

Можете да подадете NULL long* cl_mem препратка, за да разпределите паметта. Мисля, че е най-добре да използвате статичен обем памет на WI. Ако приемем, че long[200] ще се изисква за всеки работен елемент, ще използвате кода по-долу. Също така би било добра идея да разделите работата на групи, които имат еднакви (или подобни) изисквания към паметта, за да извлечете максимума от LDS паметта.

void __kernel(__local long* localMem, const int localMemPerItem
       //more args...
       )
{
  //host has 'passed' localMemPerItem*get_local_size() long values in as locamMem
  //this work item has access to all of it, but can choose to restrict
  //itself to only the portion it needs.
  //work group size will be limited to CL_DEVICE_LOCAL_MEM_SIZE/(8*localMemPerItem)
  int startIndex=localMemPerItem*get_local_id(0);
  //use localMem[startIndex+ ... ]
}
person mfa    schedule 17.02.2012
comment
Не можете да го анкетирате, но съществува ли? От известния навсякъде виждан чертеж codeproject.com/KB/showcase/Memory -Spaces/image001.jpg Предположих, че има физически отделен набор от частни регистри за всяка работна единица. Не? Надявах се по някакъв начин да се справя по-добре от ограничение CL_DEVICE_LOCAL_MEM_SIZE/(8*localMemPerItem), тъй като грубо оставя половината от ядрата неизползвани. Достъпът до глобалната памет вероятно ще бъде твърде бавен, въпреки че само увеличава брояч. - person user1111929; 18.02.2012
comment
Намерих още информация за размерите на регистрите на кипарис, кайман и ферми тук: realworldtech. com/page.cfm?ArticleID=RWT121410213827&p=11 Трябва да можете да настроите частни променливи с приличен размер в този размер. Мисля, че LDS все пак ще бъде най-добрият ви залог. - person mfa; 19.02.2012

arrow_upward
0
arrow_downward

За да отговоря колко голям е регистрационният файл в карта от серия 79xx, тъй като е базиран на GCN архитектура, той е 64KB според изображението в anandtech: http://www.anandtech.com/print/5261

За да отговорите на въпроса си как да разберете колко памет използва всяко ядро.. можете да погледнете, стартирайте AMD APP Profiler на вашето ядро, той ви казва в раздела за заетост на ядрото колко място се използва от ядрото.

person kiranputtur    schedule 20.02.2012
comment
Наистина ли? Това е странно. Мислех, че съм намерил отговора, но той е различен. В ръководството за програмиране на AMD OpenCL developer.amd.com/sdks/amdappsdk/assets/ в Приложение D има общия размер на регистрационния файл и той е посочен като 256 KB за всички съвременни устройства. Кое е правилно сега? :С - person user1111929; 20.02.2012
comment
Вярвам, че и двете са правилни. Доколкото разбирам, в архитектурата на GCN една SIMD единица има 64kb регистрационен файл и има 4 SIMD единици на изчислителна единица, т.е. 4 * 64kb = 256kb от общия регистрационен файл на изчислителна единица. - person talonmies; 20.02.2012