Как использовать async_work_group_copy в OpenCL?

Я хотел бы понять, как правильно использовать вызов async_work_group_copy() в OpenCL. Давайте рассмотрим упрощенный пример:

__kernel void test(__global float *x) {
  __local xcopy[GROUP_SIZE];

  int globalid = get_global_id(0);
  int localid = get_local_id(0);
  event_t e = async_work_group_copy(xcopy, x+globalid-localid, GROUP_SIZE, 0);
  wait_group_events(1, &e);
}

Ссылка http://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/async_work_group_copy.html говорит: "Выполните асинхронную копию элементов gentype num_elements от src до dst. Асинхронная копия выполняется всеми рабочими элементами в работе -группу, и поэтому эта встроенная функция должна быть встречена всеми рабочими элементами в рабочей группе, выполняющей ядро ​​с теми же значениями аргумента, в противном случае результаты будут undefined."

Но это не уточняет мои вопросы...

Я хотел бы знать, если следующие предположения верны:

  • Вызов async_work_group_copy() должен выполняться всеми рабочими элементами в группе.
  • Вызов должен быть таким образом, чтобы исходный адрес был идентичным для всех рабочих элементов и указывал на первый элемент области памяти, который нужно скопировать.
  • Поскольку мой исходный адрес является относительным на основе глобального идентификатора рабочего элемента первого рабочего элемента рабочей группы. Поэтому я должен вычесть локальный идентификатор, чтобы адрес был идентичным для всех рабочих элементов...
  • Действительно ли третий параметр представляет собой число элементов (не размер в байтах)?

Бонусные вопросы:

а. Можно ли использовать барьер (CLK_LOCAL_MEM_FENCE) вместо wait_group_events() и игнорировать возвращаемое значение? Если это так, возможно, это будет быстрее?

б. Локальная копия также имеет смысл для обработки на процессорах или это накладные расходы, поскольку они все равно используют кеш?

С уважением, Стефан

Ответы

Ответ 1

Одна из основных причин существования этой функции - позволить компилятору driver/kernel эффективно копировать память без необходимости разработки разработчиками предположений об оборудовании.

Вы описываете, какую память вам нужно скопировать, как если бы это была однопоточная копия, а async_work_group_copy делает это для вас, используя параллельное оборудование.

По вашим конкретным вопросам:

  • Я никогда не видел async_work_group_copy, который использовался только некоторыми рабочими элементами в группе. Я всегда предполагал, что это необходимо. Я думаю, что блокирующий характер wait_group_events заставляет все рабочие элементы быть частью копии.

  • Да. Адреса источника (и назначения) должны быть одинаковыми для всех рабочих элементов.

  • Вы можете вычесть свой локальный идентификатор, чтобы получить правильный адрес, но я обнаружил, что, основывая адрес на groupId, также решает эту проблему. (Get_group_id)

  • Да. Последним параметром является количество элементов, а не размер в байтах.

а. Нет. На основе событий вы обнаружите, что ваш барьер почти сразу поражается рабочими элементами, и данные не обязательно будут скопированы. Это имеет смысл, потому что некоторые аппаратные средства opencl могут даже не использовать вычислительные устройства вообще для выполнения операции копирования.

б. Я думаю, что реализация cpu opencl может гарантировать использование кеша L1 при использовании локальной памяти. Единственный способ узнать наверняка, если это лучше работает, - это сравнить ваше приложение с различными настройками.