Согласованные и страйдовые чтения
Почему важно соответствие потоков и адресов.
«Согласованные и страйдовые чтения» — бесплатный урок CUDA Academy на CoddyKit. Это урок 2 из 4. Ты можешь прочитать весь урок бесплатно ниже — а потом практиковать его прямо в браузере с встроенным редактором кода и ИИ-репетитором 24/7. Это часть пути обучения CUDA Academy, и твой прогресс синхронизируется между веб-версией и приложением CoddyKit. Курс CUDA Academy содержит 4 уроков всего.
Всё зависит от сопоставления
Объединение обращений зависит от того, к какому адресу обращается каждый поток. Сопоставление потока с адресом определяет, сможет ли одна транзакция обслужить весь варп.
Объединённый шаблон
Когда соседние потоки читают соседние адреса, варп охватывает один непрерывный диапазон. Такое аккуратное расположение называется объединённым доступом.
int i = blockIdx.x * blockDim.x + threadIdx.x;
float v = data[i];Почему это эффективно
Поток 0 обращается к индексу 0, поток 1 — к индексу 1 и так далее. Все 32 адреса попадают в одну выровненную строку, поэтому GPU требуется одна транзакция.
Появляется шаг
Шаг — это фиксированный промежуток между адресами, к которым обращаются последовательные потоки. Как только этот промежуток становится больше одного элемента, объединение обращений начинает нарушаться.
float v = data[i * stride];Чтения с шагом разбредаются
При шаге 2 поток 0 читает 0, поток 1 — 2, поток 2 — 4. Теперь варп охватывает вдвое больше памяти, поэтому ему требуется больше транзакций.
Затраты растут
При удвоении шага количество затронутых строк примерно удваивается. Большие шаги могут привести к множеству транзакций для варпа, хотя используется лишь малая часть каждой строки.
Распространённая ловушка
Если назначить каждому потоку целый столбец матрицы, хранящейся по строкам, получится огромный шаг. В коде это выглядит аккуратно, но незаметно разбрасывает обращения каждого варпа.
float v = matrix[threadIdx.x * width + row];Измените сопоставление
Часто достаточно поменять местами индексы, чтобы другой индекс изменялся вместе с потоком. Пусть последовательные потоки проходят по последовательным участкам памяти — и чтения снова объединятся.
float v = matrix[row * width + threadIdx.x];Смещения тоже мешают
Даже объединённый шаблон страдает, если начинается в середине строки. Невыровненное смещение переносит варп через границу и разделяет одно чтение на два.
Думайте варпами
Чтобы оценить шаблон, не представляйте один поток. Представьте все 32 линии одновременно и спросите себя, сколько строк вместе охватывают их адреса. 🔍
Практическое правило
Пусть наиболее быстро изменяющийся индекс следует за threadIdx.x. Одна эта привычка позволяет большинству чтений из глобальной памяти автоматически оставаться объединёнными.
Быстрая проверка
Выберите шаблон доступа, который лучше всего объединяется.
Итоги
Вы увидели, что объединённые чтения удерживают соседние обращения вместе, а шаг разбрасывает их по строкам. Пусть threadIdx.x управляет наиболее быстро изменяющимся индексом. 🎉
Изучай C++ с ИИ-репетитором — бесплатно
Пиши и запускай код прямо в браузере, получай мгновенную помощь от ИИ-репетитора 24/7 и продолжи учиться на сайте или в приложении.
- Курсы
- 30
- Уроки
- 120
Часто задаваемые вопросы
Урок «Согласованные и страйдовые чтения» бесплатный?
Да — полный текст урока «Согласованные и страйдовые чтения» бесплатно доступен здесь в веб-версии. Чтобы практиковать его интерактивно (встроенный редактор кода и ИИ-репетитор 24/7) и разблокировать остальной курс CUDA Academy, подпишись на CoddyKit PRO. Курс CUDA Academy содержит 4 уроков всего.
Чему я научусь в уроке «Согласованные и страйдовые чтения»?
Почему важно соответствие потоков и адресов. Ты практикуешь CUDA Academy с помощью реального кода, который запускаешь прямо в браузере, и ИИ-репетитор 24/7 отвечает на твои вопросы во время урока.
Нужен ли мне опыт, чтобы начать CUDA Academy?
Предыдущий опыт не требуется. CUDA Academy на CoddyKit структурирован для всех уровней — от новичков до продвинутых, поэтому ты можешь начать отсюда или с самого начала и учиться в своем темпе. Это урок 2 из 4.
Сколько времени занимает урок «Согласованные и страйдовые чтения»?
Большинство уроков CoddyKit занимают около 5–10 минут. Каждый из них компактный и интерактивный, поэтому ты постоянно делаешь прогресс и продолжаешь с того же места в веб-версии и приложении.
Можно ли писать и запускать код в этом уроке CUDA Academy?
Да. Каждый урок CUDA Academy включает встроенный редактор кода, поэтому ты пишешь и запускаешь реальный код прямо в браузере и получаешь моментальную обратную связь от AI — локальная установка не требуется.
Все уроки этого курса
- Что такое транзакция памяти
- Согласованные и страйдовые чтения
- Массив структур и структура массивов
- Измерение эффективной пропускной способности