Как использовать 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 при использовании локальной памяти. Единственный способ узнать наверняка, если это лучше работает, - это сравнить ваше приложение с различными настройками.