CUDA精进之路（零）：CUDA开篇

前言

着手机器视觉项目时接触到了并行编程这一概念，那时候的目的是为了在图像识别的时候通过多个线程同时对多张传入的图片进行并行处理以达到加速程序运行速度，运用的方法主要是利用了C++自带的future库以及thread库，具体实现方法可参照此博客C++11 并发指南四
此方法简单来说是利用了目前计算机为多核式结构（四核），因为在一般的串行编程中，一个程序的指令在单一的 CPU 上按照先后顺序依次执行，而并行编程则将一个程序分成独立的若干部分在一个或多个 CPU 上进行同步执行以取得更高的运算效率和性能。在我自己的项目中我将图像检测的算法写入一个函数中（算法实现较为简单），通过上述方法，同时对3张传入的待检测图进行识别（摄像机拍摄到的蚕茧图片），部分代码如下

Mat img1 = imread("new.bmp");//读入第一张图片
Mat img2 = imread("new1.bmp");//读入第二张图片
Mat img3 = imread("new2.bmp");//读入第三张图片
Mat imgResult1, imgResult2, imgResult3;//创建识别结果
//多线程运算
auto f1 = async(Pod, img1);//Pod为检测算法函数
auto f2 = async(Pod, img2);
auto f3 = async(Pod, img3);imgResult1 = f1.get();
imgResult2 = f2.get();
imgResult3 = f3.get();

但由于CPU加速效果的有限，在处理一些大数据的时候，这种并行的方式并不能从本质上解决程度处理速度慢的问题，故在5月份有幸接触到了GPU并行加速这一概念，开始学习Nvidia公司推出的基于CUDA的并行程序设计，到目前本人已经完成了一些基于OPENCV的图像处理算法关于CUDA的改编，在加速图像处理程序上有着不错的效果，故打算开此博客，将自己的一些关于CUDA的学习以及实践进行记录，如有不当之处，希望各位大佬们批评指正~

CUDA编程基础

了解电脑的同学都知道现在每台计算机在其硬件构成上都包括最重要的两部分——主机和显卡，而在CUDA的架构中引入了主机端（host）和设备（device）的概念，主机端即为CPU，而设备端则为GPU，所谓的CUDA编程即为利用GPU在大规模并行计算上的执行效率高的优势（一些对显卡要求高的大型游戏本质上来讲就是要求显卡的计算速度快），利用CPU，将一些逻辑复杂性较小的计算部分分配给GPU让其协助完成计算任务。
一个完整的CUDA程序由主机代码和设备代码两部分组成。主机端代码部分在CPU上执行，是普通的C代码；设备端代码部分在GPU上执行，此代码部分在内核kernel上遍写（.cu文件）。关于CUDA中C语言的编写规范以及一些功能网上已有许多健全的资料，在这里不在多作解释，简单利用一个CUDA程序进行作为例程进行实践：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdlib.h>
#include <iostream>
using namespace std;
//将数组A与B中的元素相加存入数组C中
__global__ void arrayAdd(int *A, int *B, int *C)
{int i = threadIdx.x;C[i] = A[i] + B[i];
}int main()
{//简单的示例，A、B、C均为长度为4的整型数组//定义主机端数组A、B、Cint A[4] = { 1, 2, 3, 4 };int B[4] = { 2, 3, 4, 5 };int C[4];//定义设备端数组d_a、d_b、d_cint *d_a;int *d_b;int *d_c;//定义他们在设备端的空间大小int size = 4 * sizeof(int);//在GPU中为他们开辟分配对应的显存空间cudaMalloc((void**)&d_a, size);cudaMalloc((void**)&d_b, size);cudaMalloc((void**)&d_c, size);利用cudaMemcpy函数在CPU端A,B的值复制到对应的GPU内存中cudaMemcpy(d_a, A, size, cudaMemcpyHostToDevice);cudaMemcpy(d_b, B, size, cudaMemcpyHostToDevice);//调用编写好的内核程序，实现数组相加功能arrayAdd << <1, 4 >> >(d_a, d_b, d_c);//利用cudaMemcpy函数将GPU端计算结果复制到CPU端cudaMemcpy(C, d_c, size, cudaMemcpyDeviceToHost);//释放所有的GPU内存（切记有借有还，不然后果自负。。。）cudaFree(d_a);cudaFree(d_b);cudaFree(d_c);system("Pause");return 0;
}

CUDA编程简单来说分为3大步，首先创建在GPU中的内存空间，将有初值的变量从CPU端复制值GPU端，接着调用在GPU内运行的内核函数进行计算，最后将内核函数的运算结果复制到CPU端的变量并且释放掉所有之前创建的GPU内存（这步不能漏，不然后面会出一堆内存问题）。
以上面的程序为例，程序目的为利用并行计算将数组A与数组B相加，首先我在主机端随意定义了3个int型数组，其数组长度均为4，接着利用cudaMalloc函数对其GPU内存进行开辟，内存大小为4倍的int型，接着利用cudaMemcpy函数对数据进行传输，这里在函数的末尾“cudaMemcpyHostToDevice”为一个cudaMemcpy函数特有的标识码，即将主机端（Host）的数据传入设备端（Device）。
然后调用了内核函数arrayADD，在解释此内核函数前，首先我将引入CUDA中Grid（网格）、Block（块）、Thread（线程）的概念，每个内核函数在启动前，程序员必须对本次计算的维度进行设计，thread是最小的计算单元，多个thread组成了一个block,而多个block组成了一个grid，这里可能有同学会问，让多线程一起计算直接用相应数量的thread不就行了，为什么还要引入block和grid，增加设计的复杂度，这是因为每台电脑的显卡在硬件设计的时候都存在的一定的物理极限，简单来说一个block中的thread不是无限的，举个例子，本人目前使用的本本中，显卡型号为Nvidia Geforce 755M，它的最大线程数为1024，超过这个值程序就无法运行，故在大计算时需要合理设计block中thread的数量，这在后续的实践中我会继续提到，关于Grid（网格）、Block（块）、Thread（线程）这三者更具体的关系可参见此博客CUDA 的Threading：Block 和Grid 的设定与Warp

由于本例数据较少，在计算任务的分配上直接利用了4个thread和一个block，如此设计在于每个thread承担一次A[i]+B[i]的任务。最后再利用cudaMemcpy函数将运算结果d_c传送给数组C，一个简单的CUDA小程序就完成了，由于本例的目的在于初步尝试CUDA编程，一些细节小问题望大家不要深究~

CUDA中线程索引threadIdx和块索引blockIdx
在上文的内核函数中，第一行中threadIdx.x为线程的索引，它的目的在于让GPU可以准确的找到目的操作数用对应的线程进行计算工作，是CUDA编程中和重要的一部分。
本人在学习CUDA的过程中，一开始经常被线程索引threadIdx和块索引blockIdx搞的崩溃，现将我在学习中总结一些的小经验进行分享：
（1）CUDA中kernel内的变量和常量
在CUDA中每个块中线程的索引变量threadIdx分为X方向上的索引threadIdx.x以及Y方向上的索引threadIdx.y，块索引变量blockIdx分为X方向上的索引blockIdx.x以及Y方向上的索引block.y，通过这两个索引我们可以准确定位到某个线程。
而kernel中的常量为blockDim以及gridDim，它们分别代表了没block的大小以及grid中block的分布。
（2）创建内核程序中block和thread

dim3 blocksPerGrid(N1,M1);//在grid中有N1*M1个block
dim3 threadsPerBlock(N2,M2);//每个block中有N2*M2个thread

（3）创建索引

int threadIndex = threadIdx.x + threadIdx.y * blockDim.x;
int blockIndex = blockIdx.x + blockIdx.y * gridDim.x;
int index = threadIndex + blockIndex * blockDim.x * blockDim.y;

起航

threadIndex为单个block中thread的索引，blockIndex为block在整个grid中的索引，index即为某个thread在整个grid中的索引。

关于CUDA图像处理加速正式开始~

CUDA精进之路（零）：CUDA开篇相关推荐

CUDA精进之路（五）：图像处理——OTSU二值算法（最大类间方差法、大津法）
引言最近在做医疗设备相关的项目,故在项目中大量用到了各类图像分割的算法,为了在图像中分割出特定目标,用到的算法可以有很多,比如阈值分割,多通道分割,边缘分割以及一些前沿的组合分割.而对大多数图像来说 ...
CUDA精进之路（四）：图像处理——Sobel算子边缘检测
引言关于图像边缘检测,记得刚开始接触图像处理时,第一个自己实现的程序是通过笔记本摄像头采集图像,利用OpenCV自带的算法库进行Canny算子边缘检测,那时候当看到程序运行后,视频窗口实时显示经Ca ...
CUDA精进之路（三）：图像处理——图像灰度化、灰度直方图统计
引言在大部分的图像处理程序中,其中必不可少的一步就是对传入的彩图进行灰度处理,将三个通道的RGB图片转化为单通道的Gray图,而对于灰度图进行直方图统计同样是观察检测图像特征的常用方法.在OpenC ...
CUDA精进之路（一）：图像处理——大图像分块处理（包括求均值、最大值）
引言在我的第一篇文章中我简单介绍了CUDA以及我的一些个人学习见解,在本文中我将开始正式开始CUDA实践之旅,众做周知CUDA目前应用的领域十分广泛,它能把一些普通的CPU代码提速几十倍甚至几百倍. ...
CUDA精进之路（二）：图像处理——形态学滤波（膨胀、腐蚀、开闭运算）
引言从这篇文章起,开始将一些较为典型的OpenCV算法通过CUDA进行实现,本文实现的为图像处理中最为常见的形态学腐蚀以及膨胀,由于本文目的在于算法移植后的验证,故在图片的选择上用小图像作为输入的示 ...
CUDA 编程上手指南：CUDA C 编程及 GPU 基本知识
作者丨科技猛兽编辑丨极市平台本文原创首发于极市平台,转载请获得授权并标明出处. 推荐大家关注极市平台公众号,每天都会更新最新的计算机视觉论文解读.综述盘点.调参攻略.面试经验等干货~ 目录 1 C ...
《Go语言精进之路，从新手到高手的编程思想、方法和技巧1》读书笔记和分享
Go语言精进之路,从新手到高手的编程思想.方法和技巧读书分享 1 本书定位 2 本书内容总览 3 选择本书的原因 4 小收获分享第7-12条真的不知道咋命名第13-19条能用--怎么用更好 ...
【CUDA 基础】3.1 CUDA执行模型概述
title: [CUDA 基础]3.1 CUDA执行模型概述 categories: CUDA Freshman tags: CUDA SM SIMT SIMD Fermi Kepler toc: t ...
Go语言精进之路：绞尽脑汁，帮你理解方法本质并选择正确的receiver类型
Go语言虽然不支持经典的面向对象语法元素,比如:类.对象.继承等,但Go语言也有方法(method).和函数相比,Go语言中的方法在声明形式上仅仅多了一个参数,Go称之为receiver参数.而rec ...

CUDA精进之路（零）：CUDA开篇

前言

CUDA编程基础

起航

CUDA精进之路（零）：CUDA开篇相关推荐

最新文章

热门文章