GPU之CUDA 第一弹

1.GUP与CUDA

GPU：Graphics Processing Units

CUDA：Compute Unified Device Architecture

2.主要概念

2.1 主机：CPU+内存

2.2 设备：GPU+显存

2.3 Grid（线程格）：一个Kernel对应一个Grid，一个Grid包含若干Block

2.4 Block（线程块）：Kernel以Block为单位执行，各Block并行运行，无法通信，没有执行顺序

2.5 thread（线程）：

可以认为线程是执行CUDA程序的最小单元。

GPU资源充裕时，所有线程都并发执行；GPU资源少于总线程个数时，部分线程需等待前面执行的线程释放资源，即变为串行化执行。

2.6 warp（线程束）：32个并行线程为warp块

2.7 Kernel（核函数）：

通过__global__定义，在主机端调用（需先分配空间再调用）。

调用时声明执行参数<<< m,n,a,b >>>，m个线程块，每个线程块有n个线程，每个线程块使用了a字节的动态分配的共享存储器，内核在流b中执行。

Kernel调用是异步的，即主机仅把要执行的Kernel顺序提交给GPU，并不等待其执行完成。这种异步调用导致Kernel无法返回函数值。

2.8 dim3结构类型：

3个由unsigned int型组成的结构体，unsigned int x;unsigned int y;unsigned z

gridDim为grid尺寸，blockDim为block尺寸，blockIdx为block在grid中的索引值，threadIdx为线程索引号。

CUDA定位公式：tid=blockIdx.x*blockDim.x+threadIdx.x

3.常用内存函数

3.1 cudaMalloc((void**)&d_a,sizeof(float)*n)：在GPU上分配内存

3.2 cudaMemcpy(d_a,a,sizeof(float)*n,cudaMemcpyHostToDevice)：在主机内存和GPU显存间互传数据（上述为主机到设备）

3.3 cudaFree(d_a)：释放内存

4.GPU内存分类

4.1 全局内存(global memory)：

可读写

通俗意义上的设备内存

访问全局内存造成的延时很高，高达访问寄存器变量延时的600倍。GPU的数据重用是达到高性能的关键。

高性能CUDA程序应充分利用GPU的内部资源（寄存器、共享内存等）来避开全局内存瓶颈。

4.2 共享内存(shared memory）：

可读写

位于设备内存，添加关键字__shared__声明

4.3 寄存器（registers）：

可读写

4.4 Local memory：

可读写

4.5 常量内存(constant memory）：

只读

4.6 纹理内存(texture memory）：

只读

总结：每个线程都有自己的registers和local memory

    同一个block中的每个线程共享一份共享内存

    所有线程（包括不同block的thread）都共享一份全局内存、常量内存、纹理内存。

    不同grid有各自的全局内存、常量内存、纹理内存。

5.SM与SP

5.1 SM（流多处理器）：

同block的线程被分配在一个SM中

当前CUDA定义中，一个SM中至多分配8个blocks 

5.2 SP（标量流处理器）：

GPU中最基本的处理单元即为SP

在 G80/G92 的架构下，有 128 个 SP，以 8 个 SP 为一组，组成 16 个 SM，再以两个 SM 为一个 TPC，共分成 8 个 TPC 来运作。

在新一代的 GT200 里，SP 则是增加到 240 个，还是以 8 个 SP 组成一个 SM，但是改成以 3 个 SM 组成一个 TPC，共 10 组 TPC。

6.错误处理

cudaErrot_t：CUDA的错误类型。

cudaError_t cudaStatus;  
cudaStatus = cudaSetDevice(0);  
if (cudaStatus != cudaSuccess) {  
        fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");  
        goto Error;  
}  

7.函数限定符

7.1 __device__：在GPU上调用，在GPU上执行。

7.2 __global__：在CPU上调用，在GPU上执行。

7.3 __host__：在CPU上调用，在CPU上执行。

8.变量限定符

8.1 __device__：声明的数据存放在显存中，所有线程均可访问，且主机也可以通过运行时库访问。

8.2 __shared__：数据存放在共享存储器上，只有所在块内的线程可以访问，其他块内的线程不能访问。

8.3 __constant__：数据存放在常量存储器上，可以被所有线程访问，主机也可以通过运行时库访问。