0Pricing
CUDA Academy · Lesson

Warps, Lanes, and Masks

The 32-thread unit and active masks.

Warps, Lanes, and Masks is a free CUDA Academy lesson on CoddyKit — lesson 1 of 4. You can read the complete lesson below for free — then practise it hands-on in the browser with a built-in code editor and a 24/7 AI tutor. It is part of the CUDA Academy learning path, one of 4 lessons in the course, and your progress syncs across the web and the CoddyKit app.

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. ✨

Frequently asked questions

Is the “Warps, Lanes, and Masks” lesson free?

Yes — the full text of “Warps, Lanes, and Masks” is free to read here on the web, and the CUDA Academy course includes 4 lessons in total. To practise it interactively (a built-in code editor and a 24/7 AI tutor) and unlock the rest of the CUDA Academy course, upgrade to CoddyKit PRO.

What will I learn in “Warps, Lanes, and Masks”?

The 32-thread unit and active masks. You practise CUDA Academy with hands-on code you run directly in the browser, and a 24/7 AI tutor answers your questions as you work through the lesson.

Do I need any experience to start CUDA Academy?

No prior experience is required. CUDA Academy on CoddyKit is structured for beginners through advanced learners; this is — lesson 1 of 4, so you can start here or from the beginning and move at your own pace.

How long does the “Warps, Lanes, and Masks” lesson take?

Most CoddyKit lessons take about 5–10 minutes. Each one is bite-sized and interactive, so you make steady progress and pick up exactly where you left off across the web and the app.

Can I write and run code in this CUDA Academy lesson?

Yes. Every CUDA Academy lesson includes a built-in code editor, so you write and run real code right in your browser and get instant AI feedback — no local setup required.

All lessons in this course

  1. Warps, Lanes, and Masks
  2. __shfl_down_sync for Reductions
  3. Ballot and Vote Functions
  4. Cooperative Groups
← Back to CUDA Academy