GPU & CPU编程

2

3

4

5

6

7

8

9

10

/*

另外：Block和Thread都有各自的ID，记作blockIdx（1D，2D），threadIdx（1D，2D，3D）

Block和Thread还有Dim，即blockDim与threadDim.
他们都有三个分量x，y，z

线程同步：void
__syncthreads(); 可以同步一个Block内的所有线程

总结来说，每个
thread 都有自己的一份 register 和 local memory 的空间。

一组thread构成一个
block，这些 thread 则共享有一份shared memory。

此外，所有的
thread(包括不同 block 的 thread)都共享一份

global
memory、constant memory、和 texture memory。

不同的
grid 则有各自的 global memory、constant memory 和 texture memory。

*/

2

3

4

5

6

7

per-thread

register                            
1 cycle

per-thread

local memory                     slow

per-block
shared memory                   1 cycle

per-grid
global memory                       500 cycle,not cached!!

constant
and texture memories            500 cycle, but cached and read-only

分配内存：cudaMalloc，cudaFree，它们分配的是global
memory

Hose-Device数据交换：cudaMemcpy

2

3

4

5

__device__  
//
GPU的global memory空间，grid中所有线程可访问

__constant__
//
GPU的constant memory空间，grid中所有线程可访问

__shared__  
//
GPU上的thread block空间，block中所有线程可访问

local       
//
位于SM内，仅本thread可访问

//
在编程中，可以在变量名前面加上这些前缀以区分。

2

3

4

5

6

7

8

9

//
内建矢量类型：

int1，int2，int3，int4，float1，float2，
float3，float4 ...

//
纹理类型：

texture<Type,
Dim, ReadMode>texRef;

//
内建dim3类型：定义grid和block的组织方法。例如：

dim3
dimGrid(2, 2);

dim3
dimBlock(4, 2, 2);

//
CUDA函数CPU端调用方法

kernelFoo<<<dimGrid,
dimBlock>>>(argument);

2

3

4

5

6

7

8

9

10

__device__
//
执行于Device，仅能从Device调用。限制，不能用&取地址；不支持递归；不支持static variable；不支持可变长度参数

__global__
//
void： 执行于Device，仅能从Host调用。此类函数必须返回void

__host__
//
执行于Host，仅能从Host调用，是函数的默认类型

//
在执行kernel函数时，必须提供execution configuration，即<<<....>>>的部分。

//  
例如：

__global__
void

KernelFunc(...);

dim3
DimGrid(100, 50); //
5000 thread blocks

dim3
DimBlock(4, 8, 8); //
256 threads per block

size_t

SharedMemBytes = 64; //
64 bytes of shared memory

KernelFunc<<<
DimGrid, DimBlock, SharedMemBytes >>>(...);

2

CUDA包含一些数学函数，如sin，pow等。每一个函数包含有两个版本，

例如正弦函数sin，一个普通版本sin，另一个不精确但速度极快的__sin版本。

2

3

4

5

/*

gridDim,
blockIdx, blockDim,

threadIdx,
wrapsize.

这些内置变量不允许赋值的

*/

2

3

4

5

6

7

/*

目前CUDA仅能良好的支持C，在编写含有CUDA代码的程序时，

首先要导入头文件cuda_runtime_api.h。文件名后缀为.cu，使用nvcc编译器编译。

目前最新的CUDA版本为5.0，可以在官方网站下载最新的工具包，网址为：

https://developer.nvidia.com/cuda-downloads

该工具包内包含了ToolKit、样例等，安装起来比原先的版本也方便了很多。

*/

2

3

4

5

6

7

8

9

10

11

12

13

14

15

16

17

18

19

20

21

22

23

24

25

26

27

28

29

30

31

32

33

34

35

36

37

38

39

40

41

42

43

44

45

46

47

48

49

50

51

52

53

54

55

56

57

58

59

60

61

62

63

64

65

66

67

68

69

70

71

72

73

74

75

76

77

78

79

80

81

82

83

84

85

86

87

88

89

90

91

92

93

1
GPU硬件

//
i GPU一个最小单元称为Streaming Processor(SP)，全流水线单事件无序微处理器，

包含两个ALU和一个FPU，多组寄存器文件（register

file，很多寄存器的组合），

这个SP没有cache。事实上，现代GPU就是一组SP的array，即SPA。

每一个SP执行一个thread

//
ii 多个SP组成Streaming Multiprocessor(SM)。

每一个SM执行一个block。每个SM包含8个SP；

2个special
function unit(SFU)：

这里面有4个FPU可以进行超越函数和插值计算

MultiThreading
Issue Unit：分发线程指令

具有指令和常量缓存。

包含shared
memory

//
iii Texture Processor Cluster(TPC) ：包含某些其他单元的一组SM

2
Single-Program Multiple-Data （SPMD）模型 

//
i CPU以顺序结构执行代码，

GPU以threads
blocks组织并发执行的代码，即无数个threads同时执行

//
ii 回顾一下CUDA的概念：

一个kernel程序执行在一个grid
of threads blocks之中

一个threads
block是一批相互合作的threads：

可以用过__syncthreads同步；

通过shared
memory共享变量，不同block的不能同步。

//
iii Threads block声明：

可以包含有1到512个并发线程，具有唯一的blockID，可以是1,2,3D

同一个block中的线程执行同一个程序，不同的操作数，可以同步，每个线程具有唯一的ID


3
线程硬件原理

//
i GPU通过Global block scheduler来调度block，

根据硬件架构分配block到某一个SM。

每个SM最多分配8个block，每个SM最多可接受768个thread

（可以是一个block包含512个thread，

也可以是3个block每个包含256个thread（3*256=768！））。

同一个SM上面的block的尺寸必须相同。每个线程的调度与ID由该SM管理。

//
ii SM满负载工作效率最高！考虑某个Block，其尺寸可以为8*8,16*16,32*32

8*8：每个block有64个线程，

由于每个SM最多处理768个线程，因此需要768/64=12个block。

但是由于SM最多8个block，因此一个SM实际执行的线程为8*64=512个线程。

16*16：每个block有256个线程，SM可以同时接受三个block，3*256=768，满负载

32*32：每个block有1024个线程，SM无法处理！


//
iii Block是独立执行的，每个Block内的threads是可协同的。

//
iv 每个线程由SM中的一个SP执行。

当然，由于SM中仅有8个SP，768个线程是以warp为单位执行的，

每个warp包含32个线程，这是基于线程指令的流水线特性完成的。

Warp是SM基本调度单位，实际上，一个Warp是一个32路SIMD指令

。基本单位是half-warp。

如，SM满负载工作有768个线程，则共有768/32=24个warp

，每一瞬时，只有一组warp在SM中执行。

Warp全部线程是执行同一个指令，

每个指令需要4个clock

cycle，通过复杂的机制执行。

//
v 一个thread的一生：

Grid在GPU上启动；

block被分配到SM上；

SM把线程组织为warp；

SM调度执行warp；

执行结束后释放资源；

block继续被分配....

4
线程存储模型

//
i Register and local memory：线程私有，对程序员透明。

每个SM中有8192个register，分配给某些block，

block内部的thread只能使用分配的寄存器。

线程数多，每个线程使用的寄存器就少了。

//
ii shared memory：block内共享，动态分配。

如__shared__
float

region[N]。

shared
memory 存储器是被划分为16个小单元，

与half-warp长度相同，称为bank，每个bank可以提供自己的地址服务。

连续的32位word映射到连续的bank。

对同一bank的同时访问称为bank
conflict。

尽量减少这种情形。


//
iii Global memory：没有缓存！容易称为性能瓶颈，是优化的关键！

一个half-warp里面的16个线程对global
memory的访问可以被coalesce成整块内存的访问，如果：

数据长度为4,8或16bytes；地址连续；起始地址对齐；第N个线程访问第N个数据。

Coalesce可以大大提升性能。

//
uncoalesced

Coalesced方法：如果所有线程读取同一地址，

不妨使用constant
memory；

如果为不规则读取可以使用texture内存

如果使用了某种结构体，其大小不是4
8 16的倍数，

可以通过__align(X)强制对齐，X=4
8 16