What is a warp in CUDA (and why warp size is 32)
Launch a block of 48 threads and you have asked the SM for two warps, not one and a half. The second one holds its slot for the life of the block with 16 of its 32 lanes switched off, and those 16 never run an instruction of your kernel.
That is not rounding in the driver. It explains much of the confusion around the most-viewed concept question in Stack Overflow's cuda tag, "How do CUDA blocks/warps/threads map onto CUDA cores?", asked in May 2012 and read 86,521 times since (https://stackoverflow.com/questions/10460742/how-do-cuda-blocks-warps-threads-map-onto-cuda-cores , checked 2026-08-30). Fourteen years of answers describe the behavior without measuring it.
This page measures it instead. You will record clock64() and __activemask() per thread into a device array, read the schedule back on the host, and look at the second warp of a 48-thread block with sixteen of its bits missing.
The unit is 32 threads, and it is not the unit you asked for
You name a block in the launch configuration. The hardware does not schedule blocks. "Each SM creates, manages, schedules, and executes threads in groups of 32 parallel threads called warps" (https://docs.nvidia.com/cuda/cuda-programming-guide/03-advanced/advanced-kernel-programming.html section 3.2.2.1, checked 2026-08-30).
A warp is what gets picked: "At each instruction issue cycle, a warp scheduler selects a warp with threads ready to execute its next instruction (the active threads of the warp) and issues the instruction to those threads" (same page, section 3.2.2.2). One instruction goes to 32 lanes. That is SIMT, and it is where the number 32 enters every other lesson in this course.
The split from block to warp is fixed, not a scheduling decision: "The way a block is partitioned into warps is always the same; each warp contains threads of consecutive, increasing thread IDs with the first warp containing thread 0", and the count is ceil(T / 32) (same section). Threads 0 to 31 are warp 0 on any GPU, on any run.
So 48 threads is one and a half warps, and there is no half. The block gets two warps, and lanes 16 to 31 of the second one hold thread IDs 48 to 63, which were never launched. The guide names this case exactly when it lists the reasons a thread can be inactive: "having exited earlier than other threads of their warp, having taken a different branch path than the branch path currently executed by the warp, or being the last threads of a block whose number of threads is not a multiple of the warp size" (same page, section 3.2.2.1.1).
The first two describe a thread that started and then left or branched away, which is warp divergence and day 22's subject. You create the third case when you choose a block size.
The cost is a hardware resource. A Tesla T4 holds 32 resident warps per SM (https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/compute-capabilities.html , Table 30, checked 2026-08-30). A 48-thread block spends two of those 32 slots to run 48 threads.
A 64-thread block spends the same two slots and runs 64. Occupancy cannot tell them apart, because it counts warps: day 2 asked the driver and measured 16 blocks, 32 warps and 100 percent for both block sizes on this T4. The metric does not show the sixteen idle lanes, so this page measures them.
Lockstep is what it does, not what it promises
The picture most people carry is that a warp's 32 threads move in lockstep, one instruction at a time, sharing a program counter. It is a good picture, it is what the hardware usually does, and on any card below compute capability 7.0 it was also a guarantee you could write code against.
The guide is precise about what changed. Below compute capability 7.0, "warps used a single program counter shared amongst all 32 threads in the warp together with an active mask specifying the active threads of the warp". From 7.0, independent thread scheduling means "the GPU maintains execution state per thread, including a program counter and call stack, and can yield execution at a per-thread granularity" (same page, section 3.2.2.1.1, checked 2026-08-30).
A schedule optimizer regroups the active threads into SIMT units, so throughput remains high but the old guarantee is gone: "Warp-synchronous code assumes that threads in the same warp execute in lockstep at every instruction, but the ability for threads to diverge and reconverge at sub-warp granularity makes such assumptions invalid."
A T4 is compute capability 7.5, so everything this page measures happens on a card with per-thread program counters. When the trace below shows 32 lanes reporting the same cycle, that is an observation about one card at one instruction, not a promise. The fix when you need the promise is __syncwarp(), the warp-scoped sibling of __syncthreads(), and the same change is why every warp primitive grew a _sync suffix and a mask argument.
__activemask() is the intrinsic that reads the mask back. "The function returns a 32-bit integer mask representing all currently active threads in the calling warp. The Nth bit is set if the Nth lane in the warp is active when __activemask() is called" (https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cpp-language-extensions.html section 5.4.6.1, checked 2026-08-30).
Read the warning on the same page before you use it: it "cannot be used to determine which warp lanes execute a given branch" and "only provides an instantaneous snapshot of the active threads within a warp". The guide gives this example of the wrong use:
if (pred) {
// Invalid: the value of 'at_least_one' is non-deterministic
// and could vary between executions.
at_least_one = __activemask() > 0;
}
The mask tells you who is here. It does not tell you who agreed with your predicate, and it does not make anyone arrive.
The documentation gives two answers about warpSize. The built-in variable is "int warpSize : A run-time value defined as the number of threads in a warp, commonly 32" (same page, section 5.4.2.2). The hardware multithreading section quoted above is firmer while defining ceil(T / 32): "Wsize is the warp size, which is equal to 32".
One document says commonly 32, while the other says equal to 32. Read warpSize at run time, compile against 32, and check that they agree. That is the first of this program's three gates.
Recording the schedule instead of printing it
Full program in code/day21-warps/warps.cu. It has one kernel and runs it five times, as a single block of 32, 33, 48, 64 and 256 threads.
All three of the rules below are about not disturbing the thing being sampled.
Nothing is printed from inside the kernel. The obvious way to see who ran when is a printf in the kernel, and it does not work: print order is not execution order, which is day 64's subject. So each thread writes its sample into its own slot of a device array and the host reads the array back.
Nothing is timed. clock64() returns "the value of a per-multiprocessor counter that is incremented every clock cycle" (https://docs.nvidia.com/cuda/archive/12.6.0/cuda-c-programming-guide/index.html section 7.13, checked 2026-08-30), which makes it an ordering probe and not a stopwatch. Two limits follow.
You can compare a reading only with another reading from the same SM, so every launch here is one block. The absolute value has no meaning, so the output shows every cycle count relative to the smallest sample in its launch.
No branch sits above the sample. A bounds guard between kernel entry and __activemask() would be part of the measurement, since the mask is a snapshot of that instant. The buffers are sized from a constexpr, a static_assert proves every block size in the sweep fits, and the kernel therefore needs no guard at all:
__global__ void recordWarpSchedule(long long* clocks, unsigned int* masks,
int* warpSizes) {
const unsigned int tid = threadIdx.x;
const long long cycle = clock64();
const unsigned int active = __activemask();
clocks[tid] = cycle;
masks[tid] = active;
warpSizes[tid] = warpSize;
}
That last line is the only place in the file where the run-time value is read, and it exists so the host can check it against the compile-time one:
// 32 on every GPU this course targets, and the number the whole page is built
// on. The built-in `warpSize` is documented as "A run-time value defined as
// the number of threads in a warp, commonly 32", so it cannot size an array or
// appear in a static_assert. This constant can do both, and the program checks
// the two against each other at run time rather than assuming they agree.
constexpr int kWarpSize = 32;
One honest caveat. Launching a single block tells you nothing about how many blocks or warps fit on an SM at once. Day 2's program asks cudaOccupancyMaxActiveBlocksPerMultiprocessor for that instead of computing it, and day 45 shows why the answer is not the goal.
Results
Measured. Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), built with nvcc -O3 -arch=sm_75. Captured 2026-08-30 on the project's verification node; full transcript in the page's evidence file.
GPU: Tesla T4 (compute capability 7.5)
warpSize: 32 from the driver, 32 compiled into this file
maxThreadsPerBlock: 1024
one block per launch, so every clock64() reading below comes from one SM
1024 slots per array, 16384 bytes of device memory in total
32 threads, 1 block: 1 warp(s), 1 distinct clock64 value(s)
warp lanes active mask cycles one mask one clock
0 32 0xffffffff 0 yes yes
33 threads, 1 block: 2 warp(s), 2 distinct clock64 value(s)
warp lanes active mask cycles one mask one clock
0 32 0xffffffff 0 yes yes
1 1 0x00000001 2 yes yes
48 threads, 1 block: 2 warp(s), 2 distinct clock64 value(s)
warp lanes active mask cycles one mask one clock
0 32 0xffffffff 0 yes yes
1 16 0x0000ffff 2 yes yes
every thread of the 48-thread block, ungrouped
thread warp lane active mask cycles
0 0 0 0xffffffff 0
1 0 1 0xffffffff 0
2 0 2 0xffffffff 0
3 0 3 0xffffffff 0
4 0 4 0xffffffff 0
5 0 5 0xffffffff 0
6 0 6 0xffffffff 0
7 0 7 0xffffffff 0
8 0 8 0xffffffff 0
9 0 9 0xffffffff 0
10 0 10 0xffffffff 0
11 0 11 0xffffffff 0
12 0 12 0xffffffff 0
13 0 13 0xffffffff 0
14 0 14 0xffffffff 0
15 0 15 0xffffffff 0
16 0 16 0xffffffff 0
17 0 17 0xffffffff 0
18 0 18 0xffffffff 0
19 0 19 0xffffffff 0
20 0 20 0xffffffff 0
21 0 21 0xffffffff 0
22 0 22 0xffffffff 0
23 0 23 0xffffffff 0
24 0 24 0xffffffff 0
25 0 25 0xffffffff 0
26 0 26 0xffffffff 0
27 0 27 0xffffffff 0
28 0 28 0xffffffff 0
29 0 29 0xffffffff 0
30 0 30 0xffffffff 0
31 0 31 0xffffffff 0
32 1 0 0x0000ffff 2
33 1 1 0x0000ffff 2
34 1 2 0x0000ffff 2
35 1 3 0x0000ffff 2
36 1 4 0x0000ffff 2
37 1 5 0x0000ffff 2
38 1 6 0x0000ffff 2
39 1 7 0x0000ffff 2
40 1 8 0x0000ffff 2
41 1 9 0x0000ffff 2
42 1 10 0x0000ffff 2
43 1 11 0x0000ffff 2
44 1 12 0x0000ffff 2
45 1 13 0x0000ffff 2
46 1 14 0x0000ffff 2
47 1 15 0x0000ffff 2
64 threads, 1 block: 2 warp(s), 2 distinct clock64 value(s)
warp lanes active mask cycles one mask one clock
0 32 0xffffffff 0 yes yes
1 32 0xffffffff 2 yes yes
256 threads, 1 block: 8 warp(s), 8 distinct clock64 value(s)
warp lanes active mask cycles one mask one clock
0 32 0xffffffff 8 yes yes
1 32 0xffffffff 10 yes yes
2 32 0xffffffff 12 yes yes
3 32 0xffffffff 14 yes yes
4 32 0xffffffff 0 yes yes
5 32 0xffffffff 2 yes yes
6 32 0xffffffff 4 yes yes
7 32 0xffffffff 6 yes yes
every gate passed: the warp size agrees three ways, every thread is in its own mask, and no launch wrote outside its block
Threads are allocated to warps in whole units of 32, and the leftovers still
cost a warp. A 33-thread block gets two warps, and the second has one active
lane with a mask of 0x00000001. Thirty-one lanes are allocated, scheduled and
idle for the block's entire life.
Every lane in a warp reports the same clock64() value, which is the
measurement behind the word "together": they are not merely started at the same
time, they execute the same instruction at the same cycle. Different warps
report different values, because they are scheduled independently.
The exercise deliberately does not use printf to establish this.
Day 64 shows that printf ordering is buffer
ordering, so it cannot tell you anything about execution order. Reading
clock64() and __activemask() into a device array and copying them back can.
Run it yourself
A free Colab T4, or any card you own. The program needs no profiler, root access, or sanitizer. It reads two intrinsics and copies an array back, so any system that can run the binary is enough.
nvcc -std=c++17 -O3 -arch=sm_75 -o warps warps.cu
./warps
There is no Compiler Explorer embed on this page: the program is about 350 lines, well past this site's 60-line embed budget. Paste it into https://godbolt.org/ (checked 2026-08-30) if that is your GPU, targeting sm_75 or lower, because the runner is a Tesla T4. With no GPU at all, the widget above answers the block-size half of this page and /setup/learn-cuda-without-a-gpu covers the rest.
Exercise
Add a 1023-thread block to the sweep, then move the __activemask() call below an if (threadIdx.x < 16) and run again. Explain what happened to the mask, and why the source does not tell you which answer you were going to get.
Time: 25 to 40 minutes. Submit: the last warp's mask at 1023 threads, the mask column before and after the branch moved, and one sentence on why both answers are correct behaviour.
Check: three gates, all real branches setting an exit status rather than asserts, because CI builds Release and NDEBUG deletes an assert out of the build that matters. The warp size the kernel read has to equal the driver's answer and the 32 compiled into the file. Every thread's own lane bit has to be set in the mask it recorded, and no launch may write a slot outside its own block.
A failure names the thread, the block size and the value. The third gate is the one your edit can trip. Harness contract at /reference/harness.
Hint 1
The mask is read at a moment, not over a region. Point at the exact instruction in your kernel where that moment happens, then list everything the hardware had to decide before reaching it.
Hint 2
A small if has two implementations available to the compiler: take a branch, or run both sides with the store predicated off for the lanes that fail. One of them leaves the warp converged and one does not. Does the source say which one you got?
Solution
At 1023 threads the block holds 32 warps and the last one carries 31 lanes, so its mask is 0x7fffffff. One thread short of 1024 costs almost nothing: both sizes occupy the same 32 warp slots and the difference is a single dead lane. The rounding hurts at the other end, where 33 threads take the same two slots as 64.
Moving the call below the branch can produce either answer. Predicate the if away and the warp never diverged, so the mask is unchanged. Branch, and the mask is a subset naming the lanes that arrived together, though the guide still refuses to promise which subset.
Nothing in the source picks between the two, and the answer can change when you edit an unrelated line.
The rule that outlives the exercise: __activemask() reports who is here, not who agreed with you. When you need the second thing, name the warp yourself and use a vote, which is what day 23 is for.
Pitfalls
Your 48-thread block is quietly a 64-thread block. Warp slots are allocated whole, so the SM reserves two of them either way and 16 lanes idle for the life of the block. Pick block sizes that are multiples of 32 unless something forces otherwise, and say why in a comment when it does. Day 10 measures what the choice costs on a real kernel.
You wrote a warp-synchronous reduction with no _sync and it passed. It passed on your card, this run. Per-thread program counters have been the model since compute capability 7.0, and code that assumes lockstep "should be revisited". Use __syncwarp() or a _sync primitive, and let day 14 show you what a barrier actually orders.
You used __activemask() to find out who took a branch. The guide bans that reading in as many words. The mask is an instantaneous snapshot for opportunistic warp-level programming, and the compiler is free to reorder around it. Compute the mask you mean, or use a warp vote, which day 23 introduces.
You recovered the execution order from printf. Print order is buffer order. Record what you want into a device array and read it back, which is what this lesson's program does and what day 64 explains.
You compared clock64() readings from two blocks. The counter is per multiprocessor, so two blocks on two SMs read two unrelated counters and their difference is neither a duration nor an ordering. Keep the comparison inside one block, or record the SM id, which needs inline PTX and arrives on day 46.
You put warpSize where a constant belongs. It is documented as a run-time value, so it cannot size a __shared__ array or appear in a static_assert. Write constexpr int kWarpSize = 32; for the compile-time uses, read warpSize once, and check that they agree instead of assuming it.
Go deeper
- CUDA Programming Guide 3.2.2.1 "SIMT Execution Model" and 3.2.2.1.1 "Independent Thread Scheduling", for lockstep and what replaced it: https://docs.nvidia.com/cuda/cuda-programming-guide/03-advanced/advanced-kernel-programming.html (checked 2026-08-30)
- CUDA Programming Guide 5.4.6.1 "Warp Active Mask", for
__activemask()and both warnings on it: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cpp-language-extensions.html (checked 2026-08-30) - CUDA C++ Programming Guide 7.13 "Time Function", the 12.6 archive, for what
clock64()counts: https://docs.nvidia.com/cuda/archive/12.6.0/cuda-c-programming-guide/index.html (checked 2026-08-30) - Turing Tuning Guide 1.4.1.1 "Instruction Scheduling", for the four warp schedulers per SM: https://docs.nvidia.com/cuda/turing-tuning-guide/index.html (checked 2026-08-30)
- Programming Massively Parallel Processors, 4th edition, chapter 4, on warps and SIMD hardware: https://shop.elsevier.com/books/programming-massively-parallel-processors/hwu/978-0-323-91231-0
Next
Day 22 takes the second reason a lane goes inactive, a branch that splits the warp, and measures what running both sides costs. Day 23 gives you the instructions that read across lanes on purpose, shuffles and votes, and the mask argument this page has been circling. Both sit in module 3, which spends ten days treating the warp as the unit rather than the thread.