CUDA Academy · Урок

Устранение расхождения варпов

Переназначайте индексы, чтобы варпы оставались занятыми.

Урок 2 из 413 шагов

«Устранение расхождения варпов» — бесплатный урок CUDA Academy на CoddyKit. Это урок 2 из 4. Ты можешь прочитать весь урок бесплатно ниже — а потом практиковать его прямо в браузере с встроенным редактором кода и ИИ-репетитором 24/7. Это часть пути обучения CUDA Academy, и твой прогресс синхронизируется между веб-версией и приложением CoddyKit. Курс CUDA Academy содержит 4 уроков всего.

Варпы работают синхронно

Варп — это 32 потока, которые одновременно выполняют одну и ту же инструкцию. Когда их пути совпадают, оборудование работает с полной скоростью.

Цена расхождения

Если потоки одного варпа выбирают разные ветви, это называется расхождением. Оборудование выполняет каждую ветвь последовательно, оставляя некоторые каналы без работы и расходуя такты.

Наивная свёртка приводит к расхождению

Простая версия использует tid % (2*s) для выбора активных потоков. Активные и бездействующие потоки чередуются внутри каждого варпа, поэтому в каждом варпе возникает сильное расхождение.

if (tid % (2 * s) == 0)
  data[tid] += data[tid + s];

Бездействующие каналы всё равно требуют ресурсов

Даже если половина потоков ничего не делает, они всё равно занимают варп. Варп не может завершить работу, пока не будут обработаны и активная, и бездействующая ветви.

Переиндексация по идентификатору потока

Исправление состоит в том, чтобы назначить активную работу потокам с наименьшими идентификаторами, а не разбросанным потокам. Вычислите индекс из идентификатора потока и шага.

int index = 2 * s * tid;
if (index < blockDim.x)
  data[index] += data[index + s];

Почему это помогает

Теперь активные потоки идут подряд: потоки 0, 1, 2 и так далее работают, а остальные бездействуют. Целые варпы становятся либо полностью активными, либо полностью бездействующими.

Полностью бездействующие варпы ничего не стоят

Варп, в котором все каналы бездействуют, просто завершается без работы. Последовательного выполнения по каналам нет, поэтому цена расхождения практически исчезает.

Ловушка операции по модулю

Скрытой причиной было условие с операцией по модулю. Оно разбросало активные потоки по каждому варпу, что и создавало расхождение.

Та же работа, более удачное распределение

Вы не изменили математику или число сложений. Вы лишь переназначили, какой поток выполняет каждое сложение, и варпы работают эффективнее.

Эффект усиливается с масштабом

При работе тысяч блоков и выполнении множества шагов устранение расхождения даёт заметное ускорение — часто в несколько раз по сравнению с наивным ядром.

Осталась ещё одна проблема

Эта версия читает соседние значения, расположенные вперемежку в разделяемой памяти, что может привести к конфликтам банков. Следующий урок устранит и эту проблему.

Быстрая проверка

Подумайте, что вызывает расхождение варпа в наивной свёртке.

Ответь в интерактивном уроке

Итоги

Вы устранили расхождение, назначив работу потокам с наименьшими идентификаторами, поэтому варпы стали полностью активными или полностью бездействующими. Та же математика, но более быстрая свёртка. Далее — конфликты банков. 🚀

Можно начать бесплатно

Изучай 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 — локальная установка не требуется.

Все уроки этого курса

← Назад к CUDA Academy