- 01OBSERVE
- 02COALESCE
- 03STAGE
- 04COMPLETE
STEP 01 · OBSERVE
THE THREADS LEFT. THE DATA DIDN'T ARRIVE.
Lesson 00 created 256 Threads. Now every Warp reads memory with a stride: compute is ready while scattered data is still in flight.CURRENT KERNEL256 THREADS · STRIDE 1601__global__ void transform(float* input) {
02int i = blockIdx.x * blockDim.x + threadIdx.x;
03float value = input[(i * 16) % 256];
04output[i] = value * 2.0f;
05}
GPU READY
WARP 00 · MEMORY REQUESTS32 THREADS TOUCH 16 SCATTERED REGIONST00→ 000R00T01→ 016R01T02→ 032R02T03→ 048R03T04→ 064R04T05→ 080R05T06→ 096R06T07→ 112R07T08→ 128R08T09→ 144R09T10→ 160R10T11→ 176R11T12→ 192R12T13→ 208R13T14→ 224R14T15→ 240R15T16→ 000R00T17→ 016R01T18→ 032R02T19→ 048R03T20→ 064R04T21→ 080R05T22→ 096R06T23→ 112R07T24→ 128R08T25→ 144R09T26→ 160R10T27→ 176R11T28→ 192R12T29→ 208R13T30→ 224R14T31→ 240R15
ACTIVE THREADS32 / WARP
REGIONS TOUCHED16
WAIT STATEMEMORY STALL