1 / 27

Control Flow/ Thread Execution

Control Flow/ Thread Execution. 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

thanos
Download Presentation

Control Flow/ Thread Execution

An Image/Link below is provided (as is) to download presentation Download Policy: Content on the Website is provided to you AS IS for your information and personal use and may not be sold / licensed / shared on other websites without getting consent from its author. Content is provided to you AS IS for your information and personal use only. Download presentation by click this link. While downloading, if for some reason you are not able to download a presentation, the publisher may have deleted the file from their server. During download, if you can't get a presentation, the file might be deleted by the publisher.

E N D

Presentation Transcript

  1. Control Flow/ Thread Execution © David Kirk/NVIDIA and Wen-mei W. Hwu, 2007-2009 ECE 498AL, University of Illinois, Urbana-Champaign

  2. 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. © David Kirk/NVIDIA and Wen-mei W. Hwu, 2007-2009 ECE 498AL, University of Illinois, Urbana-Champaign

  3. 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 physicallyin parallel in G80/GT200 • 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 kernel call of the program can take effect © David Kirk/NVIDIA and Wen-mei W. Hwu, 2007-2009 ECE 498AL, University of Illinois, Urbana-Champaign

  4. Grids and Blocks • A kernel is executed as a grid of thread blocks • All threads share global memory space • A thread block is a batch of threads that can cooperate with each other by: • Synchronizing their execution using barrier • Efficiently sharing data through a low latency shared memory • Two threads from two different blocks cannot cooperate © David Kirk/NVIDIA and Wen-mei W. Hwu, 2007-2010 ECE 408, University of Illinois, Urbana-Champaign

  5. CUDA Thread Block: Review • Programmer declares (Thread) Block: • Block size 1 to 512 concurrent threads • Block shape 1D, 2D, or 3D • Block dimensions in threads • All threads in a Block execute the same thread program • Threads share data and synchronize while doing their share of the work • Threads have thread id numbers within Block • Thread program uses thread id to select work and address shared data CUDA Thread Block Thread Id #:0 1 2 3 … m Thread program Courtesy: John Nickolls, NVIDIA © David Kirk/NVIDIA and Wen-mei W. Hwu, 2007-2010 ECE 408, University of Illinois, Urbana-Champaign

  6. How thread blocks are partitioned • Thread blocks are partitioned into warps • Thread IDs within a warp are consecutive and increasing. If block is 1D then threadID=threadIdx.x • Warp 0: thread 0, …thread 31; • Warp 1: thread 32 ,…, thread 63 • Warp n: thread 32*n ,…., thread 32(n+1)-1 • For blocks with # of threads not a multiple of 32 • Last warp padded with extra threads to fill up the 32 threads • e.g. a block: 48 threads; 2 warps; warp 1 padded with 16 extra threads © David Kirk/NVIDIA and Wen-mei W. Hwu, 2007-2009 ECE 498AL, University of Illinois, Urbana-Champaign

  7. How thread blocks are partitioned: 2D Thread Blocks 3D Thread blocks are portioned in a similar way (i.e. threads with ThreadIDx.z=1 follow threads withThreadIDx.z=0 and so on © David Kirk/NVIDIA and Wen-mei W. Hwu, 2007-2009 ECE 498AL, University of Illinois, Urbana-Champaign

  8. How thread blocks are partitioned (cont.) • Partitioning is always the same • Thus you can use this knowledge in control flow • The exact size of warps may change from generation to generation of devices • However, DO NOT rely on any ordering between warps • If there are any dependencies between threads, you must __syncthreads() to get correct results © David Kirk/NVIDIA and Wen-mei W. Hwu, 2007-2009 ECE 498AL, University of Illinois, Urbana-Champaign

  9. Control Flow Instructions • Hardware executes an instruction for all threads in the same wrap before moving to next instruction • SIMT: single instruction, multiple thread execution • Works well when all threads in a wrap follow the same control flow path • Main performance concern with branching is divergence; i.e. threads within a single warp take different paths • Different execution paths are serialized • Example: if-then-else executed in two passes • One pass for threads executing the then path • Second pass for threads executing the else path • The control paths taken by the threads in a warp are traversed one at a time until there is no more. © David Kirk/NVIDIA and Wen-mei W. Hwu, 2007-2009 ECE 498AL, University of Illinois, Urbana-Champaign

  10. Control Flow Instructions (cont.) • 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 threads in any given warp follow the same path © David Kirk/NVIDIA and Wen-mei W. Hwu, 2007-2009 ECE 498AL, University of Illinois, Urbana-Champaign

  11. Parallel Reduction • Given an array of values, “reduce” them to a single value in parallel • Examples • sum reduction: sum of all values in the array • Max reduction: maximum of all values in the array • Typically parallel implementation: • Recursively halve # threads, add two values per thread • Takes log(n) steps for n elements, requires n/2 threads © David Kirk/NVIDIA and Wen-mei W. Hwu, 2007-2009 ECE 498AL, University of Illinois, Urbana-Champaign

  12. A Vector Reduction Example • Assume an in-place reduction using shared memory • The original vector is in device global memory • The shared memory used to hold a partial sum vector • Each iteration brings the partial sum vector closer to the final sum • The final solution will be in element 0 © David Kirk/NVIDIA and Wen-mei W. Hwu, 2007-2009 ECE 498AL, University of Illinois, Urbana-Champaign

  13. A simple implementation • Assume we have already loaded array into __shared__ float partialSum[] unsigned int t = threadIdx.x; for (unsigned int stride = 1; stride < blockDim.x; stride *= 2) { __syncthreads(); if (t % (2*stride) == 0) partialSum[t] += partialSum[t+stride]; } © David Kirk/NVIDIA and Wen-mei W. Hwu, 2007-2009 ECE 498AL, University of Illinois, Urbana-Champaign

  14. Vector Reduction with Bank Conflicts Array elements 0 1 2 3 4 5 6 7 8 9 10 11 1 0+1 2+3 4+5 6+7 8+9 10+11 2 0...3 4..7 8..11 3 0..7 8..15 iterations © David Kirk/NVIDIA and Wen-mei W. Hwu, 2007-2009 ECE 498AL, University of Illinois, Urbana-Champaign

  15. Vector Reduction with Branch Divergence Thread 0 Thread 2 Thread 4 Thread 6 Thread 8 Thread 10 0 1 2 3 4 5 6 7 8 9 10 11 1 0+1 2+3 4+5 6+7 8+9 10+11 2 0...3 4..7 8..11 3 0..7 8..15 iterations Array elements © David Kirk/NVIDIA and Wen-mei W. Hwu, 2007-2009 ECE 498AL, University of Illinois, Urbana-Champaign

  16. Some Observations • In each iteration, two control flow paths will be sequentially traversed for each warp • Threads that perform addition and threads that do not • Threads that do not perform addition may cost extra cycles depending on the implementation of divergence • No more than half of threads will be executing at any time • All odd index threads are disabled right from the beginning! • On average, less than ¼ of the threads will be activated for all warps over time. • Starting with the 5th iteration, entire warps in each block will be disabled, poor resource utilization but no divergence. • This can go on for a while, up to 4 more iterations (512/32=16= 24), where each iteration only has one thread activated until all warps retire © David Kirk/NVIDIA and Wen-mei W. Hwu, 2007-2009 ECE 498AL, University of Illinois, Urbana-Champaign

  17. Short comings of the implementation • Assume we have already loaded array into • __shared__ float partialSum[] unsigned int t = threadIdx.x; for (unsigned int stride = 1; stride < blockDim.x; stride *= 2) { __syncthreads(); if (t % (2*stride) == 0) partialSum[t] += partialSum[t+stride]; } BAD: Divergence due to interleaved branch decisions © David Kirk/NVIDIA and Wen-mei W. Hwu, 2007-2009 ECE 498AL, University of Illinois, Urbana-Champaign

  18. A better implementation • Assume we have already loaded array into • __shared__ float partialSum[] unsigned int t = threadIdx.x; for (unsigned int stride = blockDim.x >> 1; stride > 1; stride >> 1) { __syncthreads(); if (t < stride) partialSum[t] += partialSum[t+stride]; } © David Kirk/NVIDIA and Wen-mei W. Hwu, 2007-2009 ECE 498AL, University of Illinois, Urbana-Champaign

  19. Execution of the revised algorithm (If blockDim.x = 32 ) Thread 0 0 1 2 3 … 13 14 15 16 17 18 19 1 0+16 15+31 3 4 © David Kirk/NVIDIA and Wen-mei W. Hwu, 2007-2009 ECE 498AL, University of Illinois, Urbana-Champaign

  20. No Divergence using revised algorithm • IfblockDim.x = 512; • In 1st iteration • Threads 0 through 255 execute add • Threads 256 through 511 do not • Pair-wise sums stored in elements 0 - 255 after • All threads in warps 1 - warp 8 execute add statement • Warps 9 - warp 15 execute all skip the add • All threads in each warp take the same path • There is no thread divergence! © David Kirk/NVIDIA and Wen-mei W. Hwu, 2007-2009 ECE 498AL, University of Illinois, Urbana-Champaign

  21. Some Observations About the New Implementation • Only the last 5 iterations will have divergence • Entire warps will be shut down as iterations progress • For a 512-thread block, 4 iterations to shut down all but one warps in each block • Better resource utilization, will likely retire warps and thus blocks faster © David Kirk/NVIDIA and Wen-mei W. Hwu, 2007-2009 ECE 498AL, University of Illinois, Urbana-Champaign

  22. A Potential Further Refinement • For last 6 loops only one warp active (i.e. tid’s 0..31) • Shared reads & writes are SIMD synchronous within a warp So skip __syncthreads() and unroll last 6 iterations unsigned int t = threadIdx.x; for (unsigned int stride = blockDim.x >> 1; stride > 1; stride >> 1) { __syncthreads(); if (t < stride) partialSum[t] += partialSum[t+stride]; } © David Kirk/NVIDIA and Wen-mei W. Hwu, 2007-2009 ECE 498AL, University of Illinois, Urbana-Champaign

  23. A Potential Further Refinement but bad idea unsigned int tid = threadIdx.x; for (unsigned int d = n>>1; d > 32; d >>= 1) { __syncthreads(); if (tid < d) shared[tid] += shared[tid + d]; } __syncthreads(); if (tid < 32) { // unroll last 6 predicated steps; Shared reads //and writes are SIMD synchronus within a wrap shared[tid] += shared[tid + 32]; shared[tid] += shared[tid + 16]; shared[tid] += shared[tid + 8]; shared[tid] += shared[tid + 4]; shared[tid] += shared[tid + 2]; shared[tid] += shared[tid + 1]; } This would not work properly if warp size decreases; need __synchthreads() between each statement! However, having ___synchthreads() in if statement is problematic. © David Kirk/NVIDIA and Wen-mei W. Hwu, 2007-2009 ECE 498AL, University of Illinois, Urbana-Champaign

  24. Predicated Execution Concept • <p1> LDR r1,r2,0 • If p1 is TRUE, instruction executes normally • If p1 is FALSE, instruction treated as NOP © David Kirk/NVIDIA and Wen-mei W. Hwu, 2007-2009 ECE 498AL, University of Illinois, Urbana-Champaign

  25. Predication Example : : if (x == 10) c = c + 1; : : : : LDR r5, X p1 <- r5 eq 10 <p1> LDR r1 <- C <p1> ADD r1, r1, 1 <p1> STR r1 -> C : : © David Kirk/NVIDIA and Wen-mei W. Hwu, 2007-2009 ECE 498AL, University of Illinois, Urbana-Champaign

  26. Predication very helpful for if-else A A B C D B C D © David Kirk/NVIDIA and Wen-mei W. Hwu, 2007-2009 ECE 498AL, University of Illinois, Urbana-Champaign

  27. If-else example : : p1,p2 <- r5 eq 10 <p1> inst 1 from B <p1> inst 2 from B <p1> : : <p2> inst 1 from C <p2> inst 2 from C : : : : p1,p2 <- r5 eq 10 <p1> inst 1 from B <p2> inst 1 from C <p1> inst 2 from B <p2> inst 2 from C <p1> : : schedule The cost is extra instructions will be issued each time the code is executed. However, there is no branch divergence. © David Kirk/NVIDIA and Wen-mei W. Hwu, 2007-2009 ECE 498AL, University of Illinois, Urbana-Champaign

More Related