CMU15-418notes(1-9)

视频：https://www.bilibili.com/video/BV1Rh4y1F7aU/?spm_id_from=333.788&vd_source=463e5b3e4b18e5453477b57388c2e427

课程主页：http://15418.courses.cs.cmu.edu/spring2016/lectures

Lecture 1: Why Parallelism

功率墙导致的摩尔定律的失效让性能的提升方向转向并行。
并行计算是同时使用多种计算资源解决计算问题的过程。
快不等同于高效。
使用加速比衡量性能提升。

Lecture 2: A Modern Multi-Core Processor

现代处理器实现并行的方式：

多核：使用多核来实现线程并行，每个核执行不同的指令流。
SIMD：在一个核的指令流中使用多个ALU。向量化可以通过编译器显式使用SIMD，也可以由硬件运行时进行。
超标量：在单指令流提高指令级并行度，并行处理多条指令。

关于访存：

提高带宽比降低延迟简单。
为了利用访存延迟的时间，可以在发生访存的Stall时切换执行另一个线程的指令(本线程的指令可能存在依赖，不能充分利用超标量下的ILP)。这需要在一个周期发射多个不同线程的指令，因此称为同时多线程，即超线程(SMT)。但是多线程会导致每个线程的存储空间减少，因此需要更高的带宽。
CPU使用大缓存，较少的线程，适中的带宽，依靠缓存和预取减少访存时间。而GPU使用小的缓存，较多线程，大带宽，通过多线程提高运算效率。所以GPU有自己的存储。

Lecture 3: Parallel Programming Models

ISPC

ISPC是intel开发的SPMD(single program multiple data)编译器。程序中的并行程序实例，实际上是转化为了SIMD向量指令。

ISPC程序中，数据可以选择交错分配给program instance，也可以选择按块分配，但由于ISPC实际上是使用SIMD指令，因此交错的方式更快。

三种通信模型

共享内存
- 需要使用同步和互斥原语
- 需要硬件支持来提高访存效率(NUMA)
消息传递模型：通过发送和接收消息传递数据
数据并行模型：有严格的要求(例如SPMD)，通常需要使用流式编程模型
- 数据流是数据的集合，数据可以独立处理，没有依赖
- 处理数据不会产生副作用

Lecture 4: Parallel Programming Basics

并行程序运行的基本过程：

分解：是程序员的工作，将问题分解为多个任务。
分配：可以由程序员静态完成(例如pthread指定线程分配任务)，也可以运行时完成(ISPC)。
编排：安排同步，通信等。
映射：由OS/硬件/编译器完成，将线程或程序实例安排给具体的硬件。

Amdahl定律：

本节的后半段是海洋模拟程序的并行计算实例：

共享内存模型，可以用多线程实现，需要锁来保护共享数据：

数据并行：

消息传递模型的例子在后续节中。

Lecture 5: Performance Optimization I: Work Distribution and Scheduling

性能优化-平衡负载

优化目标：

平衡负载
减少通信
减少额外的工作

本节主要探讨的内容是负载的平衡，从任务分配开始考虑：

静态分配：优点是没有运行时开销
动态分配：

如果task划分的太小，就会在切换task时浪费太多时间。而task划分的很大又会导致负载不平衡。因此task size要合理分配，并且size大的要先分配，保证其他线程都能同时工作，而不是放在最后，导致其他线程空闲。

上图中只用了一个队列来分配任务，由于需要使用锁，开销可能很大，改进方法是使用多个队列，这称为分布式工作队列。当队列中的任务全部完成后，线程需要从其他队列中获取任务继续工作，直到全部任务完成。这样做的另一个好处是每个队列中的任务可以是有依赖的。

fork-join程序的调度

本节的后半讨论了分治问题的并行计算。通常的并行计算是不存在数据依赖的：

数据并行的方式：

多线程(共享内存)：

但是有些问题是串行和并行交叉的，例如快排：

C++的扩展标准Cilk plus可以处理这样的并行，主要通过两个声明：

cilk_spawn foo(args); 	//逻辑上可以并行
cilk_sync;

//	foo()	and	bar()	may	run	in	parallel	
cilk_spawn	foo();	
bar();	
cilk_sync;

cilk_spawn并不是创建线程，而是提前创建线程，再给线程任务。

假设cilk_spawn的部分称为child，接下来是cilk实现的一些底层细节。

线程执行任务仍然使用work queue，以便其他线程可以去队列找到并执行可并行的待执行任务：

线程可以先执行child，也可以先执行后续部分，这两种是不一样的。假设有一个循环：

for	(int i=0; i<N; i++){	
	cilk_spawn foo(i);	
}
cilk_sync;

先执行child，那么循环会搁置在队列中；先执行后续部分，那么会先进行循环，把所有foo(i)放到队列中。

线程总是从获取最早入队的任务，对于分治问题，早入队的任务更大，相当于选取了树的靠近根节点的分支处理，这样就能尽量平衡负载。(对于快排，相当于先执行排序，将分治的另一半入队，并递归下去直到最小规模)。

对于sync，有两种策略，clik采用的是第二种策略：

stalling：调用spawn产生分支的线程继续执行sync后面的内容，需要等待其他线程完成。
greedy：最后一个完成任务的线程执行sync及之后的内容。

Lecture 6: Performance Optimization Part II: Locality, Communication, and Contention

本节课的内容比较散，直接来到summary。

Inherent vs. artifactual communication
- inherent communication：并行算法中分解和分配任务必要的通信
- artifactual communication：所有其他的通信，依赖于系统实现细节的通信，例如缓存
对涉及局部性、通信和冲突的内容进行优化：
- 通过利用局部性来提高性能
- 减少通信(通过合并消息，发送大体积的消息)
- 通过复制和交错访问资源，以及使用细粒度的锁，来减少冲突
- 提高重叠的部分，例如流水线，多线程隐藏延迟等

Lecture 7: GPU Architecture and CUDA Programming

这一节听的很模糊，最后是选看了中山大学和清华大学的并行计算课以及一些文章做的笔记。

中山大学-超级计算机原理与操作：https://www.easyhpc.net/course/27?courseTab=lessonList&activeLesson=345

清华大学-并行计算基础：https://www.easyhpc.net/course/10?courseTab=lessonList&activeLesson=110

一些其他参考链接：

https://hustcat.github.io/gpu-architecture/

http://haifux.org/lectures/267/Introduction-to-GPUs.pdf

https://zhuanlan.zhihu.com/p/123170285

GPU

GPU和CPU一样，也是一种处理器，起初其功能是这样一条流水线：

GPU与CPU的不同：

轻量级线程，可快速切换
cache容量小
较多电路用于计算，轻量级线程以SIMD方式运用
通过吞吐率提高效率，而不是单个线程的执行速度

起初的GPU编程方法是通过DirectX和OpenGL等图形API进行程序映射，后来Nvidia发布了CUDA，提供了对GPU的底层控制，GPU的可编程性逐步提高，general-purposeGPU逐渐流行起来，也成为了HPC的潮流。CUDA只能运行在Nvidia的GPU上，而OpenCL是CUDA的一个开源标准版本，可以运行在各类CPU，GPU，DSP，FPGA上。

GPU 体系结构

概念性结构

一个假想的GPU中的单核结构如下图所示：

这个核中有18个Ctx(context)组，每个组中有8个Ctx，这样这个核可以并发处理(交错)18个指令流，这样组中就有18*8个并发程序片。分组的目的是为了减少访存的延迟(超线程的思想)，分组的数量也可以改变。每个组有多个线程，每个线程处理的元素是一个fragment。(一个组执行的都是相同的指令)

现在假设使用有4个Ctx组的核，一个GPU的例子如下：

NVIDIA GeForce GTX 580

下面以GTX580为例，理解GPU的具体结构，首先对应上面的概念，一个GPU是由多个核构成的，一个GTX580包含16个核，和上面的概念性GPU是一样的。

下面是每个单核的结构，在上面的概念性阐释中，可以看到一个核是由一个decoder+多个ALU+Ctx组成的，而在GTX580中，一个核由两个decoder+多个复杂ALU(CUDA core，(也叫做Stream Processor, SP))+Ctx和共享内存构成，称为Stream Multiprocessor(SM)。

由于有两个decoder，这一个核可以并行执行两个指令。每个decoder对应16个CUDA core。所有CUDA core执行的线程被称为线程束(warp)，即SIMT(单指令流多线程)。

综上，一个GPU有多个SM，每个SM可以调度线程束warp，可以同时执行两个指令，每个warp都是执行相同指令的一组线程，这些线程通过CUDA core进行运算。此外，为了调度和运行，SM中还有warp调度器等部件。

目前CUDA的warp的大小为32，即一个线程束里面有32个线程。而每个SM中的warp数量是取决于具体的GPU的。

对于上图，标出的似乎是同时执行两个warp，但是这样每个warp就只有16个线程，单条指令。实际情况应该是和硬件架构有关的，可以同时执行多个warp，有的GPU似乎还可以在一个warp同时发射多条指令，总之，GPU core被分组为warp是由架构决定的，不是固定标准的

CUDA编程模型

CUDA程序分为CPU的HOST部分和GPU的DEVICE部分，两者有独立的存储器。GPU运行的函数被称为核函数，通过__global__声明，例如：

// Kernel definition
__global__ void VecAdd(float* A, float* B, float* C){
    int i = threadIdx.x;
    C[i] = A[i] + B[i];
}
int main(){
...
    // Kernel invocation with N threads
    VecAdd<<<1, N>>>(A, B, C);
... 
}

线程组织

CUDA中的线程组织为三个层次。最高的层次为grid，一个grid包含多个block，block可以是一维，二维或三维组织的，block中的thread也可以是一维，二维或三维组织的。

由于组织形式的不同，threadID的计算方法也不同，例如：

grid一维，block一维：threadID = blockDim.x(block的线程数)*blockIdx.x(grid中的哪个block)+threadIdx.x(block中的位置)

grid一维，block二维：threadID = blockDim.x*blockDim.y*blockIdx.x+threadIdx.y*blockDim.x+threadIdx.x

线程映射

每个核函数会产生一个grid，并有相应的block和thread组织方式，每个block最终会映射到一个SM上，而thread映射到GPU core上。同一个block中的thread是可同步的。SM上的warp调度器会对block进行调度安排。由于warp的大小一般为32，所以block所含的thread的大小一般要设置为32的倍数。