Блог

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

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

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

Смотрите Здесь объяснение (более старых) графических процессоров 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-й вектор. В этом случае это будет зависеть от того, есть ли у вас size (N, 3) или (3, N) . Поскольку в SYCL расположение данных таково, что последний индекс всегда является самым быстрым, (N, 3) на практике это соответствовало бы 3) без проблемы с заполнением. (3, N) было бы похоже на 2), но без расширенного доступа к памяти (см. Ниже)

Для 2) основная проблема с производительностью заключается в том, что вы выполняете пошаговый доступ к памяти, если x находится в [3*i] , y в [3*i 1] и т.д. Для объединения вы вместо этого хотите, чтобы x было at [i] , y at [N i] и z at [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, в целом, есть как преимущества, так и недостатки использования одного большого буфера по сравнению с несколькими меньшими (например, накладные расходы на множество буферов, но больше возможностей для стратегий оптимизации графика задач). Я ожидаю, что эти эффекты будут вторичными.