Шаблон доступа к глобальной памяти OpenCL (AMD GCN) для векторизованных данных: шаг за шагом и непрерывность

Question

Я собираюсь улучшить производительность ядра OCL и хочу уточнить, как работают транзакции памяти и какой шаблон доступа к памяти действительно лучше (и почему). Ядро снабжено векторами из 8 целых чисел, которые определены как массив: int v [8], это означает, что перед выполнением любого вычисления весь вектор должен быть загружен в GPR. Итак, я считаю, что узким местом этого кода является начальная загрузка данных.

Сначала рассмотрим некоторые основы теории.

Целевым HW является Radeon RX 480/580, который имеет 256-битную шину памяти GDDR5, на которой транзакция пакетного чтения / записи имеет степень детализации 8 слов, следовательно, одна транзакция памяти читает 2048 бит или 256 байтов. Я полагаю, что CL_DEVICE_MEM_BASE_ADDR_ALIGN относится к:

Alignment (bits) of base address: 2048.

Итак, мой первый вопрос: каков физический смысл 128-байтовой кеш-линии? Сохраняет ли он часть данных, извлекаемую при однократном чтении, но не запрашиваемую? Что будет с остальными, если мы запросим, ​​скажем, 32 или 64 байта - таким образом, остаток превышает размер строки кэша? (Полагаю, это будет просто отброшено - тогда какая часть: голова, хвост...?)

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

Но, тем не менее, я могу разместить свои данные с двумя различными схемами доступа:

1) смежный

    a[i] = v[get_global_id(0) * get_global_size(0) + i];

(который фактически выполнен как)

    *(int8*)a = *(int8*)v;

2) чередование

    a[i] = v[get_global_id(0) + i * get_global_size(0)];

Я ожидаю, что в моем случае смежность будет быстрее, потому что, как сказано выше, одна транзакция памяти может полностью заполнить 8 рабочих элементов данными. Однако я не знаю, как физически работает планировщик в вычислительном блоке: нужны ли ему все данные, чтобы быть готовыми для всех линий SIMD, или достаточно первой части для 4 параллельных элементов SIMD? Тем не менее, я полагаю, что он достаточно умен, чтобы вначале полностью обеспечить данными по крайней мере один CU, как только CU могут выполнять потоки команд независимо. В то время как во втором случае нам нужно выполнить 8 * global_size / 64 транзакций, чтобы получить полный вектор.

Итак, мой второй вопрос: верно ли мое предположение?

Теперь практика.

На самом деле, я разделил всю задачу на два ядра, потому что одна часть имеет меньшее регистровое давление, чем другая, и поэтому может использовать больше рабочих элементов. Итак, сначала я поиграл с шаблоном, как данные, хранящиеся при переходе между ядрами (используя vload8/vstore8 или приведение к int8, дают один и тот же результат), и результат был несколько странным: ядро, которое читает данные непрерывно, работает примерно на 10% быстрее (как в CodeXL и измерением времени ОС), но ядро, которое непрерывно хранит данные, работает на удивление медленнее. Общее время для двух ядер примерно одинаково. В моих мыслях оба должны вести себя по крайней мере одинаково - либо медленнее, либо быстрее, но эти обратные результаты казались необъяснимыми.

И третий вопрос: кто-нибудь может объяснить такой результат? Или может я что то не так делаю? (Или совершенно не так?)

Answer 1

Ну, на самом деле я не ответил на все мои вопросы, но некоторая информация, обнаруженная в просторах интернета, объясняет все более понятным образом, по крайней мере, для меня (в отличие от вышеупомянутого руководства по оптимизации AMD, которое кажется неясным и иногда сбивает с толку):

"Аппаратное обеспечение выполняет некоторое слияние, но это сложно...
доступ к памяти в деформации не обязательно должен быть непрерывным, но имеет значение, во сколько 32-байтовых сегментов глобальной памяти (и 128-байтовых сегментов кэша l1) они попадают. контроллер памяти может загрузить 1, 2 или 4 из этих 32-байтовых сегментов за одну транзакцию, но это читается через кэш в 128-байтовых строках.
таким образом, если каждая дорожка в деформации загружает случайное слово в диапазоне 128 байт, то штраф не применяется; это 1 транзакция и чтение на полную эффективность. но если каждая дорожка в деформации загружает 4 байта с шагом 128 байт, это очень плохо: загружается 4096 байт, но используется только 128, что дает эффективность ~3% ".

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

Answer 2

Взгляните на главу 2.1 Руководства по оптимизации AMD OpenCL. Он ориентирован в основном на карты старого поколения, но архитектура GCN не изменилась полностью, поэтому все равно должна применяться к вашему устройству (polaris).

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

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

С уважением,

Майкл