Lecture 15: CUDA Case Studies. (3) Parallel Prefix Scan on the GPU. Using Multiple Streams in CUDA.

Lecture Summary

• Last time

  • Case studies: parallel reduction on the GPU & 1D convolution

  • Looking beyond today: some more GPU computing feature, but looking for a while into optimization features

• Today

  • One more cast study: parallel prefix scan

  • Using streams in GPU computing: increasing problem size; improving execution speeds

Parallel Prefix Scan on the GPU

Definition of the algorithm

Algo 1: Hillis & Steele (1986)

• Simple, but suboptimal (O(N*log2(N)))

Algo 2: Harris-Sengupta-Owen (2007)

• Convoluted, but O(N)

• Balanced trees: A common parallel algorithm pattern

  • Upsweep from roots to the main trunk, and then down sweep from trunk to root

  • "Tree": Just a concept--the actual data structure is not used

The reduction/upsweep step

The down sweep step. Sheesh, this is just...

CUDA Streams

• A CUDA-enabled GPU has 2 engines

  • An execution engine

  • A copy engine (which contains 2 sub-engines that can work simultaneously)

    • A H2D copy sub-engine

    • A D2H copy sub-engine

• Async execution

  • Examples: Kernel launches, D2D mem copies, mem copies by functions with the Async suffix, etc

• Overlapping Host <--> Device data transfer with device execution

  • Issue: The device execution stack is FIFO

    • Addressed by the usage of CUDA "streams"

• Concurrency can be managed through streams

  • Concurrency means one of two things:

    • The copy and the execution engines of GPU working at the same time

    • Several different kernels being executed at the same time on the GPU

• A stream is a sequence of CUDA commands issued by the host that executes on the GPU in issue-order

  • CUDA operations in different streams may run concurrently

  • CUDA operations from different streams may be interleaved

• As soon as a CUDA function is invoked, a default stream (stream 0) is created

• Create using cudaStreamCreate(), destroy using cudaStreamDestroy()