- Important facts in
parallel programming:
- Although
all the threads
execute the same program:
- The progress in the program execution can be different for different threads !!!
Example:
- The main( ) function
can launch
more than 1024 threads
- In some GPU cards,
there are only 1024 processors
available
- Only a max of 1024 threads can execute simulataneously !!!
So execution progress for different threads is different because there are not enough processors to execute all the threads !!!
- The parallel executing threads
must some times
co-operate to
complete a task
correctly
Result:
- Sometimes, threads must wait for other threads to reach a certain point of execution progress before the threads can continue !!!
We call this "waiting for other threads to reach some point in the execution": synchronization
- Although
all the threads
execute the same program:
- Consider the following
problem:
- Input:
Array: A: A[0] A[1] A[2] A[3] A[4] A[5] A[N-1]
- Reduce the
input array so that:
A[0] = A[0] A[1] = A[0] + A[1] A[2] = A[0] + A[1] + A[2] A[3] = A[0] + A[1] + A[2] + A[3] ... A[N-1] = A[0] + A[1] + A[2] + .... + A[N-1]
Example:
Input: A = 1 2 3 4 Output: A = 1 3 6 10 - Input:
- CPU algorithm:
/* ------------------------------------------------ Reduce: input: a[0] a[1] a[2] ... output: a[0] = a[0] a[1] = a[0]+a[1] a[2] = a[0]+a[1]+a[2] ... ------------------------------------------------ */ void reduce(int a[], int n) { int i; for (i = 1; i < n; i++) a[i] = a[i] + a[i-1]; }
- Comment:
- The running time is O(n) (n = array size)
- Example Program:
(Demo above code)
Program file: /home/cs355001/demo/CUDA/6-reduction/cpu-array-reduce.c cs355@ghost01 (2655)> cpu-array-reduce 10 1 1 1 1 1 1 1 1 1 1 1 2 3 4 5 6 7 8 9 10
- I will use an array of 8 elements
as example to
illustrate the
parallel algorithm...
Suppose we have an array of 8 elements:
We assign 1 thread (threadID i) to compute the final value of a[i]:
- We compute the
partial array sum
in a number of rounds as follows:
- In round 1,
each thread
with threadID ≥ 1
(= 20)
computes:
a[ threadID ] = a[ threadID ] + a[ threadID - 1 ] ;Result:
- In round 2,
each thread
with threadID ≥ 2
(= 21)
computes:
a[ threadID ] = a[ threadID ] + a[ threadID - 2 ] ;Result:
- Finally, In round 3,
each thread
with threadID ≥ 4
(= 22)
computes:
a[ threadID ] = a[ threadID ] + a[ threadID - 4 ] ;Result:
- In round 1,
each thread
with threadID ≥ 1
(= 20)
computes:
- Here is a naively coded
parallel algorithm that reduce an
array of
512 elements:
__global__ void reduce(int *a, int n) { int myID = blockDim.x*blockIdx.x + threadIdx.x; if ( myID >= 1 && myID < n) a[myID] += a[myID-1]; if ( myID >= 2 && myID < n) a[myID] += a[myID-2]; if ( myID >= 4 && myID < n) a[myID] += a[myID-4]; if ( myID >= 8 && myID < n) a[myID] += a[myID-8]; if ( myID >= 16 && myID < n) a[myID] += a[myID-16]; if ( myID >= 32 && myID < n) a[myID] += a[myID-32]; if ( myID >= 64 && myID < n) a[myID] += a[myID-64]; if ( myID >= 128 && myID < n) a[myID] += a[myID-128]; if ( myID >= 256 && myID < n) a[myID] += a[myID-256]; } int main (int argc, char *argv[]) { N = 512; int K = 512; int NBlks = ceil((float) N / K ); reduce<<
- Example Program:
(Demo above code)
When you run it, you will have problems (the problem does not show up for array of < 512 elements -- the kernel must run long enough so they diverge in speed)
Program: /home/cs355001/demo/CUDA/6-reduction/array-reduce1.cu (The input of array A = 1 1 1 1 1 1 ... The correct output is: 1 2 3 4 5 6 7 .... 512 Sometimes you will get the correct output. But other times: 1 2 3 ... 474 477 478 479 480
- What went
wrong:
- Consider the
first round of
addition:
- Suppose
that thread 3 was
slower and
did not finish
when the next round starts:
- This delay will
cause a[5]
to compute an
incorrect result
in round 2:
- Consider the
first round of
addition:
-
Conclusion:
- Every thread must finish the current round before the next round can begin !!!!
- The CUDA function:
__syncthreads( );causes a thread to pause (wait) until all threads in the same thread block have executed the __syncthreads( ) method
- Note:
- Recall that only threads in the same thread block can synchronize with each other through shared memory in the (same) multi-processor
- Corrected program:
__global__ void reduce(int *a, int n) { int myID = blockDim.x*blockIdx.x + threadIdx.x; if ( myID >= 1 && myID < n) a[myID] += a[myID-1]; __syncthreads( ); if ( myID >= 2 && myID < n) a[myID] += a[myID-2]; __syncthreads( ); if ( myID >= 4 && myID < n) a[myID] += a[myID-4]; __syncthreads( ); if ( myID >= 8 && myID < n) a[myID] += a[myID-8]; __syncthreads( ); .... }
- Example Program:
(Demo above code)
Program: /home/cs355001/demo/CUDA/6-reduction/array-reduce2.cu (The input of array A = 1 1 1 1 1 1 ... Now you will always get the correct output: 1 2 3 4 5 6 7 .... 512)