Объединение памяти и векторизованный доступ к памяти - PullRequest
2 голосов
/ 10 июля 2019

Я пытаюсь понять связь между объединением памяти на графических процессорах NVIDIA / CUDA и векторизованным доступом к памяти на x86-SSE / C ++.

Насколько я понимаю, что:

  • Объединение памяти - это оптимизация времени выполнения контроллера памяти (реализована аппаратно).Сколько транзакций памяти требуется для выполнения загрузки / сохранения деформации, определяется во время выполнения.Командой загрузки / сохранения деформации может быть , которая выдается повторно , если только не происходит идеальное слияние.
  • Векторизация памяти - это оптимизация во время компиляции .Количество транзакций памяти для векторизованной загрузки / сохранения является фиксированным.Каждая векторная инструкция загрузки / сохранения выдается ровно один раз.
  • Команды загрузки / сохранения в графическом процессоре GPU более выразительны, чем инструкции векторной загрузки / сохранения SSE.Например, команда st.global.s32 PTX может храниться в 32 произвольных местах памяти (размер деформации 32), тогда как инструкция movdqa SSE может храниться только в последовательном блоке памяти.
  • Объединение памяти в CUDA, кажется, гарантирует эффективный векторизованный доступ к памяти (когда доступы коалесцируемы), тогда как в x86-SSE мы должны надеяться, что компилятор фактически векторизирует код (он может этого не сделать) илиВекторизация кода вручную с использованием встроенных функций SSE, что более сложно для программистов.

Это правильно?Я пропустил важный аспект (возможно, маскирование потоков)?

Теперь, почему графические процессоры объединяются во время выполнения?Это, вероятно, требует дополнительных аппаратных схем.Каковы основные преимущества перед объединением во время компиляции, как в процессорах?Существуют ли шаблоны доступа к приложениям / памяти, которые сложнее реализовать на процессорах из-за отсутствия объединения во время выполнения?

Ответы [ 2 ]

3 голосов
/ 10 июля 2019

предостережение: я не очень хорошо знаю / не понимаю архитектуру / микроархитектуру графических процессоров.Часть этого понимания объединена с вопросом + что другие люди написали здесь в комментариях / ответах.

Способ, которым графические процессоры позволяют одной инструкции работать с несколькими данными, очень отличается от CPU SIMD,Вот почему им вообще нужна особая поддержка объединения памяти.CPU-SIMD не может быть запрограммирован так, как это необходимо.

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


Да, объединение памяти в основном делает во время выполнения то, что делает SIMD с коротким векторным процессором во время компиляции, в пределах одного "ядра". Эквивалентом CPU-SIMD были бы сборка / разброс загрузки / хранения, которые могли бы оптимизировать для единого широкого доступа к кешу для смежных индексов. Существующие ЦП этого не делают: каждый элемент обращается к кешу отдельно всобираются.Не следует использовать загрузку сбора, если вы знаете, что многие индексы будут смежными;это будет быстрее, чтобы перетасовать 128-битные или 256-битные блоки на место.В общем случае, когда все ваши данные являются смежными, вы просто используете обычную векторную загрузочную инструкцию вместо групповой загрузки.

Смысл современного коротко-векторного CPU SIMD состоит в том, чтобы выполнять больше работы черезизвлекать / декодировать / выполнять конвейер без , что делает его более широким с точки зрения необходимости декодировать + отслеживать + выполнять больше инструкций ЦП за такт. Быстрое расширение конвейера ЦП приводит к уменьшению отдачи в большинстве случаев использованияпотому что большая часть кода не имеет большого количества ILP.

Универсальный ЦП тратит много транзисторов на механизмы планирования команд / выполнения не по порядку, поэтому просто расширяя его, чтобы иметь возможностьзапускать много больше мопов параллельно не жизнеспособно.(https://electronics.stackexchange.com/questions/443186/why-not-make-one-big-cpu-core).

Чтобы увеличить пропускную способность, мы можем повысить частоту, повысить IPC и использовать SIMD для выполнения большей работы по каждой инструкции / операции, которую должен отслеживать механизм, вышедший из строя. (И мыможет собрать несколько ядер на одном чипе, но кеш-когерентные соединения между ними + кэш-память L3 + контроллеры памяти трудны). Современные процессоры используют все эти вещи, поэтому мы получаем общую пропускную способность по частоте * IPC * SIMD и разколичество ядер, если мы многопоточны. Они не являются жизнеспособными альтернативами друг другу, это ортогональные вещи, которые вы должны сделать все из, чтобы управлять большим количеством FLOP или целочисленной работой через конвейер ЦП.

Именно поэтому ЦП SIMD имеет широкие исполнительные блоки фиксированной ширины вместо отдельной инструкции для каждой скалярной операции. Не существует механизма для гибкой подачи одной скалярной инструкции в несколько исполнительных блоков.

Чтобы воспользоваться этим, необходима векторизация во время компиляции, не только ваших загрузок / хранилищ, но и вашего ALUвычисление.Если ваши данные не являются смежными, вы должны собрать их в векторы SIMD либо со скалярными нагрузками + перемешиваниями, либо с помощью AVX2 / AVX512 - собрать нагрузки, которые принимают базовый адрес + вектор (масштабированных) индексов.


Но GPU SIMD отличается .Это для массовых параллельных задач, когда вы делаете то же самое с каждым элементом.«Конвейер» может быть очень легковесным, потому что ему не нужно поддерживать exec-order exec или переименование регистров, особенно ветвления и исключения.Это делает возможным иметь только скалярные исполнительные блоки без необходимости обрабатывать данные в фиксированных фрагментах с непрерывных адресов.

Это две очень разные модели программирования.Они оба SIMD, но детали аппаратного обеспечения, на котором они работают, очень разные.


Каждая векторная инструкция загрузки / сохранения выдается ровно один раз.

Да, это логично. На практике внутренние органы могут быть немного более сложными, например, AMD Ryzen разбивает 256-битные векторные операции на 128-битные половины, или Intel Sandybridge / IvB делает это только для загрузок + хранилищ, имея FPU шириной 256-бит.

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

В терминологии Intel моп для разделения нагрузки получает отправлено дважды, но выдает только + удаляется один раз.

Выравниваемые загрузки / хранилища, такие как movdqa или movdqu, когда память выровнена во время выполнения, представляют собой всего лишь один доступ к кэш-памяти L1d (при условии попадания в кэш). Если вы не используете процессор, который декодирует векторную инструкцию в две половины, например AMD для 256-битных векторов.


Но это все внутри ядра ЦП для доступа к кешу L1d. CPU <-> транзакции памяти выполняются целыми строками кеша, с частными кешами обратной записи L1d / L2 и общим L3 на современных процессорах x86 - Какой метод отображения кеша используется в процессоре Intel Core i7 ? (Intel со времен Nehalem, начало серии i3 / i5 / i7, AMD со времен Bulldozer, я думаю, представил для них кэш-память L3.)

В CPU это кэш L1d с обратной записью, который в основном объединяет транзакции в целые строки кэша, независимо от того, используете вы SIMD или нет.

SIMD помогает выполнять больше работы внутри ЦП, чтобы не отставать от более быстрой памяти. Или для проблем, когда данные помещаются в кэш-память L2 или L1d, чтобы быстро действительно справиться с этими данными.

0 голосов
/ 10 июля 2019

Объединение памяти относится к parallel доступам: когда каждое ядро ​​в СМ получит доступ к следующей ячейке памяти, доступ к памяти оптимизируется.

И наоборот, SIMD - это одноядерная оптимизация: когда векторный регистр заполнен операндами и выполняется операция SSE, параллелизм находится внутри ядра ЦП, при этом одна операция выполняется на каждом внутреннем логическом блоке за тактцикл.

Тем не менее, вы правы: доступ коалесцированной / не коалесцированной памяти является аспектом времени выполнения.SIMD-операции компилируются. Я не думаю, что они могут хорошо сравниваться.

Если бы я проводил параллелизм, я бы сравнил коалесцирование в графических процессорах с памятью предварительная выборка в процессорах.Это также очень важная оптимизация во время выполнения - и я считаю, что она активна и за кулисами, используя SSE.

Однако нет ничего похожего на объединение в процессорных ядрах Intel.Из-за когерентности кэша лучшее, что вы можете сделать при оптимизации параллельного доступа к памяти, - это разрешить каждому ядру доступ к независимым областям памяти.

Теперь, почему графические процессоры объединяются во время выполнения?

Графическая обработка оптимизирована для параллельного выполнения одной задачи на соседних элементах.

Например, подумайте о том, чтобы выполнить операцию с каждым пикселем изображения, назначив каждый пиксель разному ядру.Теперь ясно, что вы хотите иметь оптимальный путь для загрузки изображения, растягивающего по одному пикселю на каждое ядро.

Вот почему объединение памяти глубоко скрыто в архитектуре графических процессоров.

...