Предыдущий материал проследил путь ядра сложения векторов — c[i] = a[i] + b[i], один поток на float — от nvcc вниз до варпов. Тогда были разобраны детали запуска ядра, но многое осталось за кадром.
На этот раз разбирается путь, который проходит через железо критическая инструкция SASS (глобальная загрузка) — в данном случае, на карте RTX 4090. Подобный реверс-инжиниринг делается ради понимания производительности. Документации NVIDIA по большей части деталей этого пути недостаточно, поэтому значения определялись через замеры времени непосредственно на самом железе.
Исследуемое 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 из глобальной памяти в регистр, где значения складываются с элементами a для выполнения ядра. Одна инструкция LDG.E запрашивает по четыре байта в каждой из 32 линий (lanes). Для её обработки требуется четыре 32-байтовых сектора, одна строка кэша, одна трансляция адреса, пересечение crossbar, один из тридцати шести слайсов L2, а при промахе на всех уровнях — активация строки и четыре чтения столбцов в чипе DRAM. Именно этот путь инструкции через железо и обратно и разбирается ниже.
Для контекста: варп живёт на одном из четырёх суб-партиций SM, среди одиннадцати других резидентных варпов. Каждый такт планировщик суб-партиции выбирает один готовый к выполнению варп и выдаёт его следующую инструкцию сразу на все 32 линии. Наш варп выигрывает арбитраж дважды: первый раз для IMAD.WIDE, второй — спустя несколько тактов (когда адреса уже лежат в R4 и R5) для LDG.
История начинается с LDG.
От варпа к кэшу L1
Начнём с самой инструкции. LDG.E R4, [R4.64] — это глобальная загрузка 32 бит из 64-битного адреса, хранящегося в регистрах R4 и R5 (R5 появляется из-за аннотации .64: регистры имеют размер 32 бита), с сохранением результата в регистр R4. Чтобы загрузить сами данные, сначала нужно получить этот адрес из регистров.
Одна строка регистрового файла хранит R4 сразу для всех 32 линий (чтения предварительно проходят через operand collector — он нужен для инструкций, чьи источники делят один банк регистрового файла, поскольку банк обслуживает по одному чтению за такт; банков два, выбираемых по младшему биту номера регистра, так что соседняя пара всегда охватывает оба банка). Другая строка хранит R5. Варп читает обе записи, получая 256 байт в виде 32 отдельных 64-битных адресов — по одному адресу на линию.
Сколько стоит получение регистра
Чтение адреса добавляет максимум один такт. Загрузка из shared-memory, берущая адрес из регистра, занимает 24 такта от выдачи до первого использования, а та же загрузка с адресом как непосредственным значением — 23 такта (LDG не может принимать непосредственное значение).
После того как все адреса разрешены, инструкция выдаётся на блок load/store (LSU). LSU принимает инструкцию и адреса операндов, при необходимости выполняет арифметику над адресами (этот блок умеет добавлять непосредственные смещения — [R4.64] смещения не несёт — и задавать область видимости загрузки; LDG напрямую именует глобальное окно), а затем передаёт дальше опкод («загрузить эти адреса», в двоичном виде), 32-битную маску активных линий, вычисленные адреса и номер регистра для результата. Следующий получатель — коалесер.
Каждая инструкция LDG.E в каждой линии запрашивает 4 байта, но следующий пункт назначения, кэш L1, адресуется 32-байтовыми секторами. Задача коалесера — вычислить минимальное число секторов L1, которые нужно получить для обслуживания всех 4-байтовых запросов.
Коалесер определяет, что для 128 байт, запрошенных варпом, нужно отправить 4 непрерывных запроса секторов.
Вход в кэш L1
Запрос на четыре непрерывных 32-байтовых сектора отправляется в кэш L1.
Единица организации L1 ещё крупнее: 128-байтовые строки. Четыре непрерывных сектора представляют собой четыре части одной строки, поэтому в L1 формируется запрос именно на эту строку кэша.
Сначала нужно определить, есть ли эта строка уже в кэше. Кэш разделён на группы слотов, называемые сетами (в терминологии — L1 на 4090 является 4-way set-associative; кэши лежат на континууме между полностью ассоциативными, где строка может храниться где угодно, и «прямым отображением», где для каждой строки есть ровно одно место), и адрес строки определяет, к какому сету она принадлежит. Сет на этой карте содержит четыре слота, и каждый несёт тег, идентифицирующий хранящуюся строку. Поиск сравнивает все четыре тега с тегом искомой строки. Используемый адрес — виртуальный (предположительно, чтобы не платить за трансляцию при попадании в L1) адрес, применяемый в программе. Сет, в который попадает строка, вычисляется из виртуального адреса строки схемой хеширования (это сложная схема чётности, а не просто срез битов, — чтобы обращения со страйдом, кратным степени двойки, вроде столбцов матрицы или тензора, не попадали постоянно в одни и те же сеты, вызывая churn), которую можно восстановить реверс-инжинирингом.
Если один из четырёх тегов совпадает и нужные секторы находятся в этом слоте, данные считываются и загрузка завершена. Поскольку данные загружаются впервые, запрос промахивается и должен спуститься дальше по системе памяти.
Сколько стоит попадание в L1
Попадание в L1 возвращается примерно за 15,4 нс — 40 тактов. Число получено при прогоне одного потока по цепочке зависимостей через случайную перестановку строк, резидентных в L1.
В поисках L2: трансляция
Виртуальная память добавляет уровень косвенности между адресами, которыми оперирует программа, и адресами, по которым железо реально хранит данные. Программа получает собственное непрерывное адресное пространство, а железо раскладывает его по физическим страницам как ему удобно. Трансляция — это соответствие между ними.
L1, к которому только что обращались, был виртуально адресуемым, поэтому о трансляции заботиться не пришлось. Дальше приходится переходить на язык самого железа — промах L1 должен быть транслирован до того, как покинет SM.
Фактическое соответствие между физическими и виртуальными адресами устанавливается драйвером при аллокации: когда выделялся вектор b, драйвер выбрал для него физические страницы (по 2 МиБ) и записал в VRAM таблицы страниц, фиксирующие это соответствие.
Блок трансляции принимает виртуальный адрес и возвращает физический согласно этим таблицам. SM хранит шестнадцать последних трансляций в TLB, общем для всех варпов. Самая первая загрузка промахнётся в этом TLB.
Сколько стоит трансляция
Ни один из имеющихся замеров не показывает заметной стоимости попадания в TLB. Промах стоит около 4,4 нс — одиннадцать тактов. Та же стоимость перезаполнения держится в пределах 0,1 нс для всех страниц, которые может отобразить этот чип, и с любого SM — значит, следующий уровень кэша трансляции универсален и очень дёшев.
После трансляции наружу уходит один запрос на 128-байтовую строку: теперь с физическим адресом строки и маской нужных секторов. В нашем случае это единственный запрос со всеми четырьмя отмеченными секторами.
Запрос выходит из SM, пересекает crossbar и направляется в кэш L2.
Потерянные в L2
Запрос проходит через crossbar в один из 36 слайсов L2 по 2 МиБ, выбираемый достаточно сложной функцией от физического адреса. Любой SM может обратиться к любому слайсу. Все слайсы обслуживают запросы параллельно, поэтому суммарная пропускная способность в 36 раз выше, чем у одного слайса.
Внутри слайса структура такая же, как у L1. Каждый слайс содержит 1024 сета. Сет, к которому принадлежит строка, выбирается хешем физического адреса строки. Каждый сет теперь содержит 16 слотов: слайсы по отдельности являются 16-way set-associative. Размер строк — 128 байт, как и в L1.
Строки нет в L2, поскольку она ещё не запрашивалась (это, пожалуй, некоторое художественное допущение: загрузка вектора b из памяти хоста через PCIe могла бы закэшировать её в L2, но тогда нельзя было бы продолжить спуск к DRAM). Каждый слайс замыкается на один из 12 контроллеров памяти — по 3 слайса на контроллер. Задача каждого контроллера памяти — общаться с одним чипом GDDR6X DRAM. Запрос передаётся этому контроллеру.
Сколько это стоит
Попадание в L2 стоит около 127 нс — около 330 тактов. Каждый SM может передавать crossbar до двух запросов строк за такт, а 36 слайсов обслуживают запросы независимо. Счётчик выходного порта — l1tex__m_l1tex2xbar_req_cycles_active.
Найдено в DRAM
Задача контроллера памяти — загрузить данные из своего чипа DRAM объёмом 2 ГиБ. Делается это через выдачу команд по шине.
DRAM разделена на две отдельные шины, которыми контроллер управляет независимо, — каналы. На каждом канале расположено 16 банков: двумерных массивов ячеек памяти. Банк состоит из 65 536 строк. Железо может открыть одновременно только одну строку (активация, дорогая операция), а затем вернуть любые 32-байтовые столбцы из этой строки (чтение, дешёвое, пока строка открыта).
Адрес в последний раз разбирается на части, чтобы соответствовать этой структуре памяти. Из него выделяются канал, банк, строка и столбец. Четыре нужных сектора — это четыре столбца одной строки.
Так что для обслуживания загрузки контроллер памяти должен сначала отправить одну команду активации, а затем четыре команды чтения.
Как отвечает DRAM
Что делает чип DRAM в ответ на эти команды?
Каждая ячейка DRAM — это один конденсатор за одним транзистором. Транзисторы одной строки делят общую wordline, подключённую к их затворам. Каждый транзистор расположен между своим конденсатором и bitline, которая идёт вдоль столбца, обеспечивая путь от каждой ячейки (общий с ячейками других строк) к усилителям считывания. Биты хранятся в виде заряда конденсатора. Конденсаторы постоянно теряют заряд, поэтому чипу приходится время от времени приостанавливать каждый банк, чтобы подзарядить их.
Команда активации запускает декодер строки, который подаёт напряжение на wordline этой строки, открывая её транзисторы и направляя заряд с конденсаторов этой (и только этой) строки через bitline в усилители считывания, которые усиливают заряд до полноценных битов и удерживают их для контроллера.
Когда выдаётся команда чтения, её адрес столбца выбирает 256 из этих битов строки. Считывание с усилителей даёт сразу очень много битов, но их нужно сериализовать на выводы, передающие данные обратно по шине. На канал приходится 16 выводов данных. 256 бит, полученных при чтении, уходят по этим выводам в виде символов PAM4 (GDDR6X — это стандарт GDDR6 с добавленной сигнализацией PAM4): каждый символ — один из четырёх уровней напряжения, несущий два бита, так что 256 бит на 16 выводов дают по 16 бит на вывод — восемь символов. Тактовый сигнал передаётся по общему проводу, чтобы контроллер мог сэмплировать в нужные моменты.
Обратный путь
Эти пакеты PAM4 десериализуются в контроллере памяти и записываются в строку слайса L2. Результаты идут обратно через crossbar, обратно к своему SM и заполняют слот в L1. Они встречаются с записью, оставленной при отправке запроса, и их байты записываются в регистр R4 сразу по всем линиям.
Когда загрузка была выдана, был установлен барьер зависимости, который снимается этой записью в регистр. Варп снова становится готовым к выполнению, и на следующем такте планировщика выигрывает арбитраж. Выдаваемая инструкция — это сложение, ожидавшее b[i].
Полный цикл — L1, TLB, crossbar, L2, контроллер и обратно — стоит около 255 нс, порядка 660 тактов. Всё это время варп простаивал на барьере. Но остальной чип не бездействовал. Суб-партиция выдавала те же загрузки ещё для 11 варпов, остальная часть SM — ещё для 36, другие SM — ещё для 6096. В итоге получается какофония загрузок, в которой задержка любой конкретной операции теряется в общем шуме.
Приложение: пробы
Настройка
Все измерения проводились на одной RTX 4090 (sm_89) с ядром, зафиксированным на частоте 2,6 ГГц. Такты вычисляются из измеренных наносекунд на этой частоте. Использовались два основных инструмента:
Погоня по цепочке задержки. Чтобы измерить задержку (особенно когда изменение задержки само по себе говорит о чём-то важном), запускается указательный цикл через выбранный набор строк, совершающий 20 000 переходов, после чего измеряется среднее число наносекунд на переход. Если строки, на которые указывает цикл, помещаются в уровень кэша, они остаются резидентными, и среднее значение равно задержке попадания на этом уровне. Из-за крутизны иерархии любые загрузки, переполняющиеся на следующий уровень вниз, заметно проявляются в среднем значении. ld.global.ca (LDG.E…STRONG.SM) — для проб на уровне L1, ld.global.cg (LDG.E…STRONG.GPU) — минуя L1. Задержки попадания составляют 15,4 нс на L1, 127,4 нс на L2 и 255,4 нс на DRAM.
Аппаратные счётчики. Чтобы надёжно считывать счётчики ncu, их нужно брать как наклон по числу итераций, чтобы фиксированные накладные расходы сокращались. Счётчики секторов и запросов на выходном порту L1 и со стороны L2 используются, чтобы больше узнать о форме запросов, а счётчик секторов на уровне слайса помогает получить измерения для слайсов L2.
Функция сета L1
8 бит индекса L1 — это XOR фиксированного подмножества битов адреса. В виде битовой маски над адресом один из базисов для этих подмножеств выглядит так:
| бит | маска | бит | маска | |
|---|---|---|---|---|
| 0 | 0xc3901e00 | 4 | 0x47810400 | |
| 1 | 0x119a80a00 | 5 | 0x1b4e09180 | |
| 2 | 0x167041b00 | 6 | 0xb6405400 | |
| 3 | 0xdbc21d80 | 7 | 0xdc202c80 |
Сами маски не уникальны — любая обратимая комбинация этих восьми задаёт то же разбиение.
Таблицы страниц и TLB
16-элементный TLB — только первый уровень, но что происходит при промахе? Промах перезаполняется примерно за 4,4 нс, а попадание в L2 стоит 127 нс, обращение к VRAM — 255 нс, так что источником не могут быть они. Вывод — задержка берётся из некоего более крупного кэша трансляции на кристалле.
Стоимость постоянна в пределах 0,1 нс для всех страниц, которые может отобразить чип, и с любого SM. Дополнительное подтверждение: обход таблиц страниц через nvdebug показывает установленный бит volatile у каждой записи каталога, значит, они не кэшируются в обычной иерархии.
Функция слайса L2
Определить, какому слайсу принадлежит строка, довольно сложно. L2 индексируется физически, поэтому проба должна работать с физическими адресами устройства, полученными при обходе таблиц страниц. В Nsight Compute есть счётчик секторов на слайс, но он выдаёт только минимум, максимум, среднее и сумму по всем 36 инстансам, никогда не сообщая конкретный индекс слайса.
Тем не менее суммарного значения достаточно, чтобы понять, разделяют ли два адреса один слайс. Если оба адреса лежат на одном слайсе, после загрузки обоих счётчик максимума покажет 2, если на разных — 1. С помощью этой пробы можно получить репрезентативный адрес, попадающий на каждый из 36 слайсов.
Имея 36 представителей на руках, можно определить слайс любого нового кандидата. Если читать кандидата много раз вместе со всеми 36 представителями, каждый раз с разным числом обращений к каждому из адресов (скажем, 20001, 20002, … раз), сумма счётчика чтений кандидата и только одного из представителей совпадёт со значением счётчика максимума, и слайс можно определить по этому совпадению.
Из этого можно построить таблицу пар физический адрес — слайс. Сложность в том, чтобы перейти от такой таблицы к физически правдоподобной функции. Один из полезных приёмов — проведение тех же экспериментов на двух разных чипах, построенных на одном кристалле: 4090 и L40S, у которого на один слайс на контроллер памяти больше.
Вот вариант, полученный ранее с помощью Claude (сложно быть уверенным, что именно так устроено железо, но это правдоподобно исходя из имеющихся знаний; логика в том, что между L40S и 4090 должна быть какая-то общая логика, если предположить, что NVIDIA не делает два полностью разных функциональных пути для чипов с одинаковой топологией, но разным количеством отключённого L2, а сама функция должна быть достаточно простой на аппаратном уровне — XOR, арифметика и подобное допустимы, но если модель предлагает таблицу поиска на 4096 записей, это явно перебор):
SHIFT, OFFSET = (5, 0, 1), (1, 0, 0)
def parity(x):
return bin(x).count("1") & 1
def _state(a, N):
wide = (N == 48) # L40S: 4 слайса/контроллер, доходит до бита 35
b35 = (1 << 35) if wide else 0
# этап 1 — какой из 12 контроллеров: две чётности и mod-3 разряд
P1c = parity(a & 0x76A990400) # чётность контроллера 1 (узкая; используется на обоих чипах)
P1 = parity(a & (0x76A990400 ^ b35)) # широкая форма, нужна только для считывания на L40S
P2 = parity(a & 0x2CCF7B000) # чётность контроллера 2
A = ((a >> 15) + 2*parity(a & 0x3C9041000) + parity(a & (0x2882B0800 ^ b35)) + 2) % 3 # mod-3 разряд: (a>>15) + 2 поправки
# этап 2 — какой слайс внутри контроллера: циклический счётчик по модулю 9
g = ((a + (1 << 16)) >> 17) % 9 # значение счётчика, round(a / 2^17) mod 9
q0 = parity(a & 0x8000) # четыре поправочные чётности
q1 = parity(a & 0x5985E0500)
q2 = parity(a & (0x2354E4400 ^ b35))
q3 = parity(a & 0x3C9041000)
carry = 1 if q0 + q1 + q2 >= 2 else 0 # q0,q1,q2 как полный сумматор: перенос (мажоритарная функция)...
start = (5 + 7*q0 + 5*q1 + 2*q2 + q3 - carry) % 9 # ...задаёт, откуда стартует счётчик
o = (g - SHIFT[A] - start) % 9 # позиция в цикле длиной 9
Lf = 2 if (q0 ^ q1 ^ q2) == 0 else 1 # ...а их XOR задаёт точку разбиения
return P1c, P1, P2, A, q2, o // 3, (1 if (o % 3) >= Lf else 0) # d = o // 3, u = бит разбиения
def slice_of(a, N=36):
P1c, P1, P2, A, q2, d, u = _state(a, N)
controller = (2*P1c + P2) * 3 + A # 0..11
if N == 36: # 4090: живы 3 слайса, читаем (d, u) как три дуги Z/9
base = 2 if d == 0 else (1 if (d == 1 and u == 0) else 0)
B = ((1 - base) % 3 if q2 else base) % 3 # q2 меняет порядок дуг
B = (B + OFFSET[A]) % 3 # смещение для конкретного контроллера
return controller * 3 + B
if N == 48: # L40S: живы 4 слайса, читаем u как два индексных бита
i0, i1 = P1 ^ q2 ^ u, P1 ^ P2 ^ u
return controller * 4 + 2*i0 + i1
raise ValueError("N must be 36 or 48")Найти такую функцию очень сложно, но проверить, что найденная функция работает, — легко. При выборе 8192 резидентных в L2 строк ровно из k предсказанных слайсов:
| строки взяты из | Mзагрузок/с | относительно k=1 |
|---|---|---|
| 1 предсказанного слайса | 1 957 | 1.00× |
| 2 | 3 917 | 2.00× |
| 4 | 7 826 | 4.00× |
| 9 | 17 582 | 8.98× |
| 18 | 34 446 | 17.60× |
| всех 36 | 68 085 | 34.78× |
Индекс сета и геометрия L2
Как только функция слайса привязывает адреса к одному слайсу, можно провести такую же археологию eviction-множеств уже на этом слайсе, чтобы выяснить структуру, которая показывает, что она 16-way set-associative (цепочка из 17 элементов трэшится, а из 16 — нет).
Индекс сета внутри слайса — та же самая функция чётности, что и у L1: десять бит с той же нелинейностью (a >> 15) mod 9 в старшем бите. К сожалению, задействованные маски различаются от слайса к слайсу. Для одного слайса:
def parity(x):
return bin(x).count("1") & 1
def set_index(a): # внутри одного слайса
q = a // 1152
b0 = parity(a & 0x0bd654c80) ^ parity(q & 0x00e500)
b1 = parity(a & 0x0bd654c80) ^ parity(q & 0x010000)
b2 = parity(a & 0x07aed8b80) ^ parity(q & 0x027c00)
b3 = parity(a & 0x03e313180) ^ parity(q & 0x045500)
b4 = parity(a & 0x03e313300) ^ parity(q & 0x080300)
b5 = parity(a & 0x0bd654e80) ^ parity(q & 0x104200)
b6 = parity(a & 0x0bd654c00) ^ parity(q & 0x200b00)
b7 = parity(a & 0x044dcb880) ^ parity(q & 0x401600)
b8 = parity(a & 0x000000200) ^ parity(q & 0x804600)
b9 = parity(a & 0x13bc21180) ^ parity(q & 0x006400) ^ int((a >> 15) % 9 in (2, 6))
return sum(b << i for i, b in enumerate([b0, b1, b2, b3, b4, b5, b6, b7, b8, b9]))У этой функции есть свойства, позволяющие её проверить на здравый смысл. Например: непрерывный блок в 72 МиБ заполняет каждый слот в каждом слайсе, не вызывая трэшинга — как и следовало ожидать.
Обновление (refresh) DRAM
Ячейки DRAM теряют заряд, поэтому их нужно периодически обновлять, что усложняет некоторые виды измерений задержек. Это можно увидеть, запустив цепочку зависимых переходов, записывающую время каждого перехода в shared-memory. Большинство обращений к DRAM возвращаются с обычной задержкой, но небольшая доля занимает больше времени, равномерно распределяясь вплоть до жёсткого потолка примерно на 210 нс выше обычного. Такое равномерное распределение — признак стопора фиксированной длины. Стопор составляет около 210 нс. Ему подвержено примерно 2% обращений. Он не затрагивает весь чип сразу — эффект более локальный, — но определить точную единицу локальности не удалось.