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

Question

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

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

1. Создайте отдельный буфер для каждой координаты, например sycl::buffer<float> x, y, z;, каждый размером N. Таким образом, при доступе к ним я могу использовать sycl::id<1>передается в мое ядро ​​лямбда как индекс без математики. (Я подозреваю, что компилятор сможет это оптимизировать.)
2. Создайте один упакованный буфер для всех, например 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]. (Я не знаю, может ли компилятор это оптимизировать, и я не уверен, могут ли возникнуть проблемы с выравниванием.)
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) выровненные блоки доступа.

Answer 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, в общем, есть как преимущества, так и недостатки использования одного большого буфера по сравнению с несколькими меньшими (например, накладные расходы многих буферов, но больше возможностей для стратегий оптимизации графа задач). Я ожидаю, что эти эффекты будут вторичными.