NVIDIA GPUs (IV)
Review- Streaming Multiprocessor (SM) G80 and Fermi
G80
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.
Fermi
Shown below is a Fermi SM.
Each Fermi SM contains 32 Streaming Processors or CUDA cores in NVIDIA terminology. With enhancement to support Double FP operations, now upto 16 Double FP FMA operations can be performed per SM. Other changes include 4 SFUs (again each SFU operations often are delivered at higher latency), 32K FP32 registers / SM. 64K of configurable Shared Memory and L1 cache. 2 Warp schedulers with 2 dispatch units and 16 LD/ST units.
Memory Architecture
The memory hierarchy of the G80 and Fermi (Fermi has 32 SP per SM instead of 8) is as shown above. Essentially, there are three levels available to the programmer – local registers (1K per SP), shared memory (16KB per SM in G80 and 64KB configurable Shared Memory/L1 cache in case of Fermi), 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. Fermi follow same principle except that the per-SM L1 cache is configurable to support both shared memory and caching of private (NVIDIA calls this local) and global memory operations.
Registers
From the view of the programmer, there are 8KB of registers in each SM in the G80 (32KB for Fermi). 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 compiler/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 (Software Managed Local Scratchpad)
In G80, each SM has 16KB of shared memory. The shared memory is divided into 16 banks of 32-bit words. In Fermi, each SM has 16/48/64KB of configurable shared memory in 32 banks organization. CUDA uses shared memory as storage area visible to all threads in a thread block (both readable and writeable).
Parallel Memory Architecture
In parallel machines, it’s common for many threads to simultaneously access memory. To reduce contention and maximize bandwidth, the G80/G92/Fermi 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. Additional Bank-Conflict examples are covered in section Bank Addressing Examples section below.
L1 cache (Texture+Constant) in G80
In G80, L1 cache is a Read only cache which holds Texture and Constant data and has its own name space. It is managed by TUs and only texture instructions can use it. Constants cache provides both Immediate address and Indexed address constants caching of constants stored in DRAM. It provides an extrememly efficient way of accessing a value that is common for all threads in a block. Textures represent 2D arrays of values stored in global DRAM which are cached into L1 and L2. Caches are optimized for 2D accesses and provide better memory performance for threads in a warp that follow 2D locality
Configurable L1/Shared Memory in Fermi
In Fermi, the Shared/L1 cache memory can be dynamically configured as either 48KB Shared Memory with 16KB L1 or 16 KB Shared with 48KB L1 cache in Fermi. When configured with 48KB of shared memory, program that make extensive use of shared memory can perform upto to 3x faster. In contrast, random memory access patterns unknown beforehand can be sped up by having 48KB L1 to improve performance over direct DRAM accesses. L1 Cache in Fermi is a normal R/W cache and coherency is maintained by Software).
L2 Cache
This a Read/Write Shared Cache between Fermi Cores. Fermi provides 768 KB of L2 cache.
Namespace before Fermi and Unified Address Space in Fermi
Before Fermi, Memory had (sort of) 6 name spaces:
Offchip - (1) Global (2) Private (CUDA: Local)
L1 - (3) Texture (4) Constant
(5)Shared Memory
(6) Registers
Fermi and PTX 2.0 ISA implement a unified address space across thread private local, block shared and global for LD/ST operations. This implementation enable Fermi to support C++ programs and permits 40-bit unified address space to support a TB of addressable memory !!.
Concept of Crossbar
A crossbar is an interconnect structure which provide efficient porting and transportation of data from multiple sources to multiple destinations. It however suffers bus-type conflicts since buses (crossbar is similar to bus) can sometime can be exposed to contention.
Here, PE 2 and 4 both require A and hence both get the multicasted data from Crossbar switching-based operation. Also, since access to B is on different bank, PE 3 also gets served. However, since access to B and C by both PE 1 and 3 produces bus conflict, it is not possible to deliver C in the same cycle as B even if data is available.
Bank Addressing Examples and Practices
Ex. 1 Linear Addressing
__shared__ float shared[256]; float foo = shared[baseIndex + s * threadIdx.x];
This is bank-conflict free if s has not common factors with Number of banks. In G80/G92, #BANKS=16 in Fermi #BANKS=32, therefore, s should be odd to avoid bank conficts.
Ex. 2 Data Types and Associated Bank conflicts
foo = shared[baseIndex + threadIdx.x];
This has no bank-conflicts if type of shared is 32-bit. However if the data type is smaller it produces bank conflicts. eg. - 4 way conflict if shared was 8 bits e.g.char - 2 way conflict if shared was 16 bits e.g. short.
Ex. 3 Structs
Struct assignments compile into as many memory accesses as there are struct members. Now, consider following definitions:
struct vector { float x, y, z; };
struct myType {
float f;
int c;
};
__shared__ struct vector vectors[64];
__shared__ struct myType myTypes[64];
-As we can see, there is no bank conflicts for struct vector since each member is of size 3 words i.e. 3 accesses per thread, contiguous banks (no common factor with 16). -However, 2-way bank conflicts arise in myType because of two accesses per thread.
Ex. 4 Common Array Bank Conflict Pattern - 1D
Consider following snippet:
int tid = threadIdx.x; shared[2*tid] = global[2*tid]; shared[2*tid+1] = global[2*tid+1];
Here each thread loads 2 elements to shared memory. As shown in figure, 2-way interleaved loads result in 2-way bank conflicts:
This may make sense in traditional CPU to improve locality in cache line usage and reduce shared traffic but since shared memory in GPUs is organized in banks where there is no cache, this results in bank-conflicts.
Thus, best array access pattern involves each thread loading one element (word) in every consecutive group of blockDim elements as illustrated below.
shared[tid] = global[tid]; shared[tid + blockDim.x] = global[tid + blockDim.x];
Ex. 5 Common Array Bank Conflict Pattern - 2D
Operations on 2D arrays are very common (e.g. Image processing). Consider processing of a 16×16 2D array as shown below.
Each thread is processing each element on one particular row, this causes each thread to access the elements in each column simultaneously (e.x. shows row 1 in purple). This causes a 16-way conflict since all rows start at bank 0 as shown.
There are two practiced solutions (both shown in same diagram below): (1) Pad the rows - Add one element at end of each row (waste memory) (2) Transpose before processing (suffer bank conflicts during Transpose by possibly amortize later).
Ex. 6 Load/Store Clustering/Batching
It is possible to restructure the sequence of LD/ST and batch them up to hide memory access latency within a thread. Thus, instead of interleaving LD/ST with compute instructions we can batch LD/ST occuring on different banks to hide latency well. e.g. Instead of doing,
LD0 (long latency) Dependent MATH0 LD1 (long latency) Dependent MATH1
Do
LD0 (long latency) LD1 (short latency) Dependent MATH0 Dependent MATH1
Here, instead of wasting time on accessing consecutive memory locations (on different banks) in interleaved fashion, its better to do LD/ST in batch. The compiler optimizes this batched operation and pipelines the access by firing both banks in consecutive cycles thus enabling better latency hiding.
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.
Atomic operations
Atomic operation are a powerful data-parallel way of manipulating shared data (think histogram here) and can also be used to form synchronization and mutex primitives.
Optimized Vector Reduction to avoid Bank Conflicts
Following figure produces an example of a element-wise Vector reduction which shows how the reduction operation creates bank-conflict for each operation.
This can be eliminated by restructuring the reduce operator as shown below:
Matrix Multiplication Example
One good approach of optimizing Matrix Multiplication on Parallel GPU architecture is by Application of Blocking. We allot each sub-block Psub of size BLOCK_SIZE of the operand matrices to each Thread block and each thread in the thread block is responsible for computing one element of Psub. (In figure, it is assumed that M and N multiples of BLOCK_SIZE).
Analysis -Each Sub-block Psub computation requires 2 x WIDTH 2 x 4 bytes of shared memory storage. (Here, it assumed that the results of matrix P are held in registers during computation. - Most Optimal model for Matrix Multiplication is:
( A ) Load Subblock of M and N ( B ) Sync (can use __syncthread();) ( C ) Compute Partial Result in register ( D ) Store
After doing (A) and (B), © knows that data is in SM local memory, hence instead of using global matrix data, local address are used in © to speed up communication.
Here latency hiding is possible by filling in the slack time of operation(A) of sub-block1 by doing operation© of available block0 brought in before. Thus, interleaving memory-bound and compute-bound regions enhances AI and provides better Occupancy (provided availability of parallel work!).
