Что более эффективно в SYCL: использовать один буфер или несколько буферов?

Предположим, у меня есть массив данных, например, массив трехмерных векторов размера N. Предположим, что каждая итерация моего ядра SYCL исключительно или в основном связана только с одним вектором. Какой из следующих способов разбить это на непрерывные буферы, как правило, более эффективен — или это имеет значение?

Я понимаю, что целевое устройство сильно влияет на это, поэтому давайте предположим, что это дискретный графический процессор (т. е. данные действительно должны быть скопированы на другой чип памяти, и устройство не имеет какой-то сумасшедшей архитектуры, такой как FPGA — я в основном нацелен на GTX 1080 через CUDA, но я ожидаю, что ответ, вероятно, будет аналогичным, когда код компилируется в OpenCL или мы используем другой современный графический процессор.

  1. Создайте отдельный буфер для каждой координаты, например. sycl::buffer<float> x, y, z;, каждый размером N. Таким образом, при доступе к ним я могу использовать sycl::id<1>, переданный моему лямбда-выражению ядра, в качестве индекса без математики. (Я подозреваю, что компилятор может оптимизировать это.)
  2. Создайте один упакованный буфер для всех из них, например. sycl::buffer<float> coords; размером 3N. При доступе к ним с помощью sycl::id<1>, называемого i, я получаю координату x как buffer_accessor[3*i], координату y как buffer_accessor[3*i+1] и координату z как buffer_accessor[3*i+2]. (Я не знаю, может ли компилятор оптимизировать это, и я не уверен, что могут возникнуть проблемы с выравниванием.)
  3. Создайте один распакованный буфер, используя структуру, например. struct Coord { float x,y,z; }; sycl::buffer<Coord> coords;. Это имеет довольно тревожную стоимость увеличения использования памяти, в этом примере на 33%, из-за заполнения выравнивания, что также увеличивает время, необходимое для копирования буфера на устройство. Но компромисс заключается в том, что вы можете получить доступ к данным, не манипулируя sycl::id<1>, среда выполнения должна иметь дело только с одним буфером, и на устройстве не должно быть неэффективного выравнивания строк кэша.
  4. Используйте двумерный буфер размера (N,3) и выполняйте итерации только в диапазоне первого измерения. Это менее гибкое решение, и я не понимаю, зачем мне использовать многомерные буферы, если я не перебираю все измерения, если только для этого варианта использования не встроено много оптимизации.

Я не могу найти какие-либо рекомендации по архитектуре данных, чтобы получить представление о таких вещах. Прямо сейчас (4) кажется глупым, (3) включает неприемлемую трату памяти, и я использую (2), но задаюсь вопросом, не следует ли мне вместо этого использовать (1), чтобы избежать манипуляций с идентификатором и 3 * sizeof (float) выровненные фрагменты доступа.


c++ memory-management sycl dpc++
person sapphous    schedule 11.11.2020    source источник

Ответы (1)


arrow_upward
4
arrow_downward

Для шаблонов доступа к памяти на графических процессорах в первую очередь важно понять концепцию объединения. В основном это означает, что при определенных условиях устройство будет объединять обращения к памяти соседних рабочих элементов и вместо этого выдавать одно большое обращение к памяти. Это очень важно для производительности. Подробные требования при объединении различаются в зависимости от поставщика графических процессоров (или даже между поколениями графических процессоров одного поставщика). Но обычно требования, как правило, совпадают с

  • Определенное количество смежных рабочих элементов обращается к смежным элементам данных. Например. все рабочие элементы в подгруппе SYCL / CUDA деформируют доступ к последующим элементам данных.
  • Элемент данных, к которому обращается первый рабочий элемент, возможно, придется выровнять, например. к строке кэша.

См. здесь объяснение для (старых) графических процессоров NVIDIA: https://developer.nvidia.com/blog/how-access-global-memory-efficiently-cuda-c-kernels/

Имея это в виду, 3) тратится не только объем памяти, но и пропускная способность памяти, и если у вас есть что-то вроде my_accessor[id].x, у вас есть доступ к памяти с шагом, который предотвращает объединение.

Для 4), я не уверен, правильно ли я понимаю. Я предполагаю, что вы имеете в виду, что измерение с 3 элементами контролирует доступ к x/y/z, а измерение с N описывает n-й вектор. В этом случае это будет зависеть от того, есть ли у вас размер (N, 3) или (3, N). Поскольку в SYCL расположение данных таково, что последний индекс всегда самый быстрый, (N, 3) на практике будет соответствовать 3) без проблем с заполнением. (3, N) будет похоже на 2), но без пошагового доступа к памяти (см. ниже)

Для 2) основная проблема с производительностью заключается в том, что вы выполняете пошаговый доступ к памяти, если x находится в [3*i], y в [3*i+1] и т. д. Вместо этого для объединения вы хотите, чтобы x был в [i], y в [N+i] и z в [2N+i]. Если у вас есть что-то вроде

float my_x = data[i]; // all N work items perform coalesced access for x
float my_y = data[i+N];
float my_z = data[i+2N];

У вас хороший шаблон доступа к памяти. В зависимости от вашего выбора N и требований к выравниванию для объединенного доступа к памяти вашего устройства у вас могут возникнуть проблемы с производительностью для y и z из-за выравнивания.

Я не ожидаю, что тот факт, что вам нужно добавить смещения к вашему индексу, существенно повлияет на производительность.

Для 1) вы в основном получите гарантию того, что все данные хорошо выровнены и что доступ будет объединен. Из-за этого я ожидаю, что это будет работать лучше всего из представленных подходов.

С точки зрения среды выполнения SYCL в целом есть как преимущества, так и недостатки использования одного большого буфера по сравнению с несколькими меньшими (например, накладные расходы на множество буферов, но больше возможностей для стратегий оптимизации графа задач). Я ожидаю, что эти эффекты будут вторичными.

person illuhad    schedule 12.11.2020