Несколько лет назад Марк Харрис из NVIDIA опубликовал следующее:
Эффективная транспонирование матриц в CUDA C / C ++
, в котором он описал, как выполнять транспонирование матриц быстрее использовать общую память по сравнению с наивным подходом. В методологических целях он также реализовал версию простой матричной копии на основе тайлов совместно используемой памяти.
Несколько удивительно, что копирование через тайлы совместно используемой памяти выполнялось быстрее, чем «наивное» копирование (с двумерной сеткой): 136 ГБ / с c для простой копии, 152,3 ГБ / с c для копии на основе общих записей памяти. Это было на микроархитектуре Kepler, Tesla K20 c.
Мой вопрос: почему это имеет смысл? То есть, почему эффективная пропускная способность не была ниже , когда все, что сделано, - это объединенное чтение и запись? В частности, имело ли это отношение к тому факту, что __restrict не использовался (и, следовательно, __ldg(), вероятно, не использовался)?
Примечание. Этот вопрос не касается транспонирования . Пост был о транспозиции, и его уроки хорошо взяты В нем не обсуждалось странное явление, связанное с простым нетранспонированным копированием.