What is a lane in a CUDA warp?

A thread's position inside its warp, 0 to 31, which is what shuffle and vote instructions address.

Lane 0 of warp 1 is thread 32. Every thread carries two numbers, the one you gave it and the one the hardware gave it, and the second is the one the collective instructions speak. A shuffle names a source lane. A ballot returns one bit per lane. __activemask() hands back a word where "The Nth bit is set if the Nth lane in the warp is active", so a mask is a set of lanes and never a set of thread IDs. NVIDIA's own guidance on the warp primitives puts the mask first: "The set of threads that participates in invoking each primitive is specified using a 32-bit mask, which is the first argument of these primitives."

Getting the number is cheap and the usual two spellings are the same thing. threadIdx.x & 31 and threadIdx.x % 32 both extract the low five bits, because the operand is unsigned and 32 is a power of two, so the compiler has no division to emit either way. PTX also names it outright as the special register %laneid (section 10.3 of the ISA), which is what you reach for when reading SASS on day 46.

The trap is that neither spelling is the lane id in general. Warps are cut from the linear thread ID, x + y * blockDim.x + z * blockDim.x * blockDim.y, so threadIdx.x % 32 is a lane id only while the block is one-dimensional or its x extent is a multiple of 32. A 32 by 8 block puts one row in each warp and the shortcut survives. A 16 by 16 block does not: warp 0 holds two whole rows, threads (0,0) and (0,1) are lanes 0 and 16, and the shortcut calls them both lane 0. Compute the linear ID first, then mask it.

The reason to care is that the lane number decides what an instruction costs. Memory coalescing is a statement about lanes: consecutive lanes on consecutive addresses is the cheap case. A shared-memory bank is (byteAddress / 4) % 32, one bank per lane when the mapping works out and a bank conflict when it does not. Nothing on either list mentions block IDs, and neither does the hardware.

Measured

Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), built with nvcc -O3 -arch=sm_75. Day 21 has every thread record __activemask() into its own slot and prints one launch, the 48-thread block, thread by thread. Captured 2026-08-30 on the project's verification node.

Threads	Warp	Lanes present	Active mask each of them recorded
0 to 31	0	0 to 31	`0xffffffff`
32 to 47	1	0 to 15	`0x0000ffff`

Thread 32 is lane 0, thread 47 is lane 15, and the sixteen of them agree on a mask with sixteen bits set. Lanes 16 to 31 of that warp would be threads 48 to 63, which the launch never created, so the mask names exactly the threads that exist. Push the block down to 33 threads and the second warp is one lane wide: mask 0x00000001.

The program checks the mapping rather than printing it. For every thread of every launch in the sweep it recomputes the lane on the host and asserts that bit is set in the mask the device wrote, on the guide's rule that a calling thread is active. Every gate passed, so the lane arithmetic on this page came back off the hardware and is not a diagram of it.

Code

From code/day21-warps/warps.cu, the host-side gate. t is a thread index inside one block, so the mask and the shift are computed by two different machines and compared.

const unsigned int lane = static_cast<unsigned int>(t % kWarpSize);
if (((h_masks[t] >> lane) & 1u) == 0u) {
    std::fprintf(stderr,
                 "thread %d of the %d-thread block recorded mask "
                 "0x%08x, which does not include its own lane %u\n",
                 t, blockSize, h_masks[t], lane);
    status = EXIT_FAILURE;
    break;
}

Diagram

occupancy-stepper, preset partial-warp, lane labels on: two warp slots holding a 48-thread block, the lanes of each slot numbered 0 to 31 and the thread ID printed under each live lane, so thread 32 sits over lane 0 of the second slot.

Alt text: "Two warp slots holding 48 threads. Lane numbers restart at zero in the second warp, so thread 32 is lane 0, and the last sixteen lane positions carry no thread at all."

Related terms

Where you meet this

Day 11, memory coalescing, where what consecutive lanes read decides the bandwidth.
Day 15, shared memory bank conflicts, where lanes and banks line up one for one.
Day 21, what is a warp, the lesson that owns this term.
Day 23, warp shuffles, where you pass a lane number to an instruction for the first time.
Colab setup, which is enough card to run day 21 and read your own lane table back.

Sources

CUDA Programming Guide, C++ language extensions, for __activemask() and the Nth-bit-Nth-lane rule, and for the warp shuffle functions that take a source lane: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cpp-language-extensions.html (checked 2026-08-30)
CUDA Programming Guide, "SIMT Architecture", for warps being cut from consecutive increasing thread IDs: https://docs.nvidia.com/cuda/cuda-programming-guide/03-advanced/advanced-kernel-programming.html (checked 2026-08-30)
PTX ISA, special registers, section 10.3, which names %laneid: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html (checked 2026-08-30)
"Using CUDA Warp-Level Primitives", NVIDIA developer blog, for the mask as the first argument of every primitive: https://developer.nvidia.com/blog/using-cuda-warp-level-primitives/ (checked 2026-08-30)

Byline

Written by: pending. Reviewed by: pending. Written on: pending. Last checked on: pending. The numbers came off the verification node on 2026-08-30, and this entry stays a draft until a named author and a different named reviewer sign it.