连载:CUDA 高性能并行计算入门

[复制链接]
1192|14
 楼主 | 2019-3-5 21:31 | 显示全部楼层 |阅读模式
本帖最后由 keer_zu 于 2019-3-6 16:48 编辑

github: cuda_study

只要你的计算机有个nvidia的显卡,不管windows还是linux,都能很快开始入门。



因为要把识别算法运行在GPU环境,所以开始了解cuda的并行计算编程。

安装好nvidia的显卡及驱动,然后安装好cuda套件。就可以轻松开始了。

windows 和 linux都可以。

可以参照上面github代码。从第一个cuda程序开始:

  1. #include <stdio.h>
  2. #include <iostream>
  3. #include <iomanip>
  4. #include <cuda_runtime.h>
  5. using namespace std;

  6. void MatrixPrint(float *mat, int rows, int cols) {
  7.     for (int i = 0; i < rows; i++) {
  8.         for (int j = 0; j < cols; j++) {
  9.             cout << setw(2) << mat[i*cols+j] << " ";
  10.         }
  11.         cout << endl;
  12.     }
  13.     cout << endl;
  14. }

  15. __global__ void addone(float *a) {
  16.     int tix = threadIdx.x;
  17.     int tiy = threadIdx.y;
  18.     int bdx = blockDim.x;
  19.     int bdy = blockDim.y;
  20.     a[tix*bdy+tiy] += 1;
  21. }
  22. int main()
  23. {
  24.     int size = 5;
  25.     float *a = (float*)malloc(sizeof(float)*size*size);
  26.     for (int i = 0; i < size; i++) {
  27.         for (int j = 0; j < size; j++) {
  28.             a[i*size+j] = 1.0f;
  29.         }
  30.     }
  31.     MatrixPrint(a,size,size);
  32.     float *a_cuda;
  33.     cudaMalloc((void**)&a_cuda,sizeof(float)*size*size);
  34.     cudaMemcpy(a_cuda,a,sizeof(float)*size*size,cudaMemcpyHostToDevice);

  35.     dim3 grid(1, 1, 1), block(5, 5, 1);
  36.     addone<<<grid,block>>>(a_cuda);
  37.     cudaMemcpy(a,a_cuda,sizeof(float)*size*size,cudaMemcpyDeviceToHost);
  38.     MatrixPrint(a,size,size);
  39.     return 0;
  40. }
复制代码

任务是:将一个矩阵输入到Global内存中,利用GPU全部加1后返回到Host内存进行输出。
第一步是需要在CPU中创建一个矩阵,我们一般使用一维动态数组开辟,用二维的方式进行索引。
先利用Malloc函数在CPU上开辟一块空间,并全部赋值为1。

然后需要在GPU上同样开辟一个相同大小的空间以存放矩阵,这里使用cudaMalloc函数。

  1. float *a_cuda;
  2. cudaMalloc((void**)&a_cuda,sizeof(float)*size*size);
复制代码

接着,我们将矩阵从CPU上copy到GPU上。
  1. cudaMemcpy(a_cuda,a,sizeof(float)*size*size,cudaMemcpyHostToDevice);
复制代码
这时的a_cuda指向的是GPU上Device Memory上的一块地址。
GPU要如何才能运行这一块内存中的数据呢?
就是使用核函数,也叫作Kernel函数。
核函数的使用语法如下:

  1. Function<<<griddim,blockdim,extern shared memory,GPU stream>>>(param...);
复制代码


中间的参数可以控制核函数运行所占用的资源。
griddim表示调用的block块数
blockdim表示调用的thread数
后面两个参数分别表示动态定义共享内存大小和可使用的SM处理器数。
那说到这里,如何定义kernel呢?
kernel函数用__global__修饰符来修饰
下面我们就来定义一个矩阵每个元素都加 1 的kernel函数
在定义核函数之前先要考虑好需要调用多少block和thread,这里时5×5的矩阵,我们可以使用1个block和25个thread排列为5×5thread阵列。
核函数定义如下:
  1. __global__ void addone(float *a)
  2. {
  3.     int tix = threadIdx.x;
  4.     int tiy = threadIdx.y;
  5.     int bdx = blockDim.x;
  6.     int bdy = blockDim.y;
  7.     a[tix*bdy+tiy] += 1;
  8. }

复制代码

插曲:

在调用kernal函数时总体为一个Grid,Grid中含有Block,一个SM在运行时自动分配调用一些Block,每个Block中有大量的Thread。
GPU在运算时以一个Warp为单位,即32个Threads为单位,后面我们可以进行验证调度过程。
Block可以是一维的二维的三维的,Thread也是如此,一般我们选用二维作为调度结构,图中给出来索引的方式。

sp,sm,thread,block,grid,warp之间的关系要搞清楚:SP(streaming Process),SM(streaming multiprocessor)是硬件(GPU hardware)概念。而thread,block,grid,warp是软件上的(CUDA)概念

如下图:



17785c7e7ce5e246c.png






675835c7e7e4ab09fd.png

软件:
thread,block,grid,warp是CUDA编程上的概念,以方便程序员软件设计,组织线程,同样的我们给出一个示意图来表示。

    thread:一个CUDA的并行程序会被以许多个threads来执行。
    block:数个threads会被群组成一个block,同一个block中的threads可以同步,也可以通过shared memory通信。
    grid:多个blocks则会再构成grid。
    warp:GPU执行程序时的调度单位,目前cuda的warp的大小为32,同在一个warp的线程,以不同数据资源执行相同的指令,这就是所谓 SIMT。









GPU的硬件结构:

827545c7e7ec88fd0f.png
硬件:
SP:最基本的处理单元,streaming processor,也称为CUDA core。最后具体的指令和任务都是在SP上处理的。GPU进行并行计算,也就是很多个SP同时做处理。
SM:多个SP加上其他的一些资源组成一个streaming multiprocessor。也叫GPU大核,其他资源如:warp scheduler,register,shared memory等。SM可以看做GPU的心脏(对比CPU核心),register和shared memory是SM的稀缺资源。CUDA将这些资源分配给所有驻留在SM中的threads。因此,这些有限的资源就使每个SM中active warps有非常严格的限制,也就限制了并行能力。



@gaoyang9992006 @yyy71cj @chunyang @21小跑堂 @21ic小管家 @tyw


程序执行结果:

24915c7e7c6e7c544.png






=======================================================================

这里有个cuda计算卷积神经网络的讲解:

基于CUDA技术的卷积神经网络识别算法.pdf (311.45 KB, 下载次数: 2)
 楼主 | 2019-3-5 21:54 | 显示全部楼层
SM更像一个独立的CPU core。SM(Streaming Multiprocessors)是GPU架构中非常重要的部分,GPU硬件的并行性就是由SM决定的。以Fermi架构为例,其包含以下主要组成部分:

    CUDA cores
    Shared Memory/L1Cache
    Register File
    Load/Store Units
    Special Function Units
    Warp Scheduler

  GPU中每个sm都设计成支持数以百计的线程并行执行,并且每个GPU都包含了很多的SM,所以GPU支持成百上千的线程并行执行。当一个kernel启动后,thread会被分配到这些SM中执行。大量的thread可能会被分配到不同的SM,同一个block中的threads必然在同一个SM中并行(SIMT)执行。每个thread拥有它自己的程序计数器和状态寄存器,并且用该线程自己的数据执行指令,这就是所谓的Single Instruction Multiple Thread。
  一个SP可以执行一个thread,但是实际上并不是所有的thread能够在同一时刻执行。Nvidia把32个threads组成一个warp,warp是调度和运行的基本单元。warp中所有threads并行的执行相同的指令。一个warp需要占用一个SM运行,多个warps需要轮流进入SM。由SM的硬件warp scheduler负责调度。目前每个warp包含32个threads(Nvidia保留修改数量的权利)。所以,一个GPU上resident thread最多只有 SM*warp个。


使用道具

评论回复
 楼主 | 2019-3-5 21:56 | 显示全部楼层
分清SIMT和SIMD

  CUDA是一种典型的SIMT架构(单指令多线程架构),SIMT和SIMD(Single Instruction, Multiple Data)类似,SIMT应该算是SIMD的升级版,更灵活,但效率略低,SIMT是NVIDIA提出的GPU新概念。二者都通过将同样的指令广播给多个执行官单元来实现并行。一个主要的不同就是,SIMD要求所有的vector element在一个统一的同步组里同步的执行,而SIMT允许线程们在一个warp中独立的执行。SIMT有三个SIMD没有的主要特征:

    每个thread拥有自己的instruction address counter
    每个thread拥有自己的状态寄存器
    每个thread可以有自己独立的执行路径

  更细节的差异可以看这里。
  前面已经说block是软件概念,一个block只会由一个sm调度,程序员在开发时,通过设定block的属性,**“告诉”**GPU硬件,我有多少个线程,线程怎么组织。而具体怎么调度由sm的warps scheduler负责,block一旦被分配好SM,该block就会一直驻留在该SM中,直到执行结束。一个SM可以同时拥有多个blocks,但需要序列执行。下图显示了软件硬件方面的术语对应关系:



683775c7e7fd4235a4.png


  需要注意的是,大部分threads只是逻辑上并行,并不是所有的thread可以在物理上同时执行。例如,遇到分支语句(if else,while,for等)时,各个thread的执行条件不一样必然产生分支执行,这就导致同一个block中的线程可能会有不同步调。另外,并行thread之间的共享数据会导致竞态:多个线程请求同一个数据会导致未定义行为。CUDA提供了cudaThreadSynchronize()来同步同一个block的thread以保证在进行下一步处理之前,所有thread都到达某个时间点。
  同一个warp中的thread可以以任意顺序执行,active warps被sm资源限制。当一个warp空闲时,SM就可以调度驻留在该SM中另一个可用warp。在并发的warp之间切换是没什么消耗的,因为硬件资源早就被分配到所有thread和block,所以该新调度的warp的状态已经存储在SM中了。不同于CPU,CPU切换线程需要保存/读取线程上下文(register内容),这是非常耗时的,而GPU为每个threads提供物理register,无需保存/读取上下文。




使用道具

评论回复
 楼主 | 2019-3-6 09:42 | 显示全部楼层
总结几个重点地方:

核函数的使用方法:

  1. Function<<<griddim,blockdim,extern shared memory,GPU stream>>>(param...);
复制代码
<<<>>>里面是对block数量,thread数量的定义,还有共享内存和可使用的SM处理器数量。(block和thread如何选取接下来专门考虑)

这是kernel函数的原型:
  1. __global__ void addone(float *a)
复制代码
就是前面需要加 __global__,

然后定义block和thread数量:
  1. dim3 grid(1, 1, 1), block(5, 5, 1);
复制代码
dim3是一个CUDA内建的变量,它是三维的。
grid(1,1,1)给出了block的数量(1 x 1 x 1 = 1,一共1个block)
block(5,5,1)给出了thread的数量(5 x 5 x 1 = 25,每个block25个thread)
一共 1 x 25 = 25个thread。
为什么用25个thread?
因为我们要处理的矩阵是5x5的,而且只是对其每个元素加1.





使用道具

评论回复
 楼主 | 2019-3-6 10:46 | 显示全部楼层
再看一个稍微复杂的例子,
两个矩阵的乘法:


直观一点:



使用cpu计算很简单:

  1. void MatrixMul_host(float *a, int a_rows, int a_cols, float *b, int b_rows, int b_cols, float *c) {
  2.     for (int i = 0; i < a_rows; i++) {
  3.         for (int j = 0; j < b_cols; j++) {
  4.             float t = 0;
  5.             for (int k = 0; k < b_rows; k++) {
  6.                 t += a[i*a_cols+k]*b[k*b_cols+j];
  7.             }
  8.             c[i*b_cols+j] = t;
  9.         }
  10.     }
  11. }
复制代码




使用道具

评论回复
 楼主 | 2019-3-6 10:55 | 显示全部楼层
如何用GPU进行矩阵运算?让人一想就能想到的算法是,既然有block,既然有thread,那么每个block负责1行的计算,每个thread负责几列的计算应该是比较好的。

那么怎么实现kernel函数?
这里需要考虑一个显卡内存(Device Memory)读取速度的问题,Device Memory有一个特点,连续读取的速度远远高于随机读取的速度,那什么叫连续读取,这里就涉及到线程的调度问题。单个线程连续读取一块内存算是连续读取吗?错了,在GPU执行时,是一个线程读取完一组数据后直接调用下一个线程读取,而非一个线程连续读取,所以按照线程号排列的索引才是连续的索引。具体操作看kernel函数:

60195c7f35ae49c2e.png




使用道具

评论回复
| 2019-3-6 13:58 | 显示全部楼层
不明觉厉

使用道具

评论回复
 楼主 | 2019-3-6 14:07 | 显示全部楼层

机器上有NVIDIA显卡没有,装上驱动和cuda就可以开始了。

使用道具

评论回复
 楼主 | 2019-3-6 14:07 | 显示全部楼层
505345c7f63aef1387.png

使用道具

评论回复
 楼主 | 2019-3-6 15:18 | 显示全部楼层
764795c7f744183de0.png

使用道具

评论回复
 楼主 | 2019-3-6 15:28 | 显示全部楼层
可以调用接口获取设备信息:

296515c7f76992ecc6.png

使用道具

评论回复
 楼主 | 2019-3-6 16:22 | 显示全部楼层
选择在哪个GPU上运行程序:

919555c7f8346ea912.png

使用道具

评论回复
 楼主 | 2019-3-6 16:39 | 显示全部楼层
再说核函数参数问题:

267775c7f875063579.png

使用道具

评论回复
 楼主 | 2019-3-6 16:42 | 显示全部楼层
本帖最后由 keer_zu 于 2019-3-6 16:59 编辑

块(block)索引参数:

998785c7f8804c699c.png


使用道具

评论回复
| 2019-3-21 12:13 | 显示全部楼层
教程很全

使用道具

评论回复
扫描二维码,随时随地手机跟帖
您需要登录后才可以回帖 登录 | 注册

本版积分规则

我要发帖 投诉建议 创建版块 申请版主

快速回复

您需要登录后才可以回帖
登录 | 注册
高级模式

论坛热帖

关闭

热门推荐上一条 /5 下一条

快速回复 返回顶部 返回列表