Что более эффективно в SYCL - использовать один буфер или несколько буферов?
Предположим, у меня есть массив данных, например массив трехмерных векторов размера N. Предположим, что каждая итерация моего ядра SYCL исключительно или в первую очередь касается только одного вектора. Какой из следующих способов разбить это на непрерывные буферы, как правило, более эффективен - или это имеет значение?
Я понимаю, что целевое устройство сильно влияет на это, поэтому давайте предположим, что это дискретный графический процессор (т.е. данные действительно нужно копировать на другой чип памяти, и у устройства нет какой-то сумасшедшей архитектуры, такой как FPGA- я в основном нацелена на GTX 1080 через CUDA, но я ожидаю, что ответ, скорее всего, будет аналогичным, когда код компилируется в OpenCL или мы используем другой современный графический процессор.
- Создайте отдельный буфер для каждой координаты, например
sycl::buffer<float> x, y, z;
, каждый размером N. Таким образом, при доступе к ним я могу использоватьsycl::id<1>
передается в мое ядро лямбда как индекс без математики. (Я подозреваю, что компилятор сможет это оптимизировать.) - Создайте один упакованный буфер для всех, например
sycl::buffer<float> coords;
с размером 3н. При доступе к ним с помощьюsycl::id<1>
называетсяi
, Затем я беру координату x какbuffer_accessor[3*i]
, координата y какbuffer_accessor[3*i+1]
, а координата z какbuffer_accessor[3*i+2]
. (Я не знаю, может ли компилятор это оптимизировать, и я не уверен, могут ли возникнуть проблемы с выравниванием.) - Создайте один распакованный буфер, используя структуру, например
struct Coord { float x,y,z; }; sycl::buffer<Coord> coords;
. Это имеет довольно тревожную цену, связанную с увеличением использования памяти, в данном примере на 33%, из-за заполнения выравнивания, что также увеличивает время, необходимое для копирования буфера на устройство. Но компромисс в том, что вы можете получить доступ к данным, не манипулируяsycl::id<1>
, среда выполнения должна иметь дело только с одним буфером, и на устройстве не должно быть никаких проблем с выравниванием строк кэша. - Используйте двумерный буфер размера (N,3) и выполняйте итерацию только в диапазоне первого измерения. Это менее гибкое решение, и я не понимаю, почему я хочу использовать многомерные буферы, когда я не повторяю все измерения, если только для этого варианта использования не встроена большая оптимизация.
Я не могу найти никаких рекомендаций по архитектуре данных, чтобы получить интуитивное представление о подобных вещах. Прямо сейчас (4) кажется глупым, (3) включает недопустимую трату памяти, и я использую (2), но задаюсь вопросом, не следует ли мне использовать (1) вместо этого, чтобы избежать манипуляций с идентификатором, и 3*sizeof(float) выровненные блоки доступа.
1 ответ
Для шаблонов доступа к памяти на графических процессорах сначала важно понять концепцию объединения. По сути, это означает, что при определенных условиях устройство объединяет обращения к памяти соседних рабочих элементов и вместо этого производит один большой доступ к памяти. Это очень важно для производительности. Подробные требования, когда происходит объединение, различаются между поставщиками графических процессоров (или даже между поколениями графических процессоров одного производителя). Но обычно требования, как правило, соответствуют
- Определенное количество смежных рабочих элементов обращается к смежным элементам данных. Например, все рабочие элементы в подгруппе 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, в общем, есть как преимущества, так и недостатки использования одного большого буфера по сравнению с несколькими меньшими (например, накладные расходы многих буферов, но больше возможностей для стратегий оптимизации графа задач). Я ожидаю, что эти эффекты будут вторичными.