01 TRACE 02 SPLIT 03 OVERLAP 04 COMPLETE STEP 01 · TRACE
THE KERNEL IS FAST. THE QUEUE KEEPS WAITING. The program must process three independent chunks. The default Stream lines up input copy, Kernel, and result copy; only one kind of resource works at a time. SINGLE-QUEUE PIPELINE DEFAULT STREAM · 3 CHUNKS 01 for (int chunk = 0; chunk < 3; chunk++) {
02 cudaMemcpy(device, host, H2D);
03 filter<<<grid, block>>>(device);
04 cudaMemcpy(host, device, D2H);
05 }
DEFAULT STREAM READY 0 / 9
DEVICE TIMELINE 9 TASKS RUN IN SEQUENCE 01 02 03 04 05 06 07 08 09
H2D COPY C0 H2D IDLE IDLE C1 H2D IDLE IDLE C2 H2D IDLE IDLE
COMPUTE IDLE C0 KERNEL IDLE IDLE C1 KERNEL IDLE IDLE C2 KERNEL IDLE
D2H COPY IDLE IDLE C0 D2H IDLE IDLE C1 D2H IDLE IDLE C2 D2H
RUN SINGLE QUEUE▶ SPLIT THE WORK QUEUES→