本帖最后由 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)
|
@keer_zu :Hi~图章今年规定仅推荐时长为一周左右的新文章,所以有撤销哦,实在抱歉。
@21ic小管家 为什么要撤销图章?这个是GPU编程入门很好的教程啊。