A warp is a group of 32 threads that execute together in lockstep. It is the real unit of scheduling on NVIDIA hardware. Blocks are a programming convenience, but the hardware splits every block into warps and issues instructions one warp at a time.

// a 100-thread block still costs 4 warps: 32 + 32 + 32 + 4
kernel<<<1, 100>>>();   // the last warp runs 28 threads idle

This is why block sizes should be multiples of 32. A block of 100 threads occupies the same issue slots as a block of 128, and 28 lanes do nothing.

Every thread in a warp shares one program counter and executes the same instruction each cycle, on different data. That is the SIMT model, and it is close to SIMD on a CPU, except that each lane has its own registers and can be individually masked off. When threads in a warp take different branches, the hardware handles it by masking rather than by scheduling, which costs time (Warp divergence (CUDA)).

Warps are also the granularity of latency hiding. When a warp stalls waiting on a global memory load, the multiprocessor issues from another resident warp instead of idling. This is the entire reason GPUs tolerate memory latency measured in hundreds of cycles, and it only works if enough warps are resident (Occupancy (CUDA)).

Threads within a warp can exchange registers directly with the shuffle instructions, without going through shared memory:

int v = __shfl_down_sync(0xffffffff, value, 16);  // read from lane+16

The mask argument names which lanes take part. Since the Volta architecture, warps can diverge and reconverge independently, so the mask has to be explicit rather than assumed.