打印

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

[复制链接]
3180|15
手机看帖
扫描二维码
随时随地手机跟帖
跳转到指定楼层
楼主
本帖最后由 keer_zu 于 2019-3-6 16:48 编辑

github: cuda_study

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



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

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

windows 和 linux都可以。

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

#include <stdio.h>
#include <iostream>
#include <iomanip>
#include <cuda_runtime.h>
using namespace std;

void MatrixPrint(float *mat, int rows, int cols) {
    for (int i = 0; i < rows; i++) {
        for (int j = 0; j < cols; j++) {
            cout << setw(2) << mat[i*cols+j] << " ";
        }
        cout << endl;
    }
    cout << endl;
}

__global__ void addone(float *a) {
    int tix = threadIdx.x;
    int tiy = threadIdx.y;
    int bdx = blockDim.x;
    int bdy = blockDim.y;
    a[tix*bdy+tiy] += 1;
}
int main()
{
    int size = 5;
    float *a = (float*)malloc(sizeof(float)*size*size);
    for (int i = 0; i < size; i++) {
        for (int j = 0; j < size; j++) {
            a[i*size+j] = 1.0f;
        }
    }
    MatrixPrint(a,size,size);
    float *a_cuda;
    cudaMalloc((void**)&a_cuda,sizeof(float)*size*size);
    cudaMemcpy(a_cuda,a,sizeof(float)*size*size,cudaMemcpyHostToDevice);

    dim3 grid(1, 1, 1), block(5, 5, 1);
    addone<<<grid,block>>>(a_cuda);
    cudaMemcpy(a,a_cuda,sizeof(float)*size*size,cudaMemcpyDeviceToHost);
    MatrixPrint(a,size,size);
    return 0;
}

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

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

float *a_cuda;
cudaMalloc((void**)&a_cuda,sizeof(float)*size*size);

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

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阵列。
核函数定义如下:
__global__ void addone(float *a) 
{
    int tix = threadIdx.x;
    int tiy = threadIdx.y;
    int bdx = blockDim.x;
    int bdy = blockDim.y;
    a[tix*bdy+tiy] += 1;
}


插曲:

在调用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)概念

如下图:












软件:
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的硬件结构:


硬件:
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


程序执行结果:








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

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

基于CUDA技术的卷积神经网络识别算法.pdf (311.45 KB)

CUDA中文手册.pdf (3.41 MB)

CUDA编程模型.pdf (194.46 KB)
评论
21ic小管家 2020-1-17 09:36 回复TA
@keer_zu :Hi~图章今年规定仅推荐时长为一周左右的新文章,所以有撤销哦,实在抱歉。 
keer_zu 2020-1-16 09:10 回复TA
@21ic小管家 为什么要撤销图章?这个是GPU编程入门很好的教程啊。 
评分
参与人数 1威望 +15 收起 理由
tyw + 15 赞一个!

相关帖子

沙发
keer_zu|  楼主 | 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个。


使用特权

评论回复
板凳
keer_zu|  楼主 | 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,但需要序列执行。下图显示了软件硬件方面的术语对应关系:






  需要注意的是,大部分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,无需保存/读取上下文。




使用特权

评论回复
地板
keer_zu|  楼主 | 2019-3-6 09:42 | 只看该作者
总结几个重点地方:

核函数的使用方法:

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

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

然后定义block和thread数量:
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.





使用特权

评论回复
5
keer_zu|  楼主 | 2019-3-6 10:46 | 只看该作者
再看一个稍微复杂的例子,
两个矩阵的乘法:


直观一点:



使用cpu计算很简单:

void MatrixMul_host(float *a, int a_rows, int a_cols, float *b, int b_rows, int b_cols, float *c) {
    for (int i = 0; i < a_rows; i++) {
        for (int j = 0; j < b_cols; j++) {
            float t = 0;
            for (int k = 0; k < b_rows; k++) {
                t += a[i*a_cols+k]*b[k*b_cols+j];
            }
            c[i*b_cols+j] = t;
        }
    }
}




使用特权

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

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






使用特权

评论回复
7
keer_zu|  楼主 | 2019-3-6 14:07 | 只看该作者

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

使用特权

评论回复
8
keer_zu|  楼主 | 2019-3-6 14:07 | 只看该作者

使用特权

评论回复
9
keer_zu|  楼主 | 2019-3-6 15:18 | 只看该作者

使用特权

评论回复
10
keer_zu|  楼主 | 2019-3-6 15:28 | 只看该作者
可以调用接口获取设备信息:

使用特权

评论回复
11
keer_zu|  楼主 | 2019-3-6 16:22 | 只看该作者
选择在哪个GPU上运行程序:

使用特权

评论回复
12
keer_zu|  楼主 | 2019-3-6 16:39 | 只看该作者
再说核函数参数问题:

使用特权

评论回复
13
asimeto223| | 2019-3-21 12:13 | 只看该作者
教程很全

使用特权

评论回复
14
keer_zu|  楼主 | 2020-1-16 09:07 | 只看该作者
想玩GPU编程的可以看一下啊

使用特权

评论回复
发新帖 我要提问
您需要登录后才可以回帖 登录 | 注册

本版积分规则

1349

主题

12426

帖子

53

粉丝