cs179notes_1

2017-11-23

useful links
lecture1
1. CPU(Central Processing Unit)
2. GPU(Graphics Processing Unit)
  1. 为什么要用GPU?(GPU的优势)
lecture2 More Basics
1. 线程组织形式(Thread Organization)
  1. 关键词
  2. GPU内部构造
lecture3
1. 并行问题
2. 非并行问题
lecture4 GPU Memory Systems
lecture5 GPU Compute Programming
lecture6 同步、共享内存、矩阵转置
1. 同步

useful links

cs179
cuda c programming guide

lecture1

CPU(Central Processing Unit)

传统上，应用程序使用CPU作为主服务器计算，它有如下特点：
•通用功能
•建立技术
•通常配备8个或更少的强大的核心
•对于并发进程是最佳的，但不是大规模的并行计算

GPU(Graphics Processing Unit)

为并行化问题设计的相对较新的技术。
•最初专为图形创建
•成为更有能力的一般计算

为什么要用GPU?(GPU的优势)

举例，对于两个数组相加的操作。在CPU上这样实现：

串行实现

float *C=malloc(N*sizeof(float));
for(int i=0;i<N;++i){
    C[i]=A[i]+B[i];
}
return C;

更好一点可以使用并行实现。

(allocate memory for C)
Create # of threads equal to number of cores on processor
(around 2, 4, perhaps 8)
(Indicate portions of A, B, C to each thread...)
...
In each thread,
For (i from beginning region of thread)
C[i] <- A[i] + B[i]
//lots of waiting involved for memory reads, writes, ...
Wait for threads to synchronize...
This is slightly faster – 2-8x (slightly more with other tricks)

但是这带来一些问题：能用多少个线程？如何扩展程序性能？实际上，在CPU上允许的最大线程数量比较少。

GPU与CPU在上下文切换上的差异

CPU与GPU架构的一个主要区别就是CPU与GPU映射寄存器的方式。CPU通过使用寄存器重命名和栈来执行多线程。为了运行一个新任务，CPU需要进行上下文切换，当前所有寄存器的状态保存到栈上，然后从栈中恢复当前需要执行的新线程上次的执行状态。而这些操作通常要花费上百个CPU时钟周期，所以如果在CPU中开启过多的线程，时间几乎都将花在上下文切换过程中寄存器内容的换进换出上。
然而，GPU却恰恰相反，GPU利用多线程隐藏了内存获取与指令执行带来的延迟，因此，在GPU上开启过少的线程反而会因为等待内存事务而使得GPU处于闲置状态。另外，GPU也不实用寄存器重命名的机制，而是致力于为每一个线程都分配真实的寄存器，因此，当需要上下文切换时，所需要的操作就是将指向当前寄存器组的选择器(或者指针)更新，以指向下一个执行的线程束的寄存器组，因此几乎就是零开销。

在GPU上可以这样实现

(allocate memory for A, B, C on GPU)
Create the “kernel” – each thread will perform one (or a few)
additions
Specify the following kernel operation:
For all i‘s (indices) assigned to this thread:
C[i] <- A[i] + B[i]
Start ~20000 (!) threads
Wait for threads to synchronize...

GPU的优势在于：强调并行性意味着我们有很多核心。这允许我们同时运行多个线程而没有上下文切换。

GPU的计算步骤

• Setup inputs on the host (CPU-accessible memory)
• Allocate memory for outputs on the host
• Allocate memory for inputs on the GPU
• Allocate memory for outputs on the GPU
• Copy inputs from host to GPU
• Start GPU kernel
• Copy output from GPU to host

值得注意的是：数据的复制过程可以是异步的。

编程举例

内核代码

__global__ void cudaAddVectorsKernel(float* a,float* b,float* c){
    unsigned int index=blockIdx.x*blockDim.x+threadIdx.x;
    c[index]=a[index]+b[index];
}

CPU上调用内核的代码

void cudaAddVectors(const float* a,const float* b,float* c,size){
    float* dev_a;
    float* dev_b;
    float* dev_c;
    cudaMalloc((void**) &dev_a,size*sizeof(float));
    cudaMemcpy(dev_a,a,size*sizeof(float),cudaMemcpyHostToDevice);
    cudaMalloc((void**) &dev_b,size*sizeof(float));
    cudaMemcpy(dev_b,b,size*sizeof(float),cudaMemcpyHostToDevice);
    cudaMalloc((void**) &dev_c,size*sizeof(float));
    const unsigned int threadsPerBlock=512;
    const unsigned int blocks=ceil(size/float(threadsPerBlock));
    cudaAddVectorsKernel<<<blocks,threadsPerBlock>>>(dev_a,dev_b,dev_c);
    cudaMemcpy(c,dev_c,size*sizeof(float),cudaMemcpyDeviceToHost);
    cudaFree(dev_a);
    cudaFree(dev_b);
    cudaFree(dev_c);
}

GPU的历史

lecture2 More Basics

线程组织形式(Thread Organization)

关键词

线程(Thread)
块(Block)
网格(Grid)
流形多处理器(Streaming Multiprocessor)
Warp
Warp Divergence

GPU内部构造

把Device Memory(也叫Global Memory)当做GPU里的RAM。它比实际的RAM快速，而且还可以更快。
GPU有着许多SMs,其中每一个SM有多个处理器，但是只有一个指令单元(instruction unit)。SM中的一个处理器组在一定时间内必须运行完全相同的指令。
当内核被调用时，一个任务被分为多个线程，每个线程控制整个任务的某个部分。这些线程被分成一个个网格块(a Grid of Blocks)。Grids和Blocks都有三个维度。然而我们经常使用一个维度的Grids和Blocks。

每个block的最大线程数量是512或者1024，具体取决于机器型号。每个Grid的Block数目通常是65525个。当超过了这些限制之后会使得GPU出现问题。

每个block和一个SM对应。在SM内部，block被分成线程的Warps。每个Warp包含32个线程，这里面的每个线程在同一时刻必须执行完全相同的指令(因为这里只有一个指令单元)。一个SM中的Warps是并行的。如果想要在一个Warp中的线程运行不同的指令，这两个任务将会按照顺序执行，这种现象叫做Warp Divergence。

lecture3

并行问题

非并行问题

并不是所有的问题都是并行问题。GPU用来计算并行问题，对于非并行问题，直接用CPU计算。

lecture4 GPU Memory Systems

NVIDIA Architecture Names

cuda-gpus

延迟和吞吐量(Latency and Throughput)

延迟是由硬件的物理速度造成的延迟。
CPU有着低延迟低吞吐量。CPU clock 3GHz,main memory latency:~100ns,arithmetic instruction latency:~1+ns

GPU有着高延迟和高吞吐量。GPU clock 1GHz,main memory latency:~300+ns,arithmetic instruction latency:~10+ns。这些参数针对 Kepler GPUs(例如GTX 700)。对于 Fermi,吞吐量大概是其两倍。

计算和IO吞吐量(Compute and IO Throughput)

IO往往是吞吐量的性能瓶颈，所以需要好的IO策略。如果象牙超过900GFLOPS,则需要执行多个FLOP共享内存负载。

缓存(Cache)

缓存是位于更大的内存池和处理器之间的一块内存,它通常在硬件级别实现，具有比它更大的内存池更快的访问速度。当请求内存时，请求的内存附近的额外内存被读入一个缓存，读取的数量是缓存和内存池特定的。始终缓存在一起的内存区域称为缓存行，这使未来访问可能在缓存中找到。这种访问被称为缓存命中并允许更快的访问，如果在在缓存中找不到访问，则称为缓存未命中（显然没有性能增益）。

GPU内存

Registers

寄存器是处理器直接使用的一块内存。我们希望尽可能使用最快的内存，寄存器的速度大概是共享内存的10倍。每个SM中有着数以万计的寄存器。通常每个线程最多可以有32或者64个32位的寄存器。在内核中声明的大多数堆栈变量都存储在寄存器中。存储在堆栈中的静态索引数组有时被放入寄存器中。

Local Memory

本地内存是堆栈中不能被寄存器容纳的所有内容。本地内存的范围只是线程，本地内存存储在全局内存中，比寄存器慢得多。

Global Memory

Global Memory是独立于GPU内核的硬件(包含SM，高速缓存，等)。GPU上绝大部分内存都是全局内存，如果数据不适合全局内存，则将以块处理它使得它适合全局内存。GPU有着0.5-24GB的全局内存，现在的设备大多在2GB左右。全局内存的延时在Kepler上大约300ns,在Fermi上大约600ns。
全局内存的IO是GPU上的IO中最慢的(当然对主机内存的通信IO更慢)。由于这个原因，我们希望能够尽可能少访问全局内存。与GPU硬件搭配得很好的访问模式被称为合并内存访问。合并内存访问能够最大限度减少读入的缓存行数量。GPU高速缓存线是128字节并对齐。内存合并在现实中其实复杂得多。

Shared Memory

共享内存是SM中非常快的内存。和L1级缓存的硬件相同，大概5ns的延迟。最大容量为48KB，是用户可配置的。共享内存的范围是block。
共享内存的分配可以通过静态和动态两种方式进行。

静态方式

shared float data[1024];
在内核中定义，在主机代码中不用做任何处理。

动态方式

Host:
kernel<<>>(args);
Device:
extern shared float s[];
定义多个动态大小的变量可以参考
using shared memory cuda cc

案例分析

任务：计算字节频率计数
输入：长度为n的字节数组
输出：存储每个数字出现次数的长度为256的数组

简单方法：在全局内存建立输出，n个全局存储空间。
聪明方法：在共享内存中构建输出，并复制到全局内存中，最终需要256个全局存储空间。

计算强度(Computational Intensity)

计算强度是必须在单个数据点（FLOPs / IO）上完成的操作数的一种表示。通常和复杂度的大O表示方法一致。
如果计算强度大于1，那么相同的数据不止用来计算一次。实践中我们需要尽可能多的共享负载和尽可能少的全局负载。

内核中的常见模式

1、将数据从全局内存拷贝到共享内存
2、__syncthreads()
3、运行计算，并递增地将运行结果存储在共享内存中，如果必要还需进行线程同步。
4、将输出从共享内存拷贝到全局内存。

Bank Conflicts

共享内存被设置为32个bank,如果将共享内存划分为4个字节长的元素，则元素i的位置将在第i%32个bank。当同一个Warp中的两个线程访问同一个bank中的不同元素时将发生bank conflicts。这种冲突将导致串行存储器访问而不是并行访问。在GPU编程中串行任何内容都对性能不利。

Register spilling

It costs:
1 extra load
1 extra store
2 extra pairs of consecutive dependent instructions

L1 Cache

Fermi-caches local& global memory
Kepler,Maxwell -only caches local momory
same hardware as shared momory
configurale size(16,32,48KB)
each SM has its own L1 cache

L2 Cache

caches all global and local memory accesses
about 1MB in size
shared by all SMs

Constant Memory

常量内存是具有特殊缓存的全局内存，用于不能编译成程序的常量，运行内核之前必须从主机设置常量，用户为64KB，编译器为64KB,内核参数通过常量内存传递。
用法：
在全局范围内(内核之外，程序的顶层)：
constant in foo[1024];
主机代码：
cudaMemcpyToSymbol(foo,h_src,sizeof(int)*1024);

Texture Memory

对于通用计算来说，复杂而且只有很小的用处
有用的特点：
●通过“CUDA阵列”进行缓存的二维或三维数据位置,进入特殊的纹理缓存。
●在1D，2D或3D阵列上进行快速插补
●将整数转换为“单位”浮点数
用例：
(1)通过纹理缓存和CUDA数组读取输入数据来利用空间缓存,这是最常见的用例。
(2)通过纹理缓存和CUDA数组读取输入数据来利用。
(3)利用数字纹理功能。
(4)与OpenGL和一般计算机图形学的交互。

Read-Only Cache(CC 3.5)

许多CUDA程序不使用纹理，但是我们应该利用纹理缓存硬件。CC>=3.5 使得使用纹理缓存更加容易。许多const限制变量将自动通过纹理缓存加载(也称为只读缓存)。也可以使用__ldg内部函数强制加载缓存。其和常量内存不同，因为不需要静态索引。

lecture5 GPU Compute Programming

Warp调度

warp调度程序找到一个准备好执行下一个指令和可执行的warp核心，然后开始执行
GK110：每个SM中有4个调度程序，2个调度程序每个时钟最多4次启动指令，并在每个经纱中启动最多2条指令。

GK110(Kepler) numbers

max threads/SM=2048(64warps)
max threads/block=1024(32warps)
32bit registers /SM=64KB
max shared memory /SM=48KB
值得注意的是：在一个SM上同时运行的块的数量取决于该资源的需求块。

占用率

occupancy=warps per SM / max warps per SM
max warps /SM depends only on GPU
warps /SM depends on warps /block, registers/block, shared momory/block。

同步

在CPU上，可以使用锁，信号量，条件变量等来解决同步问题。在GPU上，这些解决方案引入了太多的内存和处理开销。其实我们有更简单的解决方案更适合并行程序。我们使用__syncthreads()函数来同步块内的线程◦但是只能在块级别运行，因为SMs是彼此分开的，所以做不到比这更好。类似于C/C++中的barrier()函数。

lecture6 同步、共享内存、矩阵转置

同步

并行的理想状态：线程之间无资源共享，无通讯。许多只需要一点点共享资源的算法仍然能够通过大规模的并行计算来实现。

需要同步的例子

1、并行BFS
2、计算列表之和
3、向GPU的共享内存载入数据

__syncthreads()

作用是同步一个block中的所有线程。值得注意的是：共享内存是在每个block之内共享。Every block that is launched will have to allocate shared memory for its own itself on its resident SM。__synchthreads()方法对于内核共享内存很有用。

原子指令：动机

一个原子指令作为一个基本单元，不能被打断。序列化访问。
CUDA上的原子指令：
atomic{Add, Sub, Exch, Min, Max, Inc, Dec, CAS, And, Or, Xor}
Syntax: atomicAdd(float* address, float val)
在全局内存和共享内存都起作用。

使用建议

做更多的便宜操作和做更少的复杂操作。
举例：计算列表的和。
Naive方法：每个线程以原子方式递增每个数字到全局内存中的累加器。
更聪明的方法是：
.每个线程在寄存器中计算自己那部分的和。
.利用warp shuffle来计算跨越warp的和。
.每个warp以原子方式递增和到全局内存中的累加器。
.将原子操作的数量减少32倍。(32为每个warp的线程数)。

warp-synchronous programming

怎样能够使得warp中的各个线程同步呢？实际上，warp中的各个线程已经是同步的。这样可以减少__syncthreads()方法的调用次数。

Warp shuffle

从warp中的另一个线程读取寄存器的值。
int __shfl(int val,int srcLane, int width=warpSize)
对于计算整个warp的和极其有用。First avalible on Kepler(no Fermi, only CC>=3.0)