Урок 15: Диагностика и исправление шагового доступа
В предыдущем уроке мы увидели, что такое объединение доступов: когда последовательные потоки в warp обращаются к последовательным адресам, одна транзакция обслуживает их все, а когда они прыгают с большим шагом (stride), warp расходует много потраченных впустую транзакций. Теперь мы освоим практичес
Представьте 4 работников, собирающих яблоки в рядах деревьев. Сначала каждый работник бежит в отдельный столбец в дальний конец сада — 4 долгих забега. Если же вместо этого каждый работник берёт следующее яблоко в одном плотном ряду, они все работают бок о бок за один проход. Мы не изменили сад — только кто-что-собирает.
- шаговый доступ (strided access)
- Шаблон, где последовательные потоки прыгают с большим шагом (например, data[threadIdx.x * N]), так что их адреса далеки друг от друга и каждый попадает в отдельную транзакцию памяти.
- объединённый против необъединённого (coalesced vs uncoalesced)
- Объединённый: адреса соседних потоков непрерывны, одна транзакция на весь warp. Необъединённый: адреса разбросаны, warp расходует много транзакций, большинство из которых потрачено впустую.
- исправление переиндексацией (index-remap fix)
- Исправление разбросанного доступа: пусть соседние потоки обращаются к соседним адресам — дайте потоку t элемент base + t вместо base + t*N, всё в глобальной памяти.
- построчный порядок (row-major order)
- Раскладка, где элементы одной строки — соседи в памяти. Обход по строкам заставляет соседние потоки обращаться к соседним адресам (объединённо); обход по столбцам прыгает на ширину (разбросанно).
