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

Question

Я пишу алгоритм в OpenCL, в котором мне нужно, чтобы каждый рабочий блок запоминал значительную часть данных, например, от long[70] до long[200] или около того на ядро.

Последние устройства AMD имеют 32 КБ __local памяти, что (для данного объема данных на ядро) достаточно для хранения информации для 20-58 рабочих единиц. Однако, насколько я понимаю из архитектуры (и особенно из этого чертежа ), каждое ядро ​​шейдера также имеет выделенный объем частной памяти. Я, однако, не могу найти его размер.

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

Мне особенно интересно узнать о HD7970, так как я планирую купить некоторые из них в ближайшее время.

Редактировать: проблема решена, ответ здесь в приложении D.

Answer 1

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

Эти значения можно найти в Приложении D Руководства по программированию AMD APP OpenCL http://developer.amd.com/sdks/amdappsdk/assets/amd_accelerated_parallel_processing_opencl_programming_guide.pdf (аналогичный документ существует для nVidia). Очевидно, что для устройств AMD размер регистра составляет 128 бит (4x32), а для всех современных устройств высокого класса - 16384 регистра, так что это замечательные 256 КБ на вычислительную единицу.

Answer 2

Чтобы ответить, насколько велик файл регистра на карте серии 79xx, поскольку он основан на архитектуре GCN, он составляет 64 КБ согласно изображению в anandtech: http://www.anandtech.com/print/5261

Чтобы ответить на ваш вопрос, как узнать, каксколько памяти использует каждое ядро. Вы можете посмотреть, как запускается AMD APP Profiler в вашем ядре, в разделе о занятости ядра указано, сколько места используется ядром.

Answer 3

Я думаю, вы ищете локальную память.Это то, на что ссылается 32 КБ локального хранилища данных.Я не думаю, что вы можете опросить устройство, чтобы получить объем приватной памяти.

Вы можете передать 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+ ... ]
}