�GPU Architecture

Kevin Garner & Spiros Tsalikis

1

Online at:

GPU Problem Breakdown

2

Whole Mesh/Matrices

CPU Thread

GPU Blocks

GPU Threads

GPU Threads

• Threads should be allocated according to problem size
• Exploit more parallelism with more blocks of smaller size (1,024 blocks with 512 threads per block vs. 2,048 blocks with 256 threads per block)

​

• Thread block may contain up to 1024 threads

​

• Threads within the same block can access shared memory data and synchronize their execution with __syncthreads() which acts as a barrier.

3

J. Cheng, M. Grossman and T. McKercher, Professional CUDA C Programming, Indianapolis, IN, USA: Wiley, 2014.

GPU Thread Memory Hierarchy

4

GPU Thread Memory Access

• In addition to global memory, there are two additional memory spaces accessible by all threads that are optimized for different memory usages:
• Constant Memory

​

• Texture Memory

5

Constant Memory

• Accesses to different addresses by threads within a warp are serialized, thus the cost scales linearly with the number of unique address reads

​

• Total of 64 KB

​

• The Constant cache is useful when threads within a warp repeatedly access the same, few address locations, allowing constant memory access to become as fast as register access

​

6

Texture Memory

• Optimized for 2D spatial locality (texture addresses that are close together will give best performance)

​

• Designed for streaming fetches with a constant latency; that is, a cache hit reduces DRAM bandwidth demand, but not fetch latency

​

7

Asynchronous Concurrent Execution

• Using asynchronous calls, device operations are queued and executed by the CUDA driver, relieving the host of the responsibility of managing the device.
• The following device operations are asynchronous with respect to the host:
• Kernel launches
• Memory copies within a single device’s memory
• Memory copies from host to device of a memory block of 64 KB or less
• Memory set function calls

8

Streams

• A sequence of commands (possibly issued by different host threads) that execute in order
• Used to manage concurrent operations
• A stream does not execute until all dependencies are met.
• The successful completion of a synchronize call guarantees that all stream commands have completed.
• Default stream is used and commands are executed in given order when stream parameter is not specified (e.g., during kernel launches and host <-> device memory copies)
• Explicit Synchronization
• Create streams with priorities using cudaStreamCreateWithPriority(), pending work with higher-priority streams take preference over low-priority streams

9

Cuda Graphs

• A new model for work submission.
• A series of operations connected by dependencies.
• Define the graph once and then launch it repeatedly.
• CPU launches of graphs can be faster than streams because most of the setup is done in advance, plus internal CUDA optimizations.
• A node can be:
• kernel
• CPU function call
• Mem Copy
• Mem set
• waiting an event
• semaphores

10

Events & Synchronous Calls

• Events can be used to:
• measure the device’s progress asynchronously
• synchronize a kernel call

​

• The control of a synchronous function which has been called won’t be returned to the host thread till it’s finished. A host thread can be defined to yield, block, or spin using cudaSetDeviceFlags(), before any other CUDA call is performed by the host.

11

Hardware Implementation: 1

• A Nvidia GPU has a scalable array of multithreaded Streaming Multiprocessors (SMs).
• A grid of thread blocks is distributed across SMs
• A SM can execute hundreds of threads concurrently.
• A SM employs the SIMT (Single-Instruction, Multiple-Thread) architecture.
• The multiprocessor creates, manages, schedules, and executes threads in groups of 32 parallel threads called warps.

12

Hardware Implementation: 2

• When a SM is given one or more thread blocks to execute, it partitions them into warps and each warp gets scheduled by a warp scheduler for execution.

​

• The way a block is partitioned into warps is always the same; each warp contains threads of consecutive, increasing thread IDs

​

• The total number of warps in a block: ceil(T / Wsize)

13

More to Research

• Unified Memory
• Dynamic Parallelism
• Profiling
• Execution Configuration Optimization (e.g. occupancy)

14

Additional References

• https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html
• https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html

​

15