Intro

shared memory algorithm design

For non-shared memory algorithms, we have to utilize bararrier to gain data transformed.

for shared memory database algorithm, we have to decide the group part of the task that shared memory a lot . Then just apply the inserting directives of omp or mpi

Design considerations

two main considerations lies in data dependence and load balance so we have to apply the following steps

• data dependence analysis
• static r dynamic and block r cyclic work assignment
• variable specification whether using shared private r reduction and row-wise r column-wise
  • shared variables cause cache coherence traffic and much lower performance
  • private and reduction variables don't need synchronization
  • dimension mapping is more relying on the cache locality

Main Consideration for data dependence analysis

RAW & WAR & WAW

all the collations should be avoided though they might run into right situations

Goal is to we should run all the dependent situation on the same processor.

loop dependence analysis

• loop-carried dependence
  • dependence exists across different iterations of loop
• loop-independent dependence
  • dependence exists within the same iteration of loop

example

iteration-space traversal graph (ITG)

• iteration-space traversal graph is a line graph showing the order of traversal in the iteration space.

• 

loop-carried dependence graph (LDG)

• Given the ITG, can determine the dependence between different loops.
• Loop-carried Dependence Graph (LDG) shows the loopcarried true/anti/output dependence relationships.
• Node in LDG is a point in the iteration space.
• Directed edge in LDG is the dependence.
• LDG helps identify parts of the loop that can be done in parallel.

examples

 

Distance and direction vector

Loop

Algorithm analysis

Intro

• Sparse matrix vector multiplication.
• Many scientific algorithms require multiplying a matrix by a vector.
  • Optimization (e.g. conjugate gradient low degree mesh.), iterative methods (solving linear systems), eigenvalue methods (e.g. graph partitioning vx), simulations (e.g. finite elements), data analysis (e.g. Pagerank the web connectivity matrix).
• The matrices are often sparse.
  • In an nxn matrix, there are \(O(n^2)\) nonzero elements.

DIA Format

ELL format

COO format

CSR format

Hybird format

what to use

ELL kernel

CSR scalar vector kernel

CSR vector kernel

COO vector kernel

• In parallel system, high performance achieved by using algorithm with high parallelism, good load balancing, memory locality, etc.
• But must also ensure multiple concurrent threads / processes operate correctly.
• Concurrency bugs can arise due to unexpected interleaving of concurrent threads.
• One of the most difficult issues to deal with in parallel / distributed computing.
  • Bugs occur at random times depending on the interleaving.
  • Bugs don’t occur during testing, but they will eventually occur in system deployed system.
  • Humans have hard time anticipating and resolving concurrency bugs.Concurrency bugs can have very serious consequences.

    Therac-25 radiation therapy system had a concurrency bug that led to radiation overdose and death of several patients.

    Space shuttle aborted 20 minutes before maiden launch due to concurrency bug in its avionics software.

eliminating concurrency bugs

• Multiple ways, each with pros and cons.
• Critical sections and locks
  • Prevent processes from accessing a block of code at the same time.
  • Easy to use, effective for some problems.
  • But cause contention, overhead and serialization.
  • Need to decide how much code to lock.
  • Too little, and may still get concurrency bug.
  • Too much, and we lose parallelism and performance.
• If processes acquire several locks, they need to coordinate to maintain correctness, avoid deadlock.
• Low priority that acquires a lock can delay high priority thread (priority inversion).
• Despite these problems, locks are still the most widely used solution.
• transactional memory
  • A block of code is defined as a transaction, i.e. the block of code either executes atomically or doesn’t execute at all.
  • Keep track of reads and writes done by a transaction. If two concurrent transactions read and write to same memory location, abort one of them, i.e. undo all the changes it made.
  • Two concurrent transactions accessing different memory locations can both commit, i.e. all the changes it made are made permanent.
  • Transactional memory can either be implemented in hardware (HTM) or software (STM).
  • HTM has limits of size and type of transactions it can handle. Implemented in e.g. Intel Haswell, IBM Power8.
  • STM is more flexible, but can be very slow.
• Write your own concurrent code, without hardware support.
  • Challenging for most programmers. Not scalable in terms of productivity.
  • Correct, efficient algorithms are often research level publications.

Mutual exclution

• Given n concurrent processes that want to perform a critical section (CS), mutual exclusion can satisfy the following properties.
  • No two processes are in CS at same time.
  • If several processes want to enter the CS, at least one succeeds in finite time (deadlock freedom).
  • If several processes want to enter the CS, every process succeeds in finite time (wait freedom).
• All (useful) mutex algorithms satisfy first and second properties.
• Some algorithms satisfy the third property, but have lower performance.

Mutual exclusion algorithms

• Mutex is provided by locks. But how are locks implemented?
  • Depends on the type of operations the underlying hardware supports.
  • First type of algorithm uses only read / write operations.
  • Second type uses hardware synchronization primitives such as test-and- set (TAS) or compare-and-swap (CAS), provided in most processors.
• TAS(x) tests if a Boolean variable x is true.
  • If x == false, it sets x to true.
  • Returns x’s value before the TS.
  • All these steps done atomically, without interruption from other threads.
  • getAndSet(x) is like TAS(x), but allows non-Boolean x.
• CAS(x,v,v’) tests if variable x currently equals v. If so, it sets x to v’. Otherwise, it doesn’t change x. It also returns x’s current value.
  • Again, all this is atomic.
  • Algorithms also depend on a processor’s memory model.
  • Some processors reorder instructions to avoid stalls and obtain higher performance. This can break many lock algorithms.
  • Most lock algorithms assume memory model is sequentially consistent, i.e. the execution order of instructions from different processes is an interleaving of the instructions of each process in program order.

first attempt

1. 
2. 

shared memory model compared with mpi

• Threads (e.g. Pthreads, Java)
  • individual sequences of instructions that can execute in parallel and access shared data
  • very general, because programmer have to manage everything.
• parallel programming language/ library
  • a parallel language or library is used to create code
  • require new compiler
• Compiler directives (OMP)
  • inserts compiler directives into a sequential program to specify parallelism and indicate shared data and the compiler translates into threads.
  • still uses threads underneath, but system manages the threads.
  • easy to program. but lose flexibility.

Processes and thread

• Process(MPI)
  • separate program with its own variable, memory, stack and instruction pointer.
  • Different programs can't access each other's memory

  

• thread(OMP)
  • Concurrent routine that shares the variables and memory space, but has its own stack and instruction pointer.

  

Reference: Cuda dev blog

why we introduce intrinsics

• Race condition: Different outcomes depending on execution order.

• Race conditions can occur in any concurrent system, including GPUs.

  • Code should print a =1,000,000.
  • But actually printed a = 88.

• GPUs have atomic intrinsics for simple atomic operations.

• Hard to implement general mutex in CUDA.

  • Codes using critical sections don’t perform well on GPUs anyway.

The atomic thread in critical section is not good in performance. The example is identified for variable d_a

cuda atomics

Can perform atomics on global or shared memory variables.

• int atomicInc(int *addr)
  • Reads value at addr, increments it, returns old value.
  • Hardware ensures all 3 instructions happen without interruption from any other thread.
• int atomicAdd(int *addr, int val)
  • Reads value at addr, adds val to it, returns old value.
• int atomicMax(int *addr, int val)
  • Reads value at addr, sets it to max of current value and val, returns old value.
• int atomicExch(int *addr1, int val)
  • Sets val at addr to val, returns old value at val.
• int atomicCAS(int *addr, old, new)
  • “Compare and swap”, a conditional atomic.
  • Reads value at addr. If value equals old, sets value to new. Else does nothing.
  • Indicates whether state changed, i.e. if your view is up to date.
  • Universal operation, i.e. can be used to perform any other kind of synchronization

examples

Finding max of array

improve it

we can improve it bu split the single global max into num_locals number of local max value.

Thread i atomically maxes with its local max. can max the local_max[locali],

a better solution is to make it into the tree DS+CA

Compute the histogram

! whether the data is stored on shared or global memory is depend on programmers

Shuffle intrinsics

! this intrinsic is not limited to cuda but all SIMD architecture

from avx256 vector elements we have

__m256 load_rotr(float *src)
{
#ifdef __AVX2__
    __m256 orig = _mm256_loadu_ps(src);
    __m256 rotated_right = _mm256_permutevar8x32_ps(orig, _mm256_set_epi32(0,7,6,5,4,3,2,1));
    return rotated_right;
#else
    __m256 shifted = _mm256_loadu_ps(src + 1);
    __m256 bcast = _mm256_set1_ps(*src);
    return _mm256_blend_ps(shifted, bcast, 0b10000000);
#endif
}

​

For Kepler architecture, we have 4 intrisics.

The goal of the shuffle intrinsics is actually for optimizing Memory-Fetch Model

butterfly operations

Assume we have at most shuffle instead of shared memory

Collective communication

• Important as the basis for many parallel algorithms.
• Implemented using multiple point to point communications, possibly in parallel.
  • Assume communicating m words (without contention) between a source and destination takes \(t_{s} + m t_{w}\) time.
  • Ignore the the distance of the message, since per hop latency th is usually small.
• With contention c, time becomes \(t_{s} + c m t_{w}\).
• Implementations depend on communication hardware architecture.
• Operations
• Broadcast and reduction.
• All-to-all broadcast and reduction.
• All-reduce, prefix sum.
• Scatter and gather.
• All-to-all scatter and reduce.
• Circular shift.

Broadcast and reduction on ring

just cut them into halves.

cost

• With p processors, log p steps.
• In each step, a processor sends a size m message.
• Total time$ (t_{s} + m t_{w}) log p$ take \((t_{s} + m t_{w})\) as const

Broadcast on mesh

algorithm

• Root first broadcasts along its row using ring algorithm.
• The nodes of root’s row then broadcast along their columns using ring algorithm.

cost

• \(log \sqrt{p} = (log p) / 2\) steps along row, same along columns.
• Total time (ts + m tw) log p.log
• p = (log p) / 2 steps along row, same along columns.
• Total time (ts + m tw) log p.

Broadcast on hypercube

algorithm

RING CAN BE PROJECTED TO THE HYPERCUBE but with more congestion.

prefix sum on hypercube