Access the Power of GPU
- 一般有三种方法来实现加速：Libraries or OpenACC or Programming Languages
- 函数库：
例如：
cuBLAS: 线性代数库
cuSPARSE: 稀疏矩阵的线性代数库
cuDNN: 深度神经网络
GPU Architecture
异构计算(Heterogeneous Computing)
Terminology:
- Host: The CPU and tis memory(host memory) 主机端(CPU)
- Device: The GPU and its memory(device memory) 设备端(GPU)
CPU与GPU的简易架构图
- CPU有很多逻辑控制单元，而GPU有很多处理核心
- CPU是缓存优化的处理器，缓存特别大，而且有很多控制单元
- GPU是一个并行吞吐优化的处理器，计算核心特别多
- 因此我们知道，GPU的单精度处理能力比CPU快很多的原因其实就是GPU将更多的晶体管用于了计算而不是缓存
程序的开发中如何选择CPU和GPU
- 从图中可以看到CPU的缓存延迟非常短，这也对应它的架构里有很多复杂的逻辑控制
- GPU则不同，图中的W是warp，1 warp = 32 Threads，GPU是以warp为单位进行切换的。我们可以通过海量的线程的切换来对缓存延迟做隐藏，如果你的程序是计算密集型的，而且有很好的并行特性，那么gpu是更好的选择。
- GPU的架构有两个部分组成，第一部分是计算单元，第二部分是内存单元
- 图中的绿色部分就是一个CUDA Core
- 从图中可以看到，GPU和CPU是通过PCle进行链接的
- 一个GPU有很多SM（流多处理器）组成，每一个SM有很多核，因此有越多的SM，那么GPU就能在同一时间内处理更多的任务。
SM的架构
- Kepler架构中的SM称为SMX，表示处理器更强劲，其逻辑控制单元是一个整体，控制192个CUDA Cores
- Maxwell架构中的SM称为SMM，其逻辑控制单元有4个，每个控制32个CUDA Cores
- 可以简单认为，一个SMM中包含了4个SMX，这样的话就避免了因为逻辑控制单元过少而造成核心的冗余，效率更高。从架构上来看，Maxwell一个CUDA核心相比于Kepler提升了35%
Memory Hierarchy
因为经常要跟内存打交道，所以需要对GPU的内存架构非常熟悉，GPU的架构分为三个层次，如图所示，和CPU类似，可以看做成一个金字塔，从下往上越来越快，首先是Global Memory，有很大的内存空间，它的空间很大，但是缓存延迟很高，因此我们需要加入一个缓存，称之为L2缓存，GPU的缓存要像Global Memory一样读取数据的话，首先我们要看这些数据是否在L2中已经有缓存，有的话称之为缓存命中，这样就不用再从Global Memory里面读数据，这样就加快了读写。L2缓存是多个缓存所共享的。再上面有很多寄存器，L1缓存等。对于不同的内存有不同的使用和优化方案。
GPU in Computer System
CUDA Programming Basics
配置CUDA开发环境
- 注意如果使用C/C++进行开发的话，我们使用的是NVCC的编译器
Heterogeneous Computing 异构编程
一般我们把GPU运行的程序构造成一个函数，在我们需要的时候进行调用，这个函数就是在GPU上运行
-
这个函数的之前和之后都是在CPU上单线程运行的，因此我们把整个程序分为：顺序执行，并行执行，顺序执行
顺序执行就是利用CPU的一个线程来执行
并行执行就是用GPU的多线程来进行计算
因此如果一部分功能是计算密集型的，我们就可以把它写成一个函数，对它进行并行，并不是把所有的代码都在GPU上运行，而是把需要的部分放到GPU上
CUDA Kernels
这里我们讲解一个概念，我们要将一部分代码在gpu上运行，这部分代码是一个函数，我们称为一个kernel（核函数），CPU（Host）执行functions，GPU（Device）执行kernels。
Hello World
//hello_world.c
#include <stdio.h>
void hello_world_kernel()
{
printf("Hello World\n");
}
int main()
{
hello_world_kernel();
}
Compile & Run:
gcc hello_world.c
./a.out
在GPU上执行:
//hello_world.cu
#include <stdio.h>
_global_ hello_world_kernel()
{
printf("Hello World\n");
}
int main()
{
hello_world_kernel<<<1.1>>>();
}
Compile & Run:
nvcc hello_world.cu
./a.out
不同之处
- 文件后缀名改为.cu
- global表示了该函数为一个核函数，标识这个函数会在gpu上运行
- “<<<...,...>>>”具体含义是GPU中的一个配置
- 编译器改为nvcc
GPU memory management
hello world程序非常简单明了，但是在编程过程中我们经常要和内存打交道，那么如何进行GPU内存管理呢？
我们很熟悉对CPU的内存管理，我们会使用
malloc();//动态分配内存
memset();//对空间进行初始化
free();//释放一段内存
对应于，在GPU中也有类似的函数，它们分别是
cudaMalloc(void** pointer, sie_t nbytes);//pointer指定空间，大小是nbytes
cudaMemset(void* pointer, int value, size_t count);
cudaFree(void* pointer);
例如：
int nbytes = 1024*sizeof(int);
int* d_a = 0;
cudaMalloc((void**)&d_a, nbytes);
cudaMemset(d_a, 0, nbytes);
cudaFree(d_a);
因此一定要分清楚，哪部分数据在cpu上，哪部分数据在gpu上
Data Copies
__host__ cudaMemcpy(void* dst, void* src, size_t nbytes, cudaMemcpyKind direction);
- global是在gpu上进行的，host就是在cpu上执行
- 四个形参：拷贝目的指针，拷贝原指针，拷贝大小，拷贝方向
- 线程阻塞的，拷贝不完成，cpu不会执行下面的代码
- 常见拷贝方向有：cudaMemcpyHostToDevice（CPU到GPU）、cudaMemcpyDeviceToHost（GPU到CPU）、cudaMemcpyDeviceToDevice（GPU到GPU）
- 设定拷贝方向时，需要注意src和dst的指针是CPU的还是GPU的，要和拷贝方向一致
- 上述函数为线程阻塞的，同样还有一个异步的函数：cudaMemcpyAsync();
三步流程
第一步：通过调用cudaMemcpy函数将CPU中的数据拷贝到GPU中，通过PCI Bus，方向就是HostToDevice
第二步：通过CPU启动kernel核函数，然后开始一个并行计算
Kernel<<<grid, block>>>();
第三步：计算完以后，再调用cudaMemcpy将数据从GPU端拷贝到CPU端，方向是DeviceToHost
例子:
代码分为五个步骤：
- 在CPU端分配n个整型变量的空间
- 在GPU端分配n个整型变量的空间
- 初始化GPU内存为0
- 将数据从GPU拷贝到CPU
- 打印变量
#include<stdio.h>
int main()
{
int dimx = 16;
int num_bytes = dimx * sizeof(int);
int *d_a = 0, *h_a = 0;//设备端和主机端的指针
//在CPU端分配n个整型变量的空间
h_a = (int*)malloc(num_bytes);
//在GPU端分配n个整型变量的空间
cudaMalloc((void**) &d_a, num_bytes);
if(0 == h_a || 0 == d_a)
{
printf("Couldn't allocate memory\n");
return 1;
}
//初始化GPU内存为0
cudaMemset(d_a,0,num_bytes);
//将数据从GPU拷贝到CPU
cudaMemcpy(h_a, d_a, num_bytes, cudaMemcpyDeviceToHost);
//打印变量
for(int i=0; i < dimx; i++)
{
printf("%d\t", h_a[i]);
}
free(h_a);//主机端释放
cudaFree(d_a);//设备端释放
return 0;
}
- 这里注意，cudaMemcpy是CPU阻塞的，也就是说，不执行完该函数CPU就不会一直向下执行。如果换成一个异步的版本，CPU会在拷贝还没有完成的时候就开始打印，打印的时候很多是错误的。
上面的示例代码缺少了GPU上面的并行计算，那么我们下一面将开始讲解如何写Kernel函数。
