What is a thread block in CUDA?
A group of threads that lands on one SM, shares that SM's shared memory, and can synchronize with __syncthreads().
Two threads in the same block can hand each other data. Two threads in different blocks cannot, and that asymmetry is the whole reason the level exists. The scheduler gives a block to one streaming multiprocessor and that is where it stays, which is what makes shared memory and __syncthreads() buildable at all: both need every participant on the same piece of silicon at the same moment. Across blocks you get neither. NVIDIA's guide is blunt about it: which blocks run on which SM "cannot be controlled or queried by the application and no ordering guarantees are made by the scheduler".
The hardware caps sit in three places and only one of them is the number you type. A block cannot exceed 1024 threads on any GPU this course targets. A T4 SM holds at most 16 resident blocks and at most 1024 resident threads, so 16 blocks of 64 and 4 blocks of 256 and 1 block of 1024 all fill it, and a 2048-thread block is not a legal way to fill it twice. The third cap is shared memory: 64 KiB per SM, but only 49,152 bytes per block by default. Day 13 asked for 53,248 bytes of dynamic shared memory and the launch came back invalid argument until cudaFuncSetAttribute opted in to the 65,536-byte ceiling.
The mistake worth naming is treating a block as a unit that runs in lockstep. It does not. A block is several warps, and only a warp is issued together, so a block-wide algorithm that skips its barrier is correct only while the block is one warp wide. Day 14 measured that exactly: a tile reverse written without __syncthreads() produced zero mismatches at 32 threads per block and 17,312 at 64. The bug is not a rare race, it is the default, and a 32-thread block hides it completely. Worse, the barrier does not announce a missing participant either: with an early return, only 96 of 256 threads reached it and the launch returned rather than deadlocking.
Everything else about a block is a choice you make on the host, and it changes who owns what. The same 2000-element array under 256-thread blocks gives element 1337 to block 5, thread 57; under 1024-thread blocks it goes to block 1, thread 313. Both launches start 2048 threads and both leave 48 of them with nothing to do. Day 10 is where that choice stops being arbitrary and starts being measured.
Measured
Everything below ran on the project's verification node on 2026-08-30: a Tesla T4 at compute capability 7.5, driver 595.84, CUDA 12.6 (V12.6.85), built with nvcc -O3 -arch=sm_75.
Day 4 ran 2000 elements through three launches:
| Launch | Threads started | Unwritten | Idle | Owner of element 1337 |
|---|---|---|---|---|
<<<7, 256>>> |
1792 | 208 | 0 | block 5, thread 57 |
<<<8, 256>>> |
2048 | 0 | 48 | block 5, thread 57 |
<<<2, 1024>>> |
2048 | 0 | 48 | block 1, thread 313 |
Day 14 took the sharing side and reversed a tile inside a block without a barrier, at four block sizes:
| Threads per block | Mismatches, no barrier | Mismatches, with barrier |
|---|---|---|
| 32 | 0 | 0 |
| 64 | 17,312 | 0 |
| 128 | 21,696 | 0 |
| 256 | 46,656 | 0 |
That first row is the trap this entry exists for. A block of 32 is one warp, one warp shares an instruction stream, and a missing barrier costs nothing until the block is wide enough to hold two. A test suite that only ever launches 32-thread blocks passes on code that is wrong.
Day 13's shared-memory cap on the same card: a 32x32 tile of floats is 4,096 bytes per block, comfortably inside the 49,152-byte default. Ask for 53,248 and the launch is refused with invalid argument, which is cudaErrorInvalidValue, code 1.
Diagram
index-tracer, preset 1d-basic, with block boundaries drawn as the outer boxes: three blocks of four threads over twelve elements, and the thread number restarting at zero at each boundary while the element number does not.
Alt text: "Three blocks of four threads over twelve elements. Thread numbers run 0 to 3 inside every block and the element numbers run 0 to 11, so the same thread number appears three times and the block number is what tells them apart."
Related terms
- grid
- streaming multiprocessor
- shared memory
__syncthreads()- resident warps per SM
- warp
- built-in index variables
- occupancy
Where you meet this
- Day 4, grid, block and thread indexing, which owns this term
- Day 10, choosing threads per block, where the block size gets measured instead of guessed
- Day 13, shared memory and tiling, for the memory a block gets to itself
- Day 14, what
__syncthreads()guarantees, which produced the mismatch table above invalid configuration argument, the synchronous refusal you get when a block asks for more than 1024 threads
Sources
- CUDA Programming Guide, on block scheduling ("no ordering guarantees are made by the scheduler, so programs cannot rely on a specific scheduling order or scheme for correct execution") and on
__syncthreads()blocking "all threads in the thread block": https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/writing-cuda-kernels.html (checked 2026-08-30) - CUDA Programming Guide, hardware multithreading, "When an SM is given one or more thread blocks to execute, it partitions them into warps": https://docs.nvidia.com/cuda/cuda-programming-guide/03-advanced/advanced-kernel-programming.html (checked 2026-08-30)
- CUDA Programming Guide, Table 30, 1024 maximum threads per block, and for compute capability 7.5 a maximum of 16 resident blocks and 32 resident warps per SM: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/compute-capabilities.html (checked 2026-08-30)
- CUDA Runtime API, the
cudaErrorenum, wherecudaErrorInvalidValue = 1: https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__TYPES.html (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.