G80 Memory Architecture
Review
The G80 is composed of several streaming multiprocessors (SM) as pictured below.
Each of these streaming multiprocessors has eight streaming processors (SP) and two super-function units (SFU). The SPs contain very simple ALUs for the most basic/common operations used in graphics processing while the SFUs are used when more complex operations are required (these should be avoided when possible). The SM performs multi-threaded instruction dispatch in vectors of 32. These vectors are referred to as warps in Nvidia terminology. There may be as many as 16 warps per a thread block for a total of 512 threads per a thread block. The large number of threads is used to hide memory access latency.
Memory Architecture
The memory hierarchy of the G80 can be seen above. Essentially, there are three levels available to the programmer – local registers (1K per SP), shared memory (16KB per SM), and global device memory. Note: The term ‘shared memory’ can be misleading. In the G80, shared means that the memory is available for use by various thread blocks. However, thread blocks are not able to use shared memory allocated to another thread block.
Registers
From the view of the programmer, there are 8KB of registers in each SM in the G80. The registers are dynamically partitioned across all thread blocks assigned to an SM. Once assigned, registers may not be accessed by threads in other thread blocks (similar to shared memory). However, there is an additional such that each thread may only access registers assigned to it.
As an example of register constraints, consider a matrix multiplication example. If each block has 16×16 (256) threads and each thread requires 10 registers, how many thread blocks can run on a single SM?
First, each thread block requires 256*10=2560 registers. There are 8192 registers available so we have 8192/2560=3 and some change. So with respect to register usage, three thread blocks can run on a single SM. What if each thread’s register requirement increases by one? Now we have 8192/2816=2 and a bunch of change so by using only one additional register, we’ve reduced the amount of parallelism to roughly 66%!
The dynamic partitioning of registers allows for an increased amount of flexibility available to the programmer – large number of threads using a small number of registers, a small number of threads using a large number of registers. This also gives the compiler the opportunity to optimize for instruction-level parallelism or thread-level parallelism.
Global Memory
There are two types of cacheable global memory, constant and texture. Constant values have the ability to be broadcast to all threads in a warp. Texture memory is optimized for 2-dimensional accesses. Both are read-only memories. Global memory can also be accessed for uncached read/write.
Shared Memory
Each SM has 16KB of shared memory. The shared memory is divided into 16 banks of 32-bit words. CUDA uses shared memory as storage area visible to all threads in a thread block (both readable and writeable). In parallel machines, it’s common for many threads to simultaneously access memory. To reduce contention and maximize bandwidth, the G80 divides shared memory into banks. Each of these banks can service one request pre-cycle. In addition, the shared memory can service as many simultaneous memory accesses as there are banks (as long as each access is to a unique bank. Multiple accesses to the same bank result in bank conflicts and result in a serialization of accesses (unless the accesses are to the same word which results in a broadcast of the word to all requestors). Successive 32-bit words are assigned to successive banks (i.e. share-memory-address % 16 == bank-number).
The illustration above shows two possible memory access patterns that avoid bank conflicts.
The illustration above shows two possible memory access patterns with bank conflicts. It’s worth noting that bank conflicts only happen within a single half-warp. If there are no bank conflicts, shared memory can be as fast as registers. There are additional memory banking examples in the slides.
Communication
How do threads communicate? Remember that this is a processor designed for graphics processing – specifically, data parallel streams representing independent vertices, triangles, fragments, pixels – these never communication with each other. However, there are some methods of communication when in compute mode, and these are designed for portability across GPU platforms. Threads belonging to the same thread block may communicate with each other through shared memory and execution kernels may communicate with each other through device memory. For synchronization, the only real way to synchronize threads in a single thread block is through an explicit CUDA call to __syncthreads(). Thread blocks are implicitly synchronized at the end of kernel execution. Remember that warps are scheduled out-of-order and you cannot and their scheduling cannot be depended upon.
