连载：CUDA 高性能并行计算入门
本帖最后由 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 << " ";
}
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 += 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 = 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 += 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计算卷积神经网络的讲解：
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个。
分清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，无需保存/读取上下文。
总结几个重点地方：
核函数的使用方法：
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.
再看一个稍微复杂的例子，
两个矩阵的乘法：
直观一点：
使用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*b;
}
c = t;
}
}
}
如何用GPU进行矩阵运算？让人一想就能想到的算法是，既然有block，既然有thread，那么每个block负责1行的计算，每个thread负责几列的计算应该是比较好的。
那么怎么实现kernel函数？
这里需要考虑一个显卡内存（Device Memory）读取速度的问题，Device Memory有一个特点，连续读取的速度远远高于随机读取的速度，那什么叫连续读取，这里就涉及到线程的调度问题。单个线程连续读取一块内存算是连续读取吗？错了，在GPU执行时，是一个线程读取完一组数据后直接调用下一个线程读取，而非一个线程连续读取，所以按照线程号排列的索引才是连续的索引。具体操作看kernel函数：
yyy71cj 发表于 2019-3-6 13:58
不明觉厉
机器上有NVIDIA显卡没有，装上驱动和cuda就可以开始了。 可以调用接口获取设备信息：
选择在哪个GPU上运行程序：
再说核函数参数问题：
教程很全 想玩GPU编程的可以看一下啊 GPU的cuda编程和C语言很像
页:
[1]