В предыдущей статье мы проследили ядро сложения векторов от nvcc до варпов. В этот раз мы следим за критической инструкцией SASS (глобальной загрузкой) через аппаратное обеспечение RTX4090. Исследуемое CUDA-ядро состоит из двух строк: __global__ void vadd(const float* a, const float* b, float* c, int n) {int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < n) c[i] = a[i] + b[i];}. Скомпилированный SASS показывает: /*0080*/ IMAD.WIDE R4, R6, R7, c[0x0][0x168] ; // &b[i] /*00a0*/ LDG.E R4, [R4.64] ; // b[i]. Эти инструкции загружают элементы вектора b из глобальной памяти в регистры.
Одна LDG.E запрашивает четыре байта в каждом из 32 каналов. Для её обслуживания требуются четыре 32-байтовых сектора, одна строка кэша, одно преобразование адреса, одно пересечение кроссбара, один из тридцати шести сегментов L2, а при промахе везде — активация и четыре чтения столбцов в чипе DRAM. Наш варп живёт на одном из четырёх подразделений SM вместе с одиннадцатью другими резидентными варпами. Инструкция выдаётся в блок загрузки/сохранения (LSU), который отправляет код операции, маску активных каналов, вычисленные адреса и номер регистра в коалесцер.
Коалесцер определяет, что для 128 байтов, запрошенных варпом, нужны 4 запроса смежных секторов. Входя в кэш L1, запрос отправляется в виртуально адресуемый кэш с 4-канальной наборно-ассоциативной организацией. Поиск сравнивает теги в наборах из четырёх слотов. Поскольку это первая загрузка, происходит промах, и запрос спускается ниже. Попадание в L1 возвращается примерно за 15,4 нс — 40 тактов. Виртуальная память требует преобразования перед выходом из SM. TLB SM хранит шестнадцать недавних преобразований. Промахи стоят около 4,4 нс — одиннадцать тактов.
После преобразования один запрос на строку 128 байтов проходит через кроссбар к одному из 36 сегментов L2 по 2 МиБ, выбираемому функцией его физического адреса. Все сегменты могут обслуживать параллельно, что даёт совокупную пропускную способность в 36 раз больше, чем у одного сегмента.
Источник: Hacker News · Сводку подготовил HeadlinesBriefing