В устройствах CUDA объединение операций записи в глобальную память так же важно, как объединение операций чтения глобальной памяти? Если да, то как это можно объяснить? Также есть ли различия между ранними поколениями устройств CUDA и самыми последними в отношении этой проблемы?
Объединение памяти при глобальных операциях записи
comment
Проблема слияния широко обсуждается в Руководстве по программированию CUDA C (раздел 5.3.2) и Руководстве по лучшим практикам CUDA C (раздел 9.2.1). Два руководства охватывают также проблему объединения для разных архитектур. Чтобы не дублировать материал, было бы более конструктивно, если бы вы ознакомились с этими документами и разместили неясные и требующие уточнения моменты.
- person Vitality schedule 25.11.2013
Ответы (2)
Объединенные записи (или их отсутствие) могут влиять на производительность, так же как и объединенные чтения (или их отсутствие).
Объединенное чтение происходит, когда запрос на чтение инициируется командой деформации, например:
int i = my_int_data[threadIdx.x+blockDim.x*blockIdx.x];
может быть удовлетворено одной транзакцией read в контроллере памяти (что, по сути, означает, что все операции чтения отдельных потоков происходят из одной строки кэша).
Объединенная запись происходит, когда запрос на запись инициируется командой деформации, например:
my_int_data[threadIdx.x+blockDim.x*blockIdx.x] = i;
может быть выполнено одной транзакцией write в контроллере памяти.
Для приведенных выше примеров, которые я показал, нет различий между поколениями.
Но есть и другие типы операций чтения или записи, которые могут объединяться (т. е. сворачиваться в единую транзакцию контроллера памяти) на более поздних устройствах, но не на более ранних. Одним из примеров является «широковещательное чтение»:
int i = my_int_data[0];
В приведенном выше примере все потоки читаются из одного и того же глобального местоположения. В более новых устройствах такое чтение будет «рассылаться» всем потокам в одной транзакции. В некоторых более ранних устройствах это приводило к сериализованному обслуживанию потоков. Такой пример, вероятно, не имеет последствий в записи, потому что несколько потоков, записывающих в одно место, дают неопределенное поведение. Однако «зашифрованная» запись может объединяться на более новых устройствах, но не на более старых:
my_int_data[(threadIdx.x+5)%32] = i;
Обратите внимание, что все описанные выше операции записи уникальны (внутри варпа) и принадлежат отдельной строке кэша, но они не удовлетворяют требованиям объединения на устройствах 1.0 или 1.1, но должны выполняться на более новых устройствах.
Если вы читали глобальную память описание доступа для устройств cc 1.0 и 1.1 и сравнение с более поздними устройствами, вы увидите некоторые требования для объединения на более ранних устройствах, которые были ослаблены на более поздних устройствах.
person
Robert Crovella
schedule
25.11.2013
Спасибо. Не могли бы вы подробнее объяснить, как кеш задействован в случае записи? Вы указали в объединенной транзакции чтения, что все чтения отдельных потоков происходят из одной строки кэша. Таким образом, в случае записи необъединенные записи занимают несколько строк кэша L2, верно?
- person Farzad; 25.11.2013
Да, не объединенная транзакция памяти охватывает более одной строки кэша, будь то чтение или запись. Сам кеш здесь не при чем. Кэшлайн — это фундаментальный квант обмена, осуществляемый контроллером памяти.
- person Robert Crovella; 25.11.2013
Мы провели этот эксперимент на курсе, который я проводил. Объединение оказалось более важным при записи, чем при чтении, возможно, потому, что кэши L1 и L2 хранят часть неиспользуемых данных для последующего использования.
person
Levi Barnes
schedule
25.11.2013
