lecture10-control-flow fall 2007.pptVIP

ECE 498ALLecture 10: Control Flow Objective To understand the implications of control flow on Branch divergence overhead SM execution resource utilization To learn better ways to write code with control flow To understand compiler/HW predication designed to reduce the impact of control flow There is a cost involved. Quick terminology review Thread: concurrent code and associated state executed on the CUDA device in parallel with other threads The unit of parallelism in CUDA Warp: a group of threads executed physically in parallel in G80 Block: a group of threads that are executed together and form the unit of resource assignment Grid: a group of thread blocks that must all complete before the next phase of the program can begin How thread blocks are partitioned Thread blocks are partitioned into warps Thread IDs within a warp are consecutive and increasing Warp 0 starts with Thread ID 0 Partitioning is always the same Thus you can use this knowledge in control flow However, the exact size of warps may change from generation to generation Covered next However, DO NOT rely on any ordering between warps If there are any dependencies between threads, you must __syncthreads to get correct results Control Flow Instructions Main performance concern with branching is divergence Threads within a single warp take different paths Different execution paths are serialized in G80 The control paths taken by the threads in a warp are traversed one at a time until there is no more. A common case: avoid divergence when branch condition is a function of thread ID Example with divergence: If threadIdx.x 2 This creates two different control paths for threads in a block Branch granularity warp size; threads 0 and 1 follow different path than the rest of the threads in the first warp Example without divergence: If threadIdx.x / WARP_SIZE 2 Also creates two different control paths for threads in a block Branch granularity is a whole multiple of warp size; all thr