()

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.

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.

GPU Control Flow

Think of threads as lanes of vector.
Masking allows you to enable/disable an instruction for a specific lane.
Masking allows control flow to diverge.

Control Flow Divergence

We would like to see all lanes executing in parallel.
However, they will not if control diverges.
The example in the slides shows that threads that satisfy the conditions of an “if” will execute in parallel. This example has execution that is between completely serial and fully parallel.

Mask Stack Enables Divergence

The mask stack allows nesting of control statements by saving enable masks when a control statement is encountered and restoring the saved mask when leaving the control statement body.
Initially, all lanes are enabled (all bits of the mask are equal to 1) and the mask stack is empty.
Lanes with a 1 in the mask get to execute, and lanes with a 0 don’t.
The slides have an animated example of how enable masks are saved and restored using the mask stack.
The stack and the enable mask are maintained by hardware.
Each SM has its own stack that can handle 4 deep nesting (i.e. the stack is 4 deep).
More than 4 deep nesting will probably not run.

Predication

Predicates can replace branches like “if” statements, by enabling/disabling specific instructions based on a conditional value calculated in each lane.
- This is similar to explicitly calculating the enable mask with software.
Using predication will result in the same threads executing as divergence, except that predication executes both paths of branches.
- Divergence only executes the chosen path, so if all lanes happen to go one way, you won’t waste computation by doing the other path.
The example in the slides shows how to do the same branching with predicates as with control divergence.

Lab 1 Review

Question 1 Breakdown

Question 1 dealt with the execution of the matrix multiply code using TSCTest.cpp to measure clock cycles.

Why take several measurements?

Because we are running on shared machines and OS has interrupts.
It was unexpected that the machines were more heavily loaded than when Mattan ran the tests, resulting in more outside influence on the test runs.

Which measurements should you report? minimum, maximum, or mean?

Minimum or mean were acceptable if reasonable explanation was given.
The minimum represents a measurement with the least interference from the OS and other users.
The mean represents typical performance on a realistic machine that has OS interrupts and\or other users.

Why does the 32X32 matrix problem size perform much worse on the first measurement?

The first run warms up the cache for the runs after it; so, the first run does not benefit from having a warmed up cache.

How can you make the runs the same?

Throw out the first run because it performs worse without a warmed up cache.
Reallocate the cache memory each time. Several groups did this to begin with and did not see the difference between the 1st and subsequent runs.
Flush the cache each run. This strategy is tricky to implement and requires use of explicit cache flush instructions to ensure a full flush.

Should you do it (make them the same)?

It depends on the usage because you could have a case where you rerun matrix multiplies several times with only a few values changed, so warming up the cache is a good and relevant thing to do.

Optimization

The following optimizations were done by the class overall:

registers
blocking
transpose B
transpose 1 block at a time
use SSE to get better register locality
prefetching + SWP
copying blocks

Register and blocking optimization were the most common and best optimizations to use. SSE was a good optimization, and transposing one block was unexpectedly good (more on this later).

Blocking Technique

The best technique was to block twice, once for the L1 cache and once for the L2 cache (and a third time for registers as explained below).

This minimizes misses from both caches.
L1 locality is important but small (~16KB).
L2 is 1MB, so it can hold more blocks.
Hierarchical blocking makes the transpose optimization irrelevant because ordering doesn’t matter when blocking.
Doing 2b³ computations results in 2b² misses for bringing in blocks.
The transpose optimization unexpectedly had good performance, but the better performance occurred when hierarchical blocking was not used.

Register Optimization Techniques

Done explicitly when using SSE.

Unroll loops 4–32 iterations. Mattan expected an explanation why a specific number of iterations was chosen.
Explicitly optimize through SSE.
Basically another level of blocking but cannot do explicitly because of the compiler.

PIN Question

PIN’s dcache counts cache hits and misses. It was modified to have two levels of cache.
Correction to dcache: Multi-line accesses were only counted as one access total. To be correct, multi-line accesses should count one access for each line accessed.
Improvement to dcache: Dirty lines evicted from the L1 cache should go to the L2 cache. This affects energy consumption.

Counting Register Accesses

Method 1: Look at the assembly code.
Method 2: Estimate based on source code (C code).
Many groups had a large disparity in number of register accesses between algorithms, which was not expected.
The number of register accesses should not have been less than the number of computations.

Algorithm Trend in PIN

Cache oblivious should behave better when changing cache size than cache aware (that doesn’t change block size).

CACTI Energy

You should have explained what number you chose and why (from CACTI energy results).
Blocking should give an energy reduction.
Leakage energy goes down because execution time goes down.
Shorter execution time leads to more power consumption.

Analytical Modeling

You should have a model to compare with measurements to see if they are consistent. Inconsistencies means something went wrong, either with the model or the measurements.

{$ \begin{align} Locality &= \frac {number\_of\_accesses\_required\_for\_computation} {number\_of\_words\_from\_a\_level\_of\_storage} \end{align} $}

c += a * b breaks down into:

t = a * b
c = c + t
This results in 6n³ accesses.

{$ \begin{align} hit\_rate &= 1 - miss\_rate \end{align} $}, (miss rate is easier to compute)

{$ \begin{align} accesses &= 6b^2N \left(\frac {N^2}{b^2}\right) \end{align} $} across the blocks of C, each of size b².

{$ \begin{align} misses &= (b^2 + 2bN) \left(\frac {N^2}{b^2}\right) \end{align} $}

To minimize misses, a small b should be used with a small cache and a large b should be use with a large cache.

{$ \begin{align} 3b^2 &= Z \end{align} $}, where Z is the cache size

If you consider line size, you just divide one of the factor of b misses by L.

Extension to multilevel memory hierarchy

The equation is the same.
The only thing that changes is the number of accesses to a level, which are the misses from the next higher level of memory.

{$ \begin{align} b &= \sqrt {\frac {Z}{3}} \end{align} $}

Therefore, {$ \begin{align} \frac {Z}{3} + 2 \left(\sqrt{\frac {Z}{b}}\right) \left(\frac {N^3}{\frac {Z}{3}}\right) \end{align} $}.