本帖最后由 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);
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, 下载次数: 7)
CUDA中文手册.pdf
(3.41 MB, 下载次数: 6)
CUDA编程模型.pdf
(194.46 KB, 下载次数: 8)
| 
楼主
