0Pricing
CUDA Academy · Урок

Варпы, lanes и маски

Единица из 32 потоков и активные маски.

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

Части этого урока еще не переведены и отображаются на английском.

Threads Run in Warps

The GPU does not schedule threads one by one. It groups them into a warp of 32 threads that march together, executing the same instruction in lockstep.

Why 32 Matters

On NVIDIA hardware a warp is always 32 threads. It is the real unit of execution, so good kernels think in groups of 32, not single threads.

Each Thread Is a Lane

Inside a warp, every thread has a position from 0 to 31 called its lane. The lane id is how warp primitives know who is talking to whom.

int lane = threadIdx.x % 32;

Finding the Lane Id

You can read the lane directly from a special register instead of computing it. The laneid register always holds 0 through 31 for the current thread.

Lockstep Has a Catch

Threads share one program counter per warp. When an if sends lanes down different paths, the warp diverges and runs each path in turn, hurting speed.

The Active Mask

Not every lane is always running. A 32-bit mask marks which lanes are active right now, with one bit per lane set to 1 when that lane participates.

Why Masks Exist

After divergence only some lanes are live. Warp primitives need to know exactly who is present, so you pass them an active mask to stay correct.

The Full Mask

When you are sure all 32 lanes are active, the mask is 0xffffffff, every bit set. This full mask is the most common value you will pass.

unsigned mask = 0xffffffff;

Building a Mask Safely

Inside a branch, do not guess the mask. Call activemask to capture exactly which lanes reached this point right now.

unsigned mask = __activemask();

Lanes Talk Without Memory

The big win is that lanes in one warp can swap data directly through registers. No shared memory and no barriers are needed for warp-local exchange.

Sync Means the Warp

The sync suffix on these intrinsics is a warp-level handshake, not a block barrier. It only coordinates the lanes named in the mask you give it.

Quick Check

Recall what a warp is and how big it is on NVIDIA GPUs.

Recap

A warp is 32 lanes running in lockstep, and a mask tracks who is active. That foundation lets lanes share data fast. Next: shuffles for reductions. ✨

Часто задаваемые вопросы

Урок «Варпы, lanes и маски» бесплатный?

Да — полный текст урока «Варпы, lanes и маски» бесплатно доступен здесь в веб-версии. Чтобы практиковать его интерактивно (встроенный редактор кода и ИИ-репетитор 24/7) и разблокировать остальной курс CUDA Academy, подпишись на CoddyKit PRO. Курс CUDA Academy содержит 4 уроков всего.

Чему я научусь в уроке «Варпы, lanes и маски»?

Единица из 32 потоков и активные маски. Ты практикуешь CUDA Academy с помощью реального кода, который запускаешь прямо в браузере, и ИИ-репетитор 24/7 отвечает на твои вопросы во время урока.

Нужен ли мне опыт, чтобы начать CUDA Academy?

Предыдущий опыт не требуется. CUDA Academy на CoddyKit структурирован для всех уровней — от новичков до продвинутых, поэтому ты можешь начать отсюда или с самого начала и учиться в своем темпе. Это урок 1 из 4.

Сколько времени занимает урок «Варпы, lanes и маски»?

Большинство уроков CoddyKit занимают около 5–10 минут. Каждый из них компактный и интерактивный, поэтому ты постоянно делаешь прогресс и продолжаешь с того же места в веб-версии и приложении.

Можно ли писать и запускать код в этом уроке CUDA Academy?

Да. Каждый урок CUDA Academy включает встроенный редактор кода, поэтому ты пишешь и запускаешь реальный код прямо в браузере и получаешь моментальную обратную связь от AI — локальная установка не требуется.

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

  1. Варпы, lanes и маски
  2. __shfl_down_sync для свёрток
  3. Функции голосования и ballot
  4. Кооперативные группы
← Назад к CUDA Academy