CPU—GPU并行处理—CUDA编程从想入门到放弃

这几天收拾准备准备GUP“库达”计算，开学搞搞计算机图形图像处理。安装问题在另一篇博文里有记载。在学习库达前，首先来阐明一些计算机体系结构概念知识，方便入手:

CPU体系架构

什么是CPU？

工作流程简化为：取指——>译码——>执行——>访存——>写回

流水线

上述的内部工作流程可以简化成一个流水线
极大地减少了时钟周期，但增加了一些芯片面积。
打个比方，工厂进行加工要走一系列流程，如果它的每个细节处理都安排员工来专门做这事，就做的越来越快，减少了制造这个产品的总时间，但是同时工厂也得雇佣更多的员工增加了工厂的规模和开支。
在流水线处理时，要注意依赖关系，并不是流水线越长越好
旁路：相互依赖的指令操作时，可以开一个（旁路）后门不进行上述的写回等工序直接引用到下一个指令，
停滞：流水线的停滞，就拿工厂来说，中间的一个步骤罢工，后面的步骤只能等待。

分支预测

不能消极的等待，所以我们做了分支预测，根据停滞前的那条指令来猜测下一条指令。
猜测还是要有一定的策略基础并非完全玄学，考试遇到不会的题，有的人猜对那可能是玄学，但有的人还是根据自己平时的积累进行排除选项提高预测概率。
具体方法给社会贡献了多少博士与论文啊（嘎嘎嘎）

超标量superscalar

提升IPC（ CPU 每一时钟周期内所执行的指令多少）,增大流水线宽度。也就如工厂一道工序进行加工的产品数目增加
但同时用的资源也多了，旁路网路增多，寄存器，存储器增多。工序占得地方大了，运载的车也就多了

乱序执行OoO（OUT of Order）

我们写的一个程序逻辑上认为计算机也会按照我们的步骤来进行，但是为了充分挖掘计算机内部的吞吐率，CPU进行重排指令。

存储器层次

计算机的主要工作占绝大数是进行访问存储器——数据的搬来搬去 。但存储器的访问速度又很慢，所以进行改良：
缓存：利用将数据放在尽可能接近的位置来供下一次更快的使用。（举个难忘的例子：我大学体育选的是乒乓球，到了期末考试，要进行一分钟55次攻防球来回，我准备了俩球，把另一个球放在了裤兜里，重要的是我还把拉链拉上了，再次取 的时候我花了好几秒，导致那次没有达成目标，第二次我直接找旁边的一女生借，虽然事先没有说明，但不得不说快多了！不出意外，有点懵逼）大概就是这么个意思。
有了前面一些简单概念的铺垫，接下来开始正式的cuda介绍。其实，cuda 的产生就是为了给CPU 分担一些计算的任务等，加快程序数据的吞吐率：

CUDA编程

1.0 编程模型

CUDA程序构架分为两部分：Host和Device。一般而言，Host指的是CPU，Device指的是GPU。在CUDA程序构架中，主程序还是由CPU来执行，而当遇到数据并行处理的部分，CUDA 就会将程序编译成GPU能执行的程序，并传送到GPU。而这个程序在CUDA里称做核（kernel）。CUDA允许程序员定义称为核的C语言函数，从而扩展了C语言，在调用此类函数时，它将由N个不同的CUDA线程并行执行N次，这与普通的C语言函数只执行一次的方式不同。执行核的每个线程都会被分配一个独特的线程ID，可通过内置的threadIdx变量在内核中访问此ID。在 CUDA 程序中，主程序在调用任何GPU内核之前，必须对核进行执行配置，即确定线程块数和每个线程块中的线程数以及共享内存大小。

---

1. 1核函数

1. 定义：在GPU进行的函数通常称为核函数
2. 一般通过__global__修饰（在核函数里，都用双下划线来修饰），调用通过<<<参数1，参数2 >>>，第一个参数代表block线程块数目，第二个参数代表线程块内含有的线程数目thread。
3. 根据2可以看出它的操作单位是总的线程格的block。
4. 调用时必须声明核函数的执行参数。
5. 正式编程前必须使用cudaMalloc（）函数先为wrap分配适当的内存空间。分配前要精打细算，因为总的内存空间一定，一个wrap分配不当的空间会导致其他线程内存不够用导致越界报错甚至死机。（字面上 还是蛮难的）。

---

任何程序入门编程：

#include<cuda_runtime.h>
#include<stdio.h>
int main()
{
printf("Hello world\n");
}

乍一看，除了头文件多了一个以外，其他毫无区别。常说cuda编程用的是c语言，难道真的一成不变？
答：cuda编程面向GPU，但如果程序里并没有用到GPU处理，那就真的是一模一样。但如果上述程序你真的不想让显卡就这么闲着，你可以这样：

// 首先，修改文件后缀为.cu
#include<stdio.h>
#include<cuda_runtime.h>
//第一步：核函数申明
__global__void kernel(void)//注意：这里的返回值只能为void
{

}
int main(void)
{
//核函数的调用在主函数
kernel<<<1,1>>> ()//启用一个block的一个thread
printf("Hello,world\n");
return 0;
}

---

1.2 dim3结构类型

• dim3是基于unit3定义的矢量类型，表示3个数据成员unsigned int x;unsigned int y;unsigned int z;可用来表示一维、二维、三维线程块。
• 相关内置变量
  2.1 gridDim,线程格的维度定义，即gridDim.x,gridDim.y,gridDim.z。
  2.2 blockDim，线程块的维度，即同上。
  2.3blockIdx,线程块ID的索引，及同上。
  2.4threadIdx,线程ID的索引，即同上。
• 一维的block，包含的线程ID，threadId=threadIdx.x，但对于二维三维，存在以下计算公式：threadID=threadIdx+threadIdx.yblockDim.x;threadID=threadIdx.x+threadIdx.yblockDim.x+threadIdx.zblockDim.yblockDim.z。
• 计算线程索引偏移量为启动线程的总数。如stride=blockDim.x*gridDim.x;threadId+=stride

---

1.3函数修饰符

kernel核函数，在主机端调用kernel在设备端创建多个线程。

---

1.4常用的内存函数

cudaMalloc()

• 用处：用来分配全局存储器（提示：少用，容易沾满）

• 函数原型：cudaError_t cudaMalloc(void **devPtr,size_t size)

• 将cudaMalloc（）分配的指针传递给设备和主机上。
  cudaMemcpy()

• 用处：进行CPU与GPU的数据拷贝

• cudaMemcpy（Md(gpu),M(原地址i)cudaMemcpyHostToDevice）

cudaFree（）

• 函数原型：cudaError_t cuda Free
• 用处：释放内存与c函数free（）一样。

---

1.5 GPU内存类型

全局内存

global memory，通常意义上的设备内存。

共享内存

• shared memory，GPU的每个线程块位于一个相同的核处理器SM，之间共享内存。
• 形式：__shared__float a[i]

常量内存

• constant memory,常量内存代替全局内存能有效的减少内存带宽。
• 要求：当我们要将数据拷贝到常量内存时，要使用cudaMemcpyToSymbol（），使用cudaMemcpy仍会拷贝到全局。

纹理内存

• Texture memory从硬件角度来说，纹理内存实际上是存储在全局内存上；唯一的区别是，当某一个变量绑定纹理内存后，在程序运行过程中会将部分信息存储在纹理缓存中，以减少线程块对全局内存的读取，进而提高程序的运行速度。
• 使用：有两种方法：纹理参考与纹理对象。
• 一维纹理内存
• • 用texture+类型声明。
• • cudaBindTexture（）绑定到纹理内存中。
• • tex1Dfetch()读取纹理内存的数据。
• 二维纹理内存
• • texture<float，2> texIn。
• • cudaBindTexture2D（）绑定到纹理内存中。
• • 通过tex2D()来读取纹理内存中的数据。
    先写这里，补作业了！

---

emmm，我太懒了！

固定内存

将数据保存在主机内存里，确保了该内存始终驻留在物理内存。

• 好处：如果知道主机内存的物理地址，可以直接通过“直接内存访问”DMA技术来在GPU和主机之间复制数据。由于DMA介入并不需要访问CPU，大大提高了访问速度。这里简单介绍一下DMA 的工作原理：CPU将DMA初始化后，DMA控制器获得权限，CPU即进入后台，由DMA控制器输出命令，直接控制I/O接口进行DMA传输。详细请戳——>百度百科DMA
• 形式：通过cudaHostAlloc（）函数来分配，通过cudaFreeHost（）释放。

---

1.6 原子性

• atomic functions,操作的执行过程分解为更小的部分。
• 函数的调用：atomicAdd(addr，y)生成一个院子的操作序列，这个操作序列包括读取addr(相应地址)出的值,将y增加到该值。结果保存到addr地址。

---

1.7 流

• 粒度是计算机领域中粒度指系统内存扩展增量的最小值。CUDA 的线程组成为Grid<-block<-thread结构。一组thread并行处理为一个block并行，一组block并行并行为一个Grid。那我们可不可以使一组Grid并行处理？因此引入–流。
• 用法：流可以实现一个设备上的多个核函数的并行运行，也可以使同一个核函数传递不同的参数。
• CUDA中的流应用cudaStream_t类型实现，用到的API有：cudaStreamCreate(cudaStream_t * s)用于创建流，cudaStreamDestroy(cudaStream_t s)用于销毁流，cudaStreamSynchronize()用于单个流同步，cudaDeviceSynchronize()用于整个设备上的所有流同步，cudaStreamQuery()用于查询一个流的任务是否已经完成。具体的含义可以查询API手册。
  这里面有很多工具的使用手册，如CUDA_GDB，Nsight，CUDA_Profiler等，方便调试程序；还有一些有用的库，如CUFFT是专门用来做快速傅里叶变换的，CUBLAS是专用于线性代数（矩阵、向量计算）的，CUSPASE是专用于稀疏矩阵表示和计算的库。这些库的使用可以降低我们设计算法的难度，提高开发效率。另外还有些入门教程也是值得一读的，你会对NVCC编译器有更近距离的接触。

     cudaStreamCreate(&stream[i]);   //创建流
      addKernel<<<1,1,0,stream[i]>>>(dev_c+i, dev_a+i, dev_b+i);    //执行流 
      ........
       cudaStreamDestroy(stream[i]);   //销毁流  

<<<>>>中的参数多了两个，前俩个没变，仍然是线程块数和线程数。第三个0表示每个block用到的共享内存大小，第四个为流对象，表示当前核函数在哪个流上运行。

2. 多个流的使用
   首先，在添加到流的顺序前，应该考虑他们之间的依赖性。就比如：大扫除，每个流程之间合理搭配才能做到省时省力。但是，分配不当，比如：你可以先打扫门窗，扫地，拖地后就能干别的事等地干了。但如果你拖完地才开始大扫门窗，导致你之后又得拖，还得刻意的等待。
   这里引用大佬的图片便于理解
   高效地运用：
   小时候我们就被数学老师熏陶“优化运筹”， 将流交叉添加，干什么一次干完：

就变成了这样。

---

1.8 线程同步

按照GPU的并行性处理原理，每一个线程块中的线程数要求相同且相互独立。但是，就拿从一个地方往另一个地方搬砖来说，如果你不知道其他伙伴搬了多少砖，也就不知道原来的地方还剩多少砖，自己每次需要怎么搬省力省时，就显得比较乱。所以可见，线程之间的交互指挥多么重要！
线程通信在CUDA中的常用实现方式：
1.共享存储器shared memory；
2.线程同步；

• block位于一个流多处理器（SM），其中，共享处理器位于SM的特殊存储器。一个SM含有若干个流处理器（SP）。还有一部分高速Cache
• 
  如图，这个SM中含有M个SP，shared memory被这M个SP共同占有。同时，指令单元也被这M个SP共享（俗称，单指令多线程SIMT架构。而CPU采用MIMD多指令多线程，从这个角度也可以看出GPU运算的快捷性）。
• 线程同步，块内线程可以同步：
• • 调用：__syncthreads创建一个barrier栅栏，每个线程在调用点等待块内所有线程执行到这个地方，然后继续后续指令。
• 
  同步从理论上来讲破坏并行性的独立性，他会导致线程的暂停，甚至会带来程序的死锁，就比如
  下面展示一下

if  (someFunc())
{
	__syncthreads();
}
else
{
	__syncthreads();
}

有的程序执行到是，有的程序执行到否，导致线程并没有达到同步更不可能达到同步，程序就会死锁。打个比方：你约好和班里的几名同学一起吃饭，不等到互相别吃饭，你在楼的西侧等他，他在楼的东侧等你，都没有等到对方，也就意味着在那一直等 ， 死锁。

---

1.9 加速比

对一个程序进行加速的时候，很多时候要预先估算出该程序使用GPU加速后的加速比，尤其做项目的时候！（从大佬那里听来的~）

如何估算程序的GPU加速比？

Amdahl阿姆达尔定律

通俗来讲，N个处理器：（（1/（P（并行）/N+S串行））P,S表示并行与串行的计算所占比例P+S=1使尽可能的系统并行化，由公式可得，程序中可并行代码的比例决定你增加处理器（总核心数）所能带来的速度提升的上限。

---

另外，GPU加速比针对的是每一个步骤。即一个程序M由P、A、B若干步骤组成，对P步骤进行加速，则GPU加速比=P步骤加速前的时间/P加速后的时间，而不是S程序加速前的时间/S加速后的时间这个非常重要，因为很多人不懂GPU加速，他们往往只关注整个程序的加速效果。而对于一个完整的程序，很多都包括一些非常耗时的操作（比如读取图片等）。这种情况下就算你对程序中的某个步骤使用GPU加速到极致，整个程序可能只感觉快了一点点，这是非常不公平的。

估算加速比

通常情况下有几种惯用的手段：

1. CPU来实现大的for循环
2. 资源的管理（内存、显存的申请与释放等）
3. 核函数级别的优化

---

（1）原程序CPU完成，程序并行化程度高，且依赖性低
这种为最理想的情况，常见于简单的图像算法。通俗来讲即for循环，循环的数据之间没有依赖性，在这种情况下，加速比容易破10+，图片越大，加速比也能达到更高。
（2）原程序GPU上完成，资源管理差
程序已经由GPU核函数或者调用GPU API实现，但是资源（显存和句柄的申请和释放等）管理比较差。比如之前在流那节讲到的高效运用多流这种情况下，看资源申请所占的时间比例，不过一般加速比能到几倍就很不错了。
（3）原程序GPU上完成，但核函数实现较差
如果只剩下对核函数进行优化，那么一般可得到的加速比就更低了，大概在2X.

---

1.10 线程调度

面临的现状是我们从软件上启动的线程数目远远大于硬件上的可用的执行单元数目。
可以参考NVIDIA公司对显卡的介绍以及cuda开发手册查看显卡的SM，sp,block数。

这里可以简单理解为：一张桌子有8张椅子，32个人要同步吃饭。就保证8为周期，进行4次，每一人恰一口。
wrap：针对的是一个block里的一组线程，32个线程为一个单位，有序号排列0-63，两组wrap，如果一个块大小不是32的倍数，最后一个wrap会被填充。
特征：一个SM在调度过程中只允许一个wrap来占用。
注意：wrap里的线程需要天然同步，wrap内部线程沿不同分支执行时，没办法为每一个ALU都设置一个逻辑复杂的机构，整个系统的速度会严重下降！
利用线程之间互相独立的代码来实现延迟掩藏。

---

1.11 并行规约

通俗来说就是将一组很大的数据通过算法化简成一个个小块，

线程之间交叉进行加和，每一轮需要的线程数目减半，但是由于wrap分割的限制（线程同步
），每一轮闲置下来的线程并不能真正的释放自由。这里改变一下线程偏移量顺序：

右面改变了线程偏移量，由1,2，4——>4,2,1.但从根本上提高了加速比。这里回忆wrap，一个wrap管理8个线程同步，右面的一半线程会提早全部结束，就可以进行内存释放等其他操作，减少了资源占用所以，要避免分支发散。
（我感兴趣的是它在医学CT等图像处理方面应用的，特意跑去知网看了看，也终于第一次看完了几篇完整的论文，里面涵盖的知识点多且片面，大都是对基于系统的介绍，怪不得老师会说这方面的论文层出不穷，但国产一个操作系统的真正实现还只是星星点点）

---

好了，框架就这基本要素。

就这？就这？之后会出一个CUDA项目实战，再偶尔看看c++，就就…开学了吧。