# 平行程式設計
> 心有重核 若計算 若邏輯 繽紛如早春繁花
> 行實多序 若同步 若分離 迅急似盛夏快雨
> [name=劉邦鋒]
> [Tutorial](https://www.youtube.com/playlist?list=PLOvZ8aEg7xDn1H78MT2oApPWrLvcxb3fc)
> [課程網頁](https://sites.google.com/view/pp2019-ntu/assignment)
#### PP Introduction
- Motivation
- 我們希望變得 faster 和 better
- 現代運算的資料量很大
- 我們希望快速地解決問題
- 有些計算是 Time Critical
- Weather forecast
- Stock market
- Radar signal processing
- Moore’s Law：電晶體密度約每兩年便會增加一倍。
- 不但數量要變多，成本也要變低，才有效益。
- 這幾年的走向越來越不穩定，業界認為接下來不會持續太多年。
- 可能可以提高效能，但成本一定會增加。
- Scientific Process
- 科學方法
- Observation
- Theory
- Experiment → Slow
- 有些科學的實驗仰賴電腦的 simulation，需要很大量的計算。
- expensive、dangerous、unfeasible、time consuming
- Definition
- 我們認為 Multicore 和 Multiprocessor 差異不大，都是在 shared memory 上面做操作。
- Multicore CPU
- shared memory, has its own cache → memory arbiter
- EX：Intel Ivy Bridge Intel Xeon
- 15 cores, 37.5 MB L3 cache.
- EX：Intel Xeon Phi
- 三個計畫所構成的電腦
- 60 cores, pentium based.
- Multiprocessor
- shared memory
- EX：TianHe
- 16000 computers(nodes)
- 2 Intel Ivy Bridge Intel Xeon
- 3 Intel Xeon Phi
- Multicomputer
- Communication Cetwork(phone call)
- Cluster Computing：
- loosely or tightly connected computers
- Computer system that are connected by a network, but not a shared memory, can be considered as a cluster.
- Massively parallel computing
- Grid computing
- Cloud computing
- Parallism
- Instruction-level parallelism
- Data parallelism
- Task parallelism
- Dependency Graph
- node → task
- edge → dependency
- may add start node, end node
- Wavefront：不存在彼此有 path 的一個點集合。
- Parallelism：最大的 wavefront 大小。
- Inequality about parallelism
- $N\ge T\ge max(L,\frac{N}{min(P,W)})$
- P = The number of processors
- N = The number of tasks
- L = The longest path in the dependency graph
- W = The maximum number of tasks in a wavefront
- T = The unit of execution time
- Distributed Computing
- 目前還沒有 Automatic Parallel Programming 的機制。
- Parallel Programming v.s. Distributed Computing
- Parallel Programming：
- 目標在減少時間(Performance)。
- Race-car
- Distributed Computing：
- 目標在增加可靠度(reliability)。
- Tank
#### PP Architecture
- Multiprocessor
- Shared global memory, UMA
- 
- Distributed shared memory, NUMA
- 
- **Logically, we do not care about the memory management unit**
- 
- However, in some programming models we need to take care of race condition
- 每個 processor 有自己的計算資源，並可以透過 shared memory 來同步
- Barrier Syncrhonization：等到所有處理器都完成某一工作再一起往後做。
- Race Condition：
- Software Solution：Critical section, lock, synchronization
- 被用來製造 hardware random number generators。
- Cache：
- 
- Hardware Solution：guarantee the consistenct between memory and cache.
- Write-through：每次寫都寫回memory。
- Write-back：別人access memory時才會寫回。
- EX：Intel Gulftown CPU
- 6 cores, Core i7-9xx, 12MB L3 cache
- EX：Nvidia Tesla
- General Purpose GPU
- 240 processors, 4096 MB of GDDR3 memory
- 
- 對應平行程式：OpenCL
- Multicomputer
- 
- Slow or very fast Network。
- 只有 local memory 沒有 shared memory。
- 需要使用一個 master 蒐集大家資訊。
- Topology of Network
- Ring
- Fat tree
- 
- Mesh and Torus
- 
- Hypercube
- FFT (butterfly)
- 
- 目前的趨勢是 Fat tree 或是高維的 Torus。
- LINPACK 作為平行處理的 benchmark
- Flynn’s Taxonomy
- SISD
- Standard von Neumann architecture
- sequential programming
- SIMD
- parallel programming on multiple processors
- data parallelism
- GPU supports
- MIMD
- functional parallelism
- different processors executes different instructions on different data
- MISD
- Fault Tolerance
- The same computation is repeated multiple times
- Google MapRedcue computation
#### PP Evaluation
- Algorithm
- 描述解決問題的過程
- 1.detailed
- 2.well-defined
- 3.temporal dependency
- Upper bound：$\exists A\forall I~~T(A,I)\le f(n)$
- Lower bound：$\forall A\exists I~~T(A,I)\ge g(n)$
- Evaluation
- Speed Up： $k=\frac{T_s}{T_p}$
- 注意要注重比較的基準。
- 不可以讓 sequential program 跑很慢來增加 speed up。
- $1\le k\le p$
- 用一個 processor 來 simulate sequential program。
- 用一個 processor 來 simulate parallel program 的所有 step 數。
- Super-linear Speedup
- 仍然可能發生 $k>p$
- 可能是因為 sub-problems 的 working set 比較少，使 cache 的效果更好
- Efficiency： $e=\frac{k}{p}$
- 單位 processor 的 speed up
- Work： $W_p=e\times T_p$
- Thoery
- Amdahl’s Law：speed up 隨著 p 增加有上限。
- 不可平行化的部分，相較於 input size 呈現比例 x
- $T_p\ge xT_s$
- $k\le \frac{1}{x}$
- Gustafson’s Law：speed up 隨著 p 增加沒有上限。
- 不可平行化的部分，相較於 input size 為 constant x
- 任意的 arbitrarily large data 都可以平行化處理
- $T_s=x+p(1-x)$
- $T_p=1$
- $k=x+p(1-x)\propto p$
- If we have a good speedup even if problem size is small, then we really have a good parallel implementation.
- Otherwise it could be a fabricated phenomenon with the use (or abuse) of Gustafson’s Law.
#### PP Principle
- Even partition
- Even workload distribution
- 希望 makespan 越小越好。
- 常常我們不知道正確的 work load 估計，這時候有一些 scheduling 的 issue。
- Proper granularity(basic unit)
- 當 granularity 小時，越容易 balance workload ，但會有 overhead
- mapping table
- scheduling
- synchronization
- 同時也要看 system 可以 spawn 出的 thread 數量決定。
- Communication reduction
- Synchronization
- 種類
- Barrier Synchronization → shared memory
- Before/After Synchronization → IPC
- Access Synchronization → semaphore
- Trade off with granularity
- Data Locality
- Temporal Locality
- Disk cache
- Translation Lookaside Buffer
- Space Locality
- Matrix Multiplication 如果要求 required data 都要在 local memory，則所有 data 都必須要在同一 processor
- Communication-to-Computation Ratio
- 越低越好，類似 Surface to Volume Ratio
- 方形的切割比片狀的切割要來的好
- Efficient Implementation
- Global synchronization
- Batch mode message passing
- Overlap communication with computation
- Explore memory hierarchy
#### PP Algorithm
- Embarrassingly Parallel
- 不需要溝通
- EX：矩陣乘法、Monte Carlo、Parameter Search。
- Divide and Conquer
- 需要大量溝通、同步。
- Summaion
- 方法一：
- 時間：$O(n/p+p)$
- 極值：$p=\sqrt n\rightarrow O(\sqrt n)$
- 方法二：
- 時間：$O(n/p+logp)$
- 極值：$p=\frac{n}{logn}\rightarrow O(logn)$
- Speed Up：$\frac{n}{n/p+logp}$
- Efficiency：$\frac{n}{n+plogp}$
- Prefix
- Sequential 的 Prefix 問題：時間只要 $O(n)$
- 常常被應用在 compact array 之中。
- [4 0 9 0 0 8 5 7] → [4 9 8 5 7]
- [1 1 2 2 2 3 4 5]
- 只要 operator 有結合律，都可以平行計算加速。
- 方法一：處理器夠多時
- **for $i = 0$ to $logn$**
- 每個 processor 加上自己往左數第 $i$ 的資料
- 時間：$O(logn)$
- 方法二：
- 1.切割資料成 $p$ 份，$p$ 個處理器分別去算 prefix sum
- 2.把每段的最後一個元素取出，額外算出一個 prefix sum
- 3.$p$ 個處理器分別去組合兩個 prefix sum 的結果
- 時間：$O(n/p+logp)$
- Sorting
- $log(n!)=\theta(nlogn)$ 因為 $\frac{n}{2}^\frac{n}{2}\le n!\le n^n$
- Sequential Bucket Sort
- 
- 總共時間複雜度為：$O(n+b)$
- Parallel Bucket Sort
- Read to Bucket
- 1.$O(\frac{n}{p})$ 把 key 拆成 $p$ 份，$p$ 個 processor 分別把 $n/p$ 個 key，記錄到自己獨立的 bucket 裡。
- Read from Bucket
- 2.$O(b)$ 每個 process 整理到一個共用的 bucket。
- 3.$O(b/p+logp)$ 做 $b$ 個元素的 prefix sum。
- 4.$O(b)$ 每個 process 把答案放到 output array。
- 總共時間複雜度為：$O(\frac{n}{p}+b+logp)$
- **需要大量的記憶體。**
- Sequential Quicker Sort(遞迴呼叫)
- 假設找到 pivot 使 input array 平均分到 $g$ 個 group 上。
- 總共時間複雜度為：$T(n)=gT(\frac{n}{g})+O(n)=O(nlogn)$
- Parallel Quicker Sort(不可遞迴呼叫→直接平行解決)
- 假設找到 pivot 使 input array 平均分到 $g$ 個 group 上。
- 1.$O(\frac{n}{p})$ 把 key 拆成 $p$ 份，每個 processor 分別把 $n/p$ 個 key，拆分到共同的 $g$ 個 group 裡。(假設可以同步)
- 2.$O(\frac{n}{p}log\frac{n}{p})$ 把每一個 bucket 裡面的元素用傳統方法 sort。(假設平均分配)
- 3.$O(logp)$ 對每段元素個數去算一個 prefix sum。
- 4.$O(\frac{n}{p})$ 每個 process 把答案放到 output array。
- 總共時間複雜度為：$O(\frac{n}{p}log\frac{n}{p}+logp)$
- Parallel Exchange Sort
- 應用在 Hypercube 上面常見的方法。
- 演算法：
- 將資料平均分到 $p=2^k$ 個處理器上，每個處理器用 bit string 表示。
- **for $i=0$ to $k$**
- 每個處理器找到第 $i$ 個 bit 不同的另一處理器，兩兩配成一對。
- 兩個處理器自行挑選一個 pivot
- 比 pivot 小的數都交給第 $i$ 個 bit $0$ 的那方
- 比 pivot 大的數都交給第 $i$ 個 bit $1$ 的那方
- 最後每個處理器在 sort 自己內部的資料。
- 這個演算法假設
- 每個回合資料都蠻平均的，每個處理器手上交換前後都是 $n/p$ 個。
- 每個回合大概都把手上的資料一半交給對方，再收到另一半的資料。
- 總共時間複雜度為：
- $O(\frac{n}{p})O(logp)+O(\frac{n}{p}log\frac{n}{p})=O(\frac{n}{p}logn)$
- Matrix Multiplication
- 時間：$O(\frac{n^3}{p})$
- Speed Up ： $p$
#### PP Thread
- **Process and Thread**
- Process 是執行中的 Program，Thread 是 Program 的一個執行緒。
- UNIX Process
- 1.Process ID, Process Group ID, User ID, Group ID
- 2.Environment Variable
- 3.Working directory
- 4.Program instructions
- 5.Registers
- 6.File descriptors
- 7.Signal actions
- UNIX Process 的資料擺放
- Stack：函數內宣告
- Heap：動態宣告
- Data：全域宣告
- Thread 和 Process 在平行運算上的差別
- Process 的溝通成本比較大
- Process 的創立成本比較大
- Threads 之間可以共享一些 Process 資源，也有自己的 private 資源
- **P Thread**
- [hello10](https://ideone.com/dFPDkx)
- [eight_single](https://ideone.com/DYPwQC)
- [eight_join](https://ideone.com/89xWUv)
- [eight_mutex](https://ideone.com/f2u3qW)
- Include 以及 Compile
- #include <pthread.h>
- gcc xxx.c -o xxx -lpthread
- 相關的資料型態和 API：
- pthread_t
- pthread_create(&thread,&attr,func,(void*)&t)
- pthread_join(thread,(void**)&num)
- pthread_exit(void* value)
- pthread_attr_t
- pthread_attr_init(&attr)
- pthread_attr_setdetachstate(&attr,PTHREAD_CREATE_JOINABLE)
- pthread_attr_destroy(&attr)
- pthread_mutex_t
- pthread_mutex_init(&lock,mutexattr)
- pthread_mutex_lock(&lock)
- pthread_mutex_unlock(&lock)
- pthread_mutex_destroy(&lock)
#### PP OpenMP
- **Basic**
- Include 以及 Compile
- #include <omp.h>
- gcc xxx.c -o xxx -fopenmp
- 先寫 sequential 的程式，再用 openmp 改成平行版本。
- 相關的使用方法：
- 基本的平行：
- #pragma omp parallel
- omp_get_num_procs()
- omp_get_num_threads()
- omp_set_num_threads()
- omp_get_thread_num()
- 迴圈的平行：
- #pragma omp parallel for
- 初始值：var = exp1
- 條件：var cond exp2
- 增加：
- var++, var--, ++var, --var
- var = var+exp3, var = var-exp3
- var+=exp3, var-=exp3
- 程式會自動等到所有 loop 跑完，做一次 barrier synchronization
- #pragma omp parallel for private(v)
- 
- 每一個 thread 會各有一個相同名稱的 private 變數，並且沒有初始化
- #pragma omp parallel for reduction(+:numSolution)
- Global 先宣告一個變數
- 每一個 thread 會各有一個 private 變數，並且初始化成 global 值
- 平行結束之後，每一個 thread 會把結果加到 global 變數中
- #pragma omp parallel for firstprivate(v)
- 
- 每一個 thread 的 private 變數，被初始化成 global 值
- #pragma omp parallel for lastprivate(v)
- 
- loop 的最後一個 iteration 執行完後，把 private 變數的值存回 global 變數
- #pragma omp critical
- 創造一個 critical section
- **Intermediate**
- Schedule Clause
- #pragma omp parallel for schedule(policy , chunk)
- static：
- round robin，每 chunk 個工作一組，分配到不同 thread 上。
- dynamic：
- 每 chunk 個工作一組，每個 thread 先要一份工作。
- 剩的工作等每個 thread 有空閒時才會再接著分配。
- guided：
- 一開始給比較大的 chunk，後來給比較少的 chunk。
- chunk size 為工作量的下限。
- runtime：
- omp_set_schedule(policy, chunk)
- omp_sched_static
- omp_sched_dynamic
- omp_sched_guided
- omp_sched_auto
- auto：
- Measure Time
- double omp_get_wtime();
- Parallel Sections
- Parallel Directives
- 程式在 #pragma omp parallel 中，會生出 multi thread 了。
```clike=
#pragma omp parallel
{
#pragma omp sections
{
#pragma omp section
#pragma omp section
}
#pragma omp for
for(...)
#pragma omp for
for(...)
}
```
- 避免把每個 thread 用 branch 分開：
```clike=
int main(void){
omp_set_num_threads(8);
#pragma omp parallel
{
int id = omp_get_thread_num();
for(int i=0;i<3;i++){
//printf("%d %d\n",i,id);
if(id<4){
#pragma omp for
for(int j=0;j<4;j++){
printf("1 --- %d %d %d\n",i,j,id);
}
}else{//some iteration doesn't done
#pragma omp for
for(int j=4;j<20;j++){
printf("2 --- %d %d %d\n",i,j,id);
}
}
}
}
return 0;
}
```
```clike=
int main(void){
omp_set_num_threads(8);
#pragma omp parallel
{
int id = omp_get_thread_num();
for(int i=0;i<3;i++){
if(id<4){ //block forever
#pragma omp for
for(int j=0;j<16;j++){
printf("%d %d %d\n",i,j,id);
}
}
}
sleep(id);
printf("%d %d\n",id,omp_get_num_threads());
}
return 0;
}
```
- **Advanced**
- **Private Variable**
- 不同 thread 去存取 global 變數的時間，遠大於去存取 private 變數的時間。
- **Thread Creation/Destroying**
- Thread 的生成和銷毀很耗資源。
- **nowait(上課講義為甚麼會對??)**
- 讓 parallel for 的結束，不特別做 barrier synchronization
```clike=
int main(void){
#pragma omp parallel
{
int id = omp_get_thread_num();
for(int i=0;i<3;i++){
#pragma omp for // nowait
for(int j=0;j<8;j++){
sleep(id);
printf("%d %d %d\n",i,j,id);
}
}
}
return 0;
}
```
```clike=
int main(void){
#pragma omp parallel
{
int id = omp_get_thread_num();
#pragma omp for // nowait
for(int i=0;i<7;i++){
sleep(id);
printf("%d %d\n",i,id);
}
#pragma omp sections
{
#pragma omp section
printf("section 1 %d\n",id);
#pragma omp section
printf("section 2 %d\n",id);
#pragma omp section
printf("section 3 %d\n",id);
}
}
return 0;
}
```
- Critical Region
- 包含越少變數越好，否則會互相等來等去
- 如果需要大量 access 同一個變數：
- 1.設成 private 最後再整理在一起。
- 2.開額外的空間，最後再整理在一起。
- 3.用 reduction 是最好的方法。
- Double Buffer(Game Of Life)
#### PP OpenCL
- **OpenCL：**
- 在 **heterogeneous platforms** 上面執行 的 **framework**
- 底層可以是任何處理器：CPU、GPU、DSP、FPGA 等
- 
- Hardware
- Host：C/C++，單執行緒在CPU上執行，控制流程
- Device：OpenCL平行程式，多執行緒，計算
- compute devices
- compute units
- processing elements(PE)
- Memory
- 每個 PE 有 private memory
- 同一個 compute unit 的每個 PE 共享 local memory
- 不同個 compute unit 的每個 PE 共享 global memory
- **資訊查詢：**
- [Platform ID](https://www.khronos.org/registry/OpenCL/sdk/1.0/docs/man/xhtml/clGetPlatformIDs.html)：
- [Platform Information](https://www.khronos.org/registry/OpenCL/sdk/1.0/docs/man/xhtml/clGetPlatformInfo.html)：
- [Device ID](https://www.khronos.org/registry/OpenCL/sdk/1.0/docs/man/xhtml/clGetDeviceIDs.html)：
- [Device Information](https://www.khronos.org/registry/OpenCL/sdk/1.0/docs/man/xhtml/clGetDeviceInfo.html)：
```clike=
cl_int clGetPlatformIDs();
cl_int clGetPlatformInfo();
cl_int clGetDeviceIDs ();
cl_int clGetDeviceInfo();
```
- **編譯與執行：**
- [Create Context](https://www.khronos.org/registry/OpenCL/sdk/1.0/docs/man/xhtml/clCreateContext.html)：
- [Create Command Queue](https://www.khronos.org/registry/OpenCL/sdk/1.0/docs/man/xhtml/clCreateCommandQueue.html)：
- [Create Program](https://www.khronos.org/registry/OpenCL/sdk/1.0/docs/man/xhtml/clCreateProgramWithSource.html)：
- [Build Program](https://www.khronos.org/registry/OpenCL/sdk/1.0/docs/man/xhtml/clBuildProgram.html)：
- [Create Kernel](https://www.khronos.org/registry/OpenCL/sdk/1.0/docs/man/xhtml/clCreateKernel.html)：
- [Create Buffer](https://www.khronos.org/registry/OpenCL/sdk/1.0/docs/man/xhtml/clCreateBuffer.html)：
- [Set Kernel Argument](https://www.khronos.org/registry/OpenCL/sdk/1.0/docs/man/xhtml/clSetKernelArg.html)：
- [Set Shape And Run](https://www.khronos.org/registry/OpenCL/sdk/1.0/docs/man/xhtml/clEnqueueNDRangeKernel.html)：
- [Read Buffer](https://www.khronos.org/registry/OpenCL/sdk/1.0/docs/man/xhtml/clEnqueueReadBuffer.html)：
```clike=
cl_context clCreateContext();
cl_command_queue clCreateCommandQueue();
cl_program clCreateProgramWithSource();
cl_int clBuildProgram();
cl_kernel clCreateKernel();
cl_mem clCreateBuffer();
cl_int clSetKernelArg();
cl_int clEnqueueNDRangeKernel();
cl_int clEnqueueReadBuffer();
```
- **清空記憶體：**
- clReleaseContext
- clReleaseCommandQueue
- clReleaseProgram
- clReleaseKernel
- clReleaseMemObject
- 
- **Kernel Function 操作性質**：
- **可以印出記憶體位置，來區分該變數到底是 (__global, __const)、__local、__private，這三種的記憶體位置會存放在很不一樣的 virtual memory，而每個 thread 會根據自己的 page table 來對應到正確的 physical memory**
- 參數：
- 指標：可以用 __global 或是 __local
- 變數：可以不特別指定或是 __private
- 變數：
- __local
- __private
- 不可以用 __global
- 同步：
- 只能在同一個 work group 內同步
- barrier(CLK_LOCAL_MEM_FENCE);
- barrier(CLK_GLOBAL_MEM_FENCE);
- **Buffer 的性質：clCreateBuffer 時可以設定的參數**
- 
- 若是 **ALLOC** 或是 **COPY** ：Host 需要使用 clEnqueueWriteBuffer 或 clEnqueueReadBuffer 來操作 Device 的記憶體
- 若是 **USE** ：Device 會自動存回 Host 所指定的記憶體，因此速度比較慢。
- **注意：如果裝置沒有支援 Unified memory for Host and Device，則此設定無效，仍然需要 clEnqueueReadBuffer 把資料讀回來。**
- [**clFinish**](https://www.khronos.org/registry/OpenCL/sdk/1.0/docs/man/xhtml/clFinish.html)：
```clike=
cl_int clFinish(cl_command_queue command_queue);
```
- clEnqueueNDRangeKernel 是一個 nonblocking 的操作
- 使用 clFinish 會等待所有的 command queue 都執行完。
- **Index Domain**：
```clike=
size_t get_global_id(uint dimindx);
size_t get_local_id(uint dimindx);
size_t get_group_id(uint dimindx);
```
```clike=
size_t get_global_size(uint dimindx);
size_t get_local_size(uint dimindx);
size_t get_num_groups(uint dimindx);
```
- 每一個 work item 就是平行的執行 kernel function
- 可以透過以上的 API ，來 specify 不同 work item 要做不同事情
- **矩陣乘法：**
```clike=
__kernel void mul(__global int matrixA[N][N],
__global int matrixB[N][N],
__global int matrixC[N][N]){
int row = get_global_id(0);
int col = get_global_id(1);
int sum = 0;
for(int i = 0; i < N; i++)
sum += matrixA[row][i] * matrixB[i][col];
matrixC[row][col] = sum;
}
```
```clike=
#define N 1024
#define Blk 64
#define BSIDE (N/Blk)
__kernel void mul(__global int A[N][N],
__global int B[N][N],
__global int C[N][N]){
int globalRow = get_global_id(0);
int globalCol = get_global_id(1);
int localRow = get_local_id(0);
int localCol = get_local_id(1);
__local int ALocal[BSIDE][BSIDE];
__local int BLocal[BSIDE][BSIDE];
int sum = 0;
for(int i = 0; i < Blk; i++){
ALocal[localRow][localCol] = A[globalRow][i*BSIDE + localCol];
BLocal[localRow][localCol] = B[i*BSIDE + localRow][globalCol];
barrier(CLK_LOCAL_MEM_FENCE);
for(int k = 0; k < BSIDE; k++)
sum += ALocal[localRow][k] * BLocal[k][localCol];
barrier(CLK_LOCAL_MEM_FENCE);
}
C[globalRow][globalCol] = sum;
}
```
- **Time Measurement**
- [分析 Event](https://www.khronos.org/registry/OpenCL/sdk/1.0/docs/man/xhtml/clGetEventProfilingInfo.html)：
- CL_PROFILING_COMMAND_QUEUED
- CL_PROFILING_COMMAND_SUBMIT
- CL_PROFILING_COMMAND_START
- CL_PROFILING_COMMAND_END
```clike=
clCreateCommandQueue(..., CL_QUEUE_PROFILING_ENABLE, ...);
clEnqueueNDRangeKernel(..., &event);
clWaitForEvents(1 , &event);
cl_int clGetEventProfilingInfo(event, param_name,...);
```
- **Multiple Devices**：
- 如果有兩張 GPU 卡，若是 task 是 independent 的，其實可以同時操作這些 device 來加速
- EX：[20019](https://judgegirl.csie.org/problem/0/20019) / [Main](https://ideone.com/5JMZuu) / [Kernel](https://ideone.com/ywmXXe)
- **Dependency**：
- 方法一：用 clFinish 確保該 command queue 完全執行結束。
```clike=
cl_int clFinish(cl_command_queue command_queue);
```
- 方法二：後來的 kernel 在執行時，設定說要確保等到前面的(兩個) event 執行結束才開始執行。
```clike=
cl_int clEnqueueNDRangeKernel(..., 2, events ,...);
```
- **Group Size**：
- **GPU Streaming multiprocessors 架構：**
- 每一張 GPU 卡片約有 20 個 Streaming multiprocessors (SM)。
- 每一個 SM 有 32 個 processing cores。
- 每一個 core 共用記憶體，執行 SIMD 指令。
- 每次會抓 32 個 thread 形成一個 warp，丟到同一個 SM 上面執行
- 因此最多 640 條 thread 同時運行。
- $GPU$ 同時可以服務的 groups 數量有個上限
- 最多就是 20 個 SM 同時執行
- $GPU$ 的 group size 有個上限
- 規定最多只能有 1024 個 item
- 根據 processing cores，最多只有 32 個 thread 同時執行
- 假設現在工作的數量固定為 $N$，在 group size 不超過 1024 得前提下：
- group size 小，group 數多，平行度不好
- group size 大，group 數少，平行度好
- **Makefile 編譯與執行：**
```clike=
LIBS = /usr/lib/x86_64-linux-gnu/libOpenCL.so -lrt
LINKOPT = -Wl,-rpath,/usr/lib/x86_64-linux-gnu
all:
gcc main.c $(LIBS) $(LINKOPT) -o main-cl -fopenmp
```
#### PP CUDA
- EX：[20029](https://judgegirl.csie.org/problem/0/20029) / [Main](https://ideone.com/AzNqjt)
- **設置記憶體空間：**
```clike=
cudaSetDevice(0);
cudaMalloc(Cu[0], 100 * sizeof(UINT));
cudaSetDevice(1);
cudaMalloc(Cu[1], 200 * sizeof(UINT));
```
- **Host 與 Device 記憶體搬動：**
- [cudaMemcpy](http://horacio9573.no-ip.org/cuda/group__CUDART__MEMORY_g48efa06b81cc031b2aa6fdc2e9930741.html)
- cudaMemcpyHostToHost
- cudaMemcpyHostToDevice
- cudaMemcpyDeviceToHost
- cudaMemcpyDeviceToDevice
```clike=
cudaMemcpy(To, From, Size, Direction);
```
- **呼叫 Kernel Function：**
- 一維：
```clike=
int N = 640; // 共有 640 件工作
int blocks = 20; // 分成 20 個 block
int threads = 32; // 每個 block 包含 32 件工作
fun<<<blocks,threads>>>(N, ...);
```
- 最多三維：
```clike=
int N = 1024*1024; // 共有 1024*1024 件工作
dim3 blocks(16,16); // 分成 16*16 個 block
dim3 threads(64,64); // 每個 block 包含 64*64 件工作
fun<<<blocks,threads>>>(N, ...);
```
- **Kernel Function 寫法：**
```clike=
__global void fun(int n, int A[][MAXN], int B[][MAXN])
{
// 同一個 block 可以 shared 記憶體
__shared__ int Alocal[BSIDE][BSIDE];
__shared__ int Blocal[BSIDE][BSIDE];
// 查詢自己的 block id 以及 thread id
int localRow = threadIdx.x;
int localCol = threadIdx.y;
int globalRow = blockIdx.x * BSIDE + threadIdx.x;
int globalCol = blockIdx.y * BSIDE + threadIdx.y;
for(int iter=0;iter<100;iter++){
//do something
__syncthreads();
//do something
__syncthreads();
}
}
```
- **Makefile 編譯與執行：**[-D_FORCE_INLINES](https://github.com/tsiv/ccminer-cryptonight/issues/13)
```clike=
nvcc -Xcompiler "-O2 -fopenmp" -D_FORCE_INLINES main.cu -o main
```
#### PP MapReduce / Hadoop
- [Google File System](https://zh.wikipedia.org/wiki/Google%E6%AA%94%E6%A1%88%E7%B3%BB%E7%B5%B1)
- Google 自行研發的一種分散式系統
- 假設：
- 系統是由大量便宜、不穩定的硬體組成
- 檔案數量不多，但每個檔案都很大
- 每個 File 被 Divide 成 64 MB blocks
- 每個 File 有三份 Replica
- Read 動作大多都是 sequential 的
- Write 動作大多都是 sequential 的接在檔案最後
- 重要的是高的頻寬，而非開始的延遲
- MapReduce
- 概念：
- Programmer 只需要撰寫簡單的 Map 以及 Reduce 的函數，專注在計算本身而非平行化
- 系統會自動在 large cluster 上平行的執行
- 
- Map Function
- 把 input 變成 (key, value)
- 運行 Map function 的機器稱為 Mapper
- 運行完會把結果寫入 local disk
- Reduce Function
- 把 (key, value) 變成 output
- 運行 Reduce function 的機器稱為 Reducer
- 運行前要使用 remote procedure call 去向 Mapper 附近的 Disk 拿資料
- Shuffle Stage
- Runtime System 自動對中間的 (key, value) 做整理
- 把 (key, value) 從 map function 透過網路傳輸到 reduce function
- 此外在送給 reduce function 前，會先 sorting 以及 grouping
- Partition
- 預設是使用 Hash Function，決定每個 (key, value) 要送給哪個 Reducer
- 使用者也可以自行設定對應關係
- Combiner
- 在 Mapper 寫入 local disk 時，預先跑 combiner 的整理程式
- 目的是降低網路傳輸時間，也減輕 reducer 的工作負擔
- Fault Tolerance
- Master 會一段時間固定去 ping worker
- Mapper 發生 Fail，他原已完成的工作要全部重做
- Reducer 發生 Fail，他原已工作已寫入 GFS，不用重做
- Master 發生 Fail，系統需要重新啟動
- [Hadoop](https://hadoop.apache.org/)
- 基於 GFS 以及 MapReduce，由 Apache 開發，使用 java 所撰寫的開源軟體架構
Sign in with Wallet
Connect another wallet