University:
Boston University
Course:
GRS CS 910 | Directed Study: Systems
Academic year:

2021
Views:

342

Pages:

9
Author:

Josue Stephenson

GPU Computing with CUDA More Optimizations 1 Outline of lecture ‣ Recap ‣ Further optimization techniques ‣ Instruction optimization ‣ Memory as a limiting factor ‣ Thread and block heuristics 2 Recap ‣ Control flow - Hardware serializes divergent thread execution within same warp (if else instructions) - All threads of warp have to complete execution before warp ends ‣ Loop divergence may drag run time of the whole warp - Recommendations ‣ Try to make all threads of same warp do the same thing ‣ If you have branching code, cut your branch between warps 3 Recap ‣ Memory coalescing - Memory is accessed in chunks of 32, 64 or 128 bytes - Protocol ‣ Find memory segment that contains address requested by the lowest numbered active thread ‣ Find other active threads whose requested address lies in same segment, and reduce transaction size if possible ‣ Do transaction. Mark serviced threads as inactive. ‣ Repeat until all threads are serviced 4 Recap ‣ Memory coalescing - Recommendations - Have contiguous data storage patterns so to waste the least amount of bandwidth possible - Try to use a SoA programming model - Take advantage of cudaMalloc that guarantees contiguous data in memory - If your problem is compute bound don’t worry to much! ‣ Latency hiding - Do as much work per memory transaction as possible 5 Further optimization techniques ‣ Data prefetching - Can be used with tiling to hide latency Loop { Load tile to shared memory __syncthreads(); Load first tile to registers from global memory Loop { Load tile to shared memory from registers Compute tile; __syncthreads(); __syncthreads; Load next tile to registers from global memory } Compute tile; No prefetching __syncthreads; } With prefetching 6 Further optimization techniques ‣ Loop unrolling - Loops involve branching code that slows things down! for (int k=0; k

GPU Computing with CUDA - More Optimization