深度学习（三十六）异构计算CUDA学习笔记（1）

异构计算CUDA学习笔记（1）

原文地址：http://blog.csdn.net/hjimce/article/details/51506207

作者：hjimce

近日因为感觉自己在深度学习工程化之路比较薄弱，故此开始学习CUDA编程，弥补自己在这方面的缺陷。写笔记以记录自己对cuda编程的一些简单理解。

个人感觉学习CUDA，最重要两点是：

1、理解Grid、Block、线程之间的层次关系；

2、理解存储器的层次关系（共享存储器、显卡等）。

一、硬件知识：

GPU硬件知识：Grid网格(每个显卡的个数)-》block-》线程

1、Grid：一个Grid代表一块GPU芯片，所有的线程共享显存数据；每个grid就相当于一块显卡。

2、Block：在每一个GPU芯片里面包含着多个block,每个block包含了512或者1024个线程。

每个线程的ID号可以通过一维0~1024索引，也可以通过二维dx*dy=1024索引，或者通过三维dx*dy*dz=1024。这个就像图像opencv访问某个像素点一样，可以通过一维访问、或者二维访问：i+width*j。

每个块里各自有一个共享数据存储的区域，只有块内的线程可以访问；在一个块内，共享变量的修改，可能需要用到等待所有的线程处理完毕，然后再修改共享变量，可以采用syncthreads（）函数用于等待。

3、Thread：每个block包含多个thread。

GPU存储空间：

(1)block中的每个线程都有自己的寄存器和local memory；

(2)block中的所有线程共享一个shared memory；

(3)一个grid共享一个global memory（或者称之为显存）、常量存储器、纹理存储器。

根据这些存储器的不同，我们后面定义的变量的时候，也要使用限定符，告诉程序，我们所要定义的变量是位于那个存储器，具体后面再解释。

二、CUDA编程步骤：

1、设置显卡编号：cudaSetDevice；

2、为显卡开辟变量内存:cudaMalloc；

3、把cup上的数据拷贝到GPU上：cudaMemcpy；

4、调用内核函数__global__类型函数；

5、把计算结果拷贝到CPU上：cudaMemcpy；

6、释放显存空间cudaFree；

示例代码：

//计算a、b相加，得到c，size输入向量的维度
cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size)
{int *dev_a = 0;int *dev_b = 0;int *dev_c = 0;cudaError_t cudaStatus;//选择显卡cudaStatus = cudaSetDevice(0);if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");goto Error;}// 在显存上，开辟空间，存储变量ccudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaMalloc failed!");goto Error;}
// 在显存上，开辟空间，存储变量acudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaMalloc failed!");goto Error;}
// 在显存上，开辟空间，存储变量bcudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaMalloc failed!");goto Error;}//把数据a、b拷贝到显存上cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaMemcpy failed!");goto Error;}cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaMemcpy failed!");goto Error;}// 设置核函数addKernel<<<1, size>>>(dev_c, dev_a, dev_b);// Check for any errors launching the kernelcudaStatus = cudaGetLastError();if (cudaStatus != cudaSuccess) {fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));goto Error;}// cudaDeviceSynchronize waits for the kernel to finish, and returns// any errors encountered during the launch.cudaStatus = cudaDeviceSynchronize();if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);goto Error;}//把计算结果拷贝到cpucudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaMemcpy failed!");goto Error;}Error:cudaFree(dev_c);//释放空间cudaFree(dev_a);cudaFree(dev_b);return cudaStatus;
}

三、CUDA内置函数库

1、Cuda的函数命名规则：以cuda开头+该函数功能名，函数功能名每个单词的第一个字母都是大写，比如：cudaSetDevice，也就是cuda+Set+Device。一些常用的函数功能名基本上都是和c语言一样，比如：cudaMalloc、cudaMemcpy。

2、GPU称之为设备device；device(0)表示设置显卡号码，多显卡，在CUDA程序中，我们可以采用:cudaSetDevice(0)函数，表示选用第一块显卡进行计算。

3、CPU称之为主机host；

所以cuda在定义函数、变量的时候，前面会有个限定词：host、device、global，三者分别表示定义的函数：在cpu调用执行、在gpu调用执行、cpu调用gpu执行。

四、CUDA自定义函数

因为程序是在gpu、cpu不同的设备上混合使用的，所以在自定义函数的时候，需要加入函数限定词__device__ , __global__, __host__；这些是用来告诉程序，你定义的这个函数是要在GPU上调用执行，还是要在CPU上调用执行。

(1)__device__：表示从GPU上调用，在GPU上执行；

(2)__global__：表示在CPU上调用，在GPU上执行，也就是所谓的内核(kernel)函数；内核主要用来执行多线程调用。

(3)__host__：表明在CPU上调用，在CPU上执行，这是默认时的情况，也就是传统的C函数。

示例：

__global__ void addKernel(int *c, const int *a, const int *b)
{
}

定义了一个函数名为addkernel的函数，该函数限定词为global，表示该函数由cpu调用，由GPU执行。

四、重要类型变量

1、dim3

这个类型是线程索引必备数据结构，是一个向量类，类似于opencv的vec3；

在定义dim3变量的时候，为指定的分量都自动被初始化为1，dim3类型的变量定义完毕后，我们可以采用.x,.y.z访问每个维度的数值。

例子：

dim3 bb(10,20);
std::cout <<bb.x<<","<<bb.y<<","<<bb.z<<std::endl;

这个时候可以看到输出bb.z=1；

切记上面会自动为未被分配的分量初始化为1。这个对于后面<<< >>>的配置，理解非常重要，因为<<< …… >>>的前两个输入参数是dim3类型，如果你输入参数是一个数值a，那么会被自动转换成dim(a,1,1)向量。

2、cudaError、cudaDeviceProp等程序信息

除了int、float等还有，cuda内部定义的一些结构体：

cudaError#让我们获得相关信息；

cudaDeviceProp#获得设备的相关参数，比如GPU线程个数、显存大小；

3、变量类型

int3、int2……等可以用于定义一个变量是3维、2维整型向量，当然还有其他float向量等

五、配置内核函数

内核函数的调用格式：函数名<<<Dg,Db,Ns,s>>>(函数参数);

内核函数的输入参数就是我们开辟显存，然后从CPU上把数据拷贝到GPU的变量

调用内核函数都需要配置参数，<<<dim3 Dg,dim3 Db,size_t Ns,cudaStream_t s>>>内核函数配置参数：

Dg、Db都是dim3类型，可能我们在使用的时候直接输入数值<<<2，50>>>，这样其实系统会自动进行类型转换，把2转换成dim3 Dg(2,1,1)。

示例：

<<<2，50>>>，表示采用2个block，每个block启用50个线程进行计算，这样算下来一共有100个线程；

<<<1, size>>>表示运行时配置符号，里面1表示只分配一个线程组（又称线程块、Block）、size表示每个线程组有size个线程。size个线程都会调用这个核函数，我们可以根据threadid获取当前调用该函数的线程。这样设置参数表示我们只用了一组线程块，同时调用了该线程块的size个线程。

<<<size，1>>>，那么就会有size个线程块，每个线程块只启用了一个线程计算。

六、CUDA内置变量

Cuda内置为我们定义了几个经常用到的变量，这些变量基本都是dim3类型。

1、gridDim：利用这个变量，我们可以通过gridDim.x、gridDim.y、gridDim.z，知道网格三个维度的尺寸，这个变量除非是GPU集群，单显卡这种肯定用不到，忽略。

blockDim：表示块在三个方向的长宽高

上面gridDim、blockDim其实等于在我们调用内核函数的时候，配置<<< >>>所需要的参数

2、blockldx、threadldx:分别用于索引当前线程所在的块，块中的线程编号

七、CUDA变量自定义

我们在定义变量的时候，在程序中，如果不声明限定符的话，那么默认都是定义在cpu内存上。现在gpu还有显存、共享存储器、缓存等，所有在定义变量的时候，要声明限定符，告诉程序我们所定义的变量要存在哪里。

1、__device__：表明声明的数据存放在显存中，所有的线程都可以访问

2、__shared__：表示数据存放在共享存储器在，只有在所在的块内的线程可以访问，其它块内的线程不能访问

3、__constant__：表明数据存放在常量存储器中，可以被所有的线程访问，也可以被主机通过运行时库访问

参考文献：

1、《NVIDIA CUDA计算统一设备架构》

********************转载请保留原文地址、作者信息**************************

深度学习（三十六）异构计算CUDA学习笔记（1）相关推荐

36.深度解密三十六：网络舆情监测之另类“免费监测”方法详解
网络营销推广技术.技巧深度解密(三十六)指南: 1.本文档适合零基础以及互联网营销推广工作者,主要讲解免费网络舆情监测的问题. 2.原创版权文档,任何抄袭或者全部.部分模仿都是侵权行为. 3.敬畏法律 ...
推荐系统遇上深度学习(三十六)--Learning and Transferring IDs Representation in E-commerce...
本文介绍的文章题目为<Learning and Transferring IDs Representation in E-commerce>,下载地址为:https://arxiv.org ...
Java多线程学习三十六：主内存和工作内存的关系
CPU 有多级缓存,导致读的数据过期由于 CPU 的处理速度很快,相比之下,内存的速度就显得很慢,所以为了提高 CPU 的整体运行效率,减少空闲时间,在 CPU 和内存之间会有 cache 层,也就 ...
torch学习 (三十四)：迁移学习之微调
文章目录引入 1 微调 2 热狗识别 2.1 数据集载入 2.2 数据集预处理 2.3 定义和初始化模型 2.4 微调模型致谢引入场景: 从图像中识别出不同种类的椅子,然后将购买链接推 ...
JavaWeb学习 (二十六)————监听器(Listener)学习(二)
一.监听域对象中属性的变更的监听器域对象中属性的变更的事件监听器就是用来监听 ServletContext, HttpSession, HttpServletRequest 这三个对象中的属性变更信 ...
OpenCV学习笔记（三十六）——Kalman滤波做运动目标跟踪 OpenCV学习笔记（三十七）——实用函数、系统函数、宏core OpenCV学习笔记（三十八）——显示当前FPS OpenC
OpenCV学习笔记(三十六)--Kalman滤波做运动目标跟踪 kalman滤波大家都很熟悉,其基本思想就是先不考虑输入信号和观测噪声的影响,得到状态变量和输出信号的估计值,再用输出信号的估计误差加 ...
知识图谱论文阅读（八）【转】推荐系统遇上深度学习(二十六)--知识图谱与推荐系统结合之DKN模型原理及实现
学习的博客: 推荐系统遇上深度学习(二十六)–知识图谱与推荐系统结合之DKN模型原理及实现知识图谱特征学习的模型分类汇总知识图谱嵌入(KGE):方法和应用的综述论文: Knowledge Gra ...
推荐系统遇上深度学习(三十九)-推荐系统中召回策略演进！
推荐系统中的核心是从海量的商品库挑选合适商品最终展示给用户.由于商品库数量巨大,因此常见的推荐系统一般分为两个阶段,即召回阶段和排序阶段.召回阶段主要是从全量的商品库中得到用户可能感兴趣的一小部分候选 ...
花书+吴恩达深度学习（十六）序列模型之双向循环网络 BRNN 和深度循环网络 Deep RNN
目录 0. 前言 1. 双向循环网络 BRNN(Bidirectional RNN) 2. 深度循环网络 Deep RNN 如果这篇文章对你有一点小小的帮助,请给个关注,点个赞喔~我会非常开心的~ 花 ...

深度学习（三十六）异构计算CUDA学习笔记（1）

深度学习（三十六）异构计算CUDA学习笔记（1）相关推荐

最新文章

热门文章