__shfl_down_sync для свёрток
Выполняйте свёртки варпа без барьеров.
«__shfl_down_sync для свёрток» — бесплатный урок CUDA Academy на CoddyKit. Это урок 2 из 4. Ты можешь прочитать весь урок бесплатно ниже — а потом практиковать его прямо в браузере с встроенным редактором кода и ИИ-репетитором 24/7. Это часть пути обучения CUDA Academy, и твой прогресс синхронизируется между веб-версией и приложением CoddyKit. Курс CUDA Academy содержит 4 уроков всего.
Части этого урока еще не переведены и отображаются на английском.
Shuffle Moves Registers
A shuffle lets one lane read another lane's register value directly. Data hops between threads with no shared memory and no barrier in between.
Meet shfl_down_sync
The workhorse for reductions is shfl_down_sync. Each lane grabs a value from a lane a fixed number of positions higher in the warp.
float v = __shfl_down_sync(mask, val, offset);Reading the Arguments
The call takes an active mask, the value to share, and an offset. Lane i receives the value held by lane i plus offset.
Add What You Receive
A warp sum just adds the shuffled value back in. Lane i takes its neighbor's number and folds it into its own running total.
val += __shfl_down_sync(mask, val, offset);Halve the Offset
Like a tree reduction, the offset starts at 16 and halves each step: 16, 8, 4, 2, 1. After five steps the whole warp is summed.
for (int o = 16; o > 0; o >>= 1)
val += __shfl_down_sync(mask, val, o);Why Five Steps
A warp has 32 lanes, and 2 to the fifth power is 32. So five halving steps are exactly enough to combine all 32 values into one.
Answer Lands in Lane 0
After the loop, the complete warp total sits in lane 0. The other lanes hold partial junk, so only lane 0 should write the result out.
No Barriers Needed
Because exchange happens through registers in lockstep, you skip syncthreads entirely. The sync suffix already coordinates the masked lanes.
Faster Than Shared
A warp shuffle reduction beats the shared-memory version: fewer instructions, no bank conflicts, and no barrier stalls. It is the fast path for the last 32.
Pass the Right Mask
If some lanes already exited, a full mask is wrong. Capture the live set with activemask so each shuffle only touches present lanes.
Two-Level Reductions
Big reductions often reduce each warp with shuffles, then combine warp results in shared memory. Shuffles handle the cheap inner level beautifully.
Quick Check
Think about which lane holds the answer when the loop finishes.
Recap
You summed a warp with shfl_down_sync: halve the offset five times, and lane 0 holds the total, no barriers needed. Next: ballot and vote. ✨
Часто задаваемые вопросы
Урок «__shfl_down_sync для свёрток» бесплатный?
Да — полный текст урока «__shfl_down_sync для свёрток» бесплатно доступен здесь в веб-версии. Чтобы практиковать его интерактивно (встроенный редактор кода и ИИ-репетитор 24/7) и разблокировать остальной курс CUDA Academy, подпишись на CoddyKit PRO. Курс CUDA Academy содержит 4 уроков всего.
Чему я научусь в уроке «__shfl_down_sync для свёрток»?
Выполняйте свёртки варпа без барьеров. Ты практикуешь CUDA Academy с помощью реального кода, который запускаешь прямо в браузере, и ИИ-репетитор 24/7 отвечает на твои вопросы во время урока.
Нужен ли мне опыт, чтобы начать CUDA Academy?
Предыдущий опыт не требуется. CUDA Academy на CoddyKit структурирован для всех уровней — от новичков до продвинутых, поэтому ты можешь начать отсюда или с самого начала и учиться в своем темпе. Это урок 2 из 4.
Сколько времени занимает урок «__shfl_down_sync для свёрток»?
Большинство уроков CoddyKit занимают около 5–10 минут. Каждый из них компактный и интерактивный, поэтому ты постоянно делаешь прогресс и продолжаешь с того же места в веб-версии и приложении.
Можно ли писать и запускать код в этом уроке CUDA Academy?
Да. Каждый урок CUDA Academy включает встроенный редактор кода, поэтому ты пишешь и запускаешь реальный код прямо в браузере и получаешь моментальную обратную связь от AI — локальная установка не требуется.
Все уроки этого курса
- Варпы, lanes и маски
- __shfl_down_sync для свёрток
- Функции голосования и ballot
- Кооперативные группы