Review: Multiply using several blocks
How addresses map to banks on G80
- G80 has 16 banks of shared memory
- Each bank has a bandwidth of 32 bits per SRAM clock cycle
- Successive 32-bit words are assigned to successive banks
- An address will map to a bank number computed as bank # = address % 16
- This allows the possibility for zero bank conflicts
- Each SM has 8 core SPs (Streaming processors)
- GTS 8800 has 128 SPs in total
- The core clock is 650 MHz while SPs run at 1350 MHz
- Therefore, 2 cycles per SM are required, fewer than that won’t support the required BW
Fermi
- Fermi has 32 banks of shared memory
- It has 16 load store units
These numbers do change with advancement in versions. Design is always first optimized for graphics. To add a feature, market should be there. That’s because, if one needs a feature, and if that feature is incorporated, all will have to share the cost of adding that feature.
- Warps are never explicitly exposed to the programmers in the CUDA programming model, but they’re there
- One should focus on managing thread performance
- One has to only ensure that no bank conflicts are there within a warp
- Any bank conflict within a warp will slowdown access to shared memory
- One instruction cannot have multiple load/store.
- Also, loads are not non*blocking
Shared Memory Bank Conflicts
As long as there are no bank conflicts, shared memory acts as fast as registers.
- If all threads of a half*warp access different banks, there is no bank conflict
- If all threads of a half*warp access the identical address, even then there is no bank conflict (broadcast)
- However, if multiple threads in the same half*warp access the same bank, it’ll lead to bank conflicts
- Accesses would be serialized
- The cost of such serialized accesses will be determined by the maximum number of simultaneous accesses to a single bank
Data Type & Bank Conflicts
Consider the arrangement as shown above where there are 16 threads and 16 banks. Depending on the datatype in use, there can be different possibilities, some of them leading to bank conflicts.
foo = shared[baseIndex + threadIdx.x]
- If the type of shared is 32*bits, then there won’t be bank conflicts
- But if the data type is smaller, there will be bank conflicts.
- 4*way bank conflicts would occur if the data type is 8*bits, i.e. char. __shared__ char shared[];
- 2*way bank conflicts would occur if the data type is 16*bits, i.e. short __shared__ short shared[];
Structs & Bank Conflicts
Struct assignments compile into as many memory accesses as there are struct members
Consider two struct definitions as:
struct vector { float x, y, z; };
struct myType { float f; int c; };
Now, vectors are created for these struct types as,
__shared__ struct vector vectors[64]; __shared__ struct myType myTypes[64];
In the first case, there won’t be any bank conflicts as struct size is 3 word. There will be 3 accesses per thread, contiguous banks (no common factor with 16) struct vector v = vectors[baseIndex + threadIdx.x];Emphasized
However, the second struct type will cause 2*way bank conflicts for myType. Here, there will be only 2 accesses per thread, which is a factor of 16. struct myType m = myTypes[basedIndex + threadIdx.x]Emphasized
Common Array Bank Conflicts Patterns (1D)
Consider the case when each thread loads 2 elements into shared memory. In this case, a 2-way-interleaved load will result in 2-way bank conflicts:
‘ ’ ‘int tid = threadIdx.x; shared[2*tid] = global[2*tid]; shared[2*tid+1] = global[2*tid+1];′ ’
Such a pattern makes sense for traditional CPU threads, locality in cache line usage and reduced sharing traffic. But it doesn’t suit for shared memory usage where there is no cache line effects but banking effects are there and conflicts must be avoided.
An improved 1 D access pattern
A solution to the problem encountered above can be attained by ensuring that each thread loads one element in every consecutive group of blockDim elements as shown in the image.
‘ ’ ‘shared[tid] = global[tid]; shared[tid + blockDim.x] = global[tid + blockDim.x];′ ’
There is generally no advantage of giving multiple banks to a single thread. However, double precision ones go into 2 banks.
Common Array Bank Conflicts Patterns (2D)
Similar to access patterns explored in 1 D, there can be 2 D access patterns. Operating on 2D arrays becomes essential in applications like Image Processing.
Here, consider the example of operating on 2D array of floats in shared memory, say with 16×16 blocks.
Here, each thread processes a row. So threads in a block access the elements in each column simultaneously (as shown by row 1 in purple in the image above)
Such an obvious access pattern, however, leads to 16-way bank conflicts as all rows start at bank 0.
There can be 2 solutions to this problem.
First is to pad the row, i.e. add one float to the end of each row., as shown in the image (box below). Second can be to transpose the matrix before processing. Though transpose operation itself will suffer from bank conflicts, any later conflicts will be saved.
Load/Store (Memory Read/Write) clustering
Another way of speeding things up is to cluster load/stores together.
That’s important because memory accesses are actually performed in clusters. So, if in the instruction sequence the loads are not together, and they access consecutive memory locations, memory will be accessed multiple times to read the same data. clustering the loads/stores. Compiler can take care of this, but, this requires that there are sufficient non*dependent Load/store accesses.
Thus, a sequence of LD0 → MATH0 → LD1 → MATH1 can be replaced by LD0 → LD1 → MATH0 → MATH1
Memory coalescing
Modern DRAM channels return a large burst of data on each access. It is 64 Bytes in the case of FERMI. Unused data in a burst degrades effective (and scarce) memory throughput.
To explain it further, consider a simple example. When M[0] is accessed, the read is actually performed on 64 bytes for that one access. Now, if another instruction accesses M[1], it will have to be re-read unless the instructions are sequenced in a way that enables using the value which is already read.
L → A → L → A → L → A 0 1 2
So, all 0, 1, ,2 bring the same data; Duplication of efforts.
GPU has hardware to perform memory coalescing. The same accesses will be replaced by
L → L → L → A → A → A
Where L represents a load and A represents a compute operation.
When all 3 loads done together, memory coalescing is achieved. But, one needs to have enough registers for that. This is called load clustering.
Thus, a wasteful duplication of efforts can be avoided by memory coalescing. Memory coalescing hardware minimizes the number of DRAM transactions from a single warp. Other than this explicit optimizing for spatial locality, caches can be used, as is the case with Fermi.
Fermi uses caches for coalescing. Up to 128 bytes are L1 cacheable and 32 bytes are L2 cacheable.
Another important thing is that coalescing works across threads in a warp. When there is no out of order execution, skipping is ruled out.
Also, there is very limited buffering for coalescing. So, when there are many load instructions separated by some other instructions in between, it drastically reduces the probability of it getting coalesced. That is why clustering becomes important to be taken care of by programmer.
0, 1, 2, 3, .. – good
0, 100, 200, … - bad
Reordering of memory operations is difficult. NVDIA attempts coalescing at the level of the core, not at the level of the channel. Channel receives memory requests from the core. Things which go on different cores do not get coalesced in NVDIA design. Some coalescing between warps on a core, but none among the cores. Another thing to keep in mind is that while bank conflicts are for shared memory; coalescing is for global memory.
In a nutshell, If L1 cache is active, there is no need to worry for coalescing. Even without a cache, accesses can be quick if coalescing is there. Caching could still help as it reduces no. of memory accesses. Caches in a way extend the size.
Host Memory, PCI Express, Memory Mapping
GPUs cannot run OS. They cannot run device drivers. They can only act as coprocessor. They invariably need a host to function correctly.
PCI express lets you talk from one DRAM to another directly (through memory controller, without the intervention of CPU/GPU). PCI Express hardware doesn’t support compression, though.
Device driver is responsible for all data movement. CPU and GPU need to agree on what data is there. They’re all memory mapped. It is possible for the GPU kernel to ask for a location in its memory which is mapped to the host memory. But such accesses are very slow.
Device transfers eat up physical memory. They account for a significant amount of time. The no. of pages that can be pinned is limited.
GPU L2 is snooped on PCI transfers. Anytime the GPU does the PCI express transfer, to see if new data is available there. L2 is shared among cores, so it’s coherent with respect to them and is also coherent.
Normal virtual memory allocation involves Page locked pages (virtual to physical lock, that cannot be paged out). There is a need to make the transfer faster.
AMD provides good programmable fusion GPUs (on the same die as CPU), while Intel has graphic only GPUs.
For GPU to access CPU memory, it’s a very slow affair. Currently large barrier b/w CPU n’ GPU is due to device driver even if they’re on the same chip. AMD is moving towards integrating them seamlessly.
Bandwidth of GeForce 9800 GTX
Frequency * 600 MHz with ALUs running at 1.2 GHz
ALU bandwidth (GFLOPs) * (1.2 GHz) X (16 SM) X ((8 SP)X(2 MADD) + (2 SFU)) = ~400 GFLOPs
Register BW * (1.2 GHz) X (16 SM) X (8 SP) X (4 words) = 2.5 TB/s
Shared Memory BW * (600 MHz) X (16 SM) X (16 Banks) X (1 word) = 600 GB/s
Device memory BW * 2 GHz GDDR3 with 256 bit bus: 64 GB/s
Host memory BW * PCI*express: 1.5GB/s or 3GB/s with page locking
Bandwidth of GeForce GTX480/Tesla c2050
GTX 480 – 15 cores, Tesla C050 – 14 cores
Frequency - 700/650 MHz with ALUs running at 1.4/1.3 GHz
ALU bandwidth (GFLOPs)
(1.4 GHz) X (15 SM) X ((32 SP)X(2 MADD)) = ~1.3 TFLOPs (1.3 GHz) X (14 SM) X ((32 SP)X(2 MADD)) = ~1.1 TFLOPs
Register BW
(1.4 GHz) X (15 SM) X ((32 SP)X(4 words)) = ~2.6 TFLOPs (1.3 GHz) X (14 SM) X ((32 SP)X(4 words)) = ~2.2 TB/s
Shared Memory BW
(700 MHz) X (15 SM) X (32 Banks) X (1 word) = 1.3 GB/s (650 MHz) X (14 SM) X (32 Banks) X (1 word) = 1.1 GB/s
Device memory BW
GTX480: 3.6 Gb/s GDDR5 with 6 64*bit channels: 177 GB/s C2050: 3 Gb/s GDDR5 with 6 64*bit channels: 144 GB/s (less with ECC on)
Host memory BW
PCI-express 2.0: Effective bandwidth is 1.5GB/s or 3GB/s with page locking
Communication
Q: How do threads communicate?
GPUs are designed for graphics processing – specifically, data parallel streams representing independent vertices, triangles, fragments, pixels – these are never in communication with each other.
Some communication is allowed in compute mode to allow portability across GPU platforms. There is Shared memory for threads in a thread block. Within warps, no special communication is allowed. Also, there is no communication among thread blocks, expect atomics and derivatives thereof.
Execution kernels can communicate through global device memory. This is a mechanism designed to ensure portability.
Synchronization
Q: Do threads need to synchronize?
Threads in a block share memory, therefore, they need sync. That is also because warps are scheduled out of order, hence their order cannot be relied upon. Hence, a barrier command must be inserted for all threads in a block to ensure all threads reach the same point. A __synchthreads() is implicitly there.
Q: Do blocks need to synchronize?
Blocks should not synchronize, as there is always implicit synchronization at the end of kernel.
Atomic operation are a powerful data*parallel way of manipulating shared data. Using atomics and fences allows one to form one’s own primitives. This, however, is not safe considering GPU uses streaming concurrency. Parallel concurrency may lead to deadlocks if care is not taken to prevent them.
Memory Fences
Fences not same as a barrier. “At this point, I’m guaranteeing that all my memory operations are done”, is what the fence says.
3 types of fences are there.
Thread fence – block : All of thread references to shared memory are complete before the next instruction gets run. All writes done to shared memory by this thread, are done, all memory ops of shared memory are done. It is not exactly a barrier. Everyone is able to see what was written, though.
Thread fence – device : All writes from this thread to shared memory are done. All writes this thread did to global memory are done. All writes by this thread are visible to other threads in the device. This fence Implies an L1 flush.
Thread fence – system: Host can see all writes
Putting a fence doesn’t affect other threads.
There are 2 ways of taking advantage of concurrency.
- Parallel execution
- Streaming execution
Parallel execution means everything is live at the same time; while Streaming means, some, say 4 running at the same time. Streaming is easy to convert to parallel but the converse is not true.
GPU is designed for the streaming model. Streaming between threads within a block is allowed and encouraged. Communication in streaming mode can lead to deadlocks, hence never performed.
It is not safe to assume which CTAs are live at any given point, not even how many CTAs are “live”.
Atomic Operations
With atomic ops one can form synchronization. Atomic operations are what forms the exception to communication among blocks i.e. this is the only form of communication allowed among blocks.
Atomic ops are not cached, therefore, lot of atomics will slow down. For example, with an L2 bank associated with each memory controller, 6 memory controllers so 6 atomic ops can be allowed.
They can be of the following subtypes.
Atomic read*modify*write
- Shared memory
Global memory (Executed in each L2 bank) Simple ALU operations
- Add, subtract, AND, OR, min, max, inc, dec
Exchange operations
- Compare and swap, exchange
