首页 > 代码库 > [CUDA] CUDA to DL
[CUDA] CUDA to DL
又是一枚祖国的骚年,阅览做做笔记:http://www.cnblogs.com/neopenx/p/4643705.html
这里只是一些基础知识。帮助理解DL tool的实现。
“这也是深度学习带来的一个全新领域,它要求研究者不仅要理论强,建模强,程序设计能力也要过硬,不能纸上谈兵。”
- CUDA的广泛应用造就了GPU计算专用Tesla GPU的崛起。
- 随着显卡的发展,GPU越来越强大,而且GPU为显示图像做了优化。在计算上已经超越了通用的CPU。如此强大的芯片如果只是作为显卡就太浪费了,因此NVidia推出CUDA,让显卡可以用于图像计算以外的目的。
- 目前只有G80、G92、G94、G96、GT200、GF100、GF110、GK100、GK104、GK107平台(即GeForce 8~Gecorce GTX780Ti)的NVidia显卡才能使用CUDA,工具集的核心是一个C语言编译器。G80中拥有128个单独的ALU(Arithmetic Logic Unit,算术逻辑单元),因此非常适合并行计算,而且数值计算的速度远远优于CPU。
Tesla GPU
Ref: http://bbs.gpuworld.cn/forum.php?mod=viewthread&tid=199
A. GeForce系列GPU追求“速度”,并不会对“数据正确性”进行“再确认”,因为“显示”就算有1%~2%的错误,也无伤大雅,反正刷屏速度60Hz,肉眼也分辨不出,错就错了也无所谓
B. 但对于“运算”需求,是容不得丝毫错误的,必须达到99.999999999....%正确率的,就有非常高的要求,这也是Tesla GPU必须采用ECC显存,以确保运算正确率的原因,导致价格比GeForce高出很多。
如果,你对“数据正确性”要求不高,那 GeForce 卡绝对会让你很开心。但如果你的计算结果容不下一点点错误,那还是得咬着牙选择 Tesla 专业卡,否则你还要花更多成本去面对难以预期的风险。
最后的caffe性能对比,自觉脑补。
OpenCL
OpenCL(全称Open Computing Language,开放运算语言)是第一个面向异构系统通用目的并行编程的开放式、免费标准,也是一个统一的编程环境,便于软件开发人员为高性能计算服务器、桌面计算系统、手持设备编写高效轻便的代码,而且广泛适用于多核心处理器(CPU)、图形处理器(GPU)、Cell类型架构以及数字信号处理器(DSP)等其他并行处理器,在游戏、娱乐、科研、医疗等各种领域都有广阔的发展前景。
CUDA到底是个什么?
为了泛型编程(C、C++、Fortran多语言)、以及榨取更多的计算力,NVIDIA对OpenCL进行的改装,贴合自己的GPU硬件架构,量身定做出CUDA。
较游戏程序员不同,CUDA程序员主要工作,就是把握硬件架构,在算法理论时间复杂度下,将算法串行执行体系,改组为并行执行体系。(能并行的并行化)
NVIDIA为它在不同成长阶段卖出的产品,规定了计算能力体系:
- 计算能力1.0是跑CUDA的最低条件,这一时期的代表作是8800GT家族。
- Fermi架构的计算能力是2.0,
- Kepler是3.0,
- Maxwell是4.0。
GPU会按照负载均衡的原则,将任务平均分至各个SM阵列 <---- Stream Multiprocessors(流多处理器),民间多译为SM计算阵列
对于每个SM阵列,就调度它手下那一伙CUDA核心干活。流处理器(SP)改名为"CUDA核心"。
2.
由于每个SM阵列的CUDA核心有限,NVIDIA规定,Fermi架构,每个SM最多并行执行1024个线程。
当然,实际任务中,每个SM会分到几百万个线程,这时候,就只能小部分并行,然后再串行了。
- Fermi 1.0架构,官方设计是16组SM,512SP,然而旗舰GTX480最后只弄出了15组,480SP,顺次阉割出GTX470、GTX460。
- Fermi 2.0架构,旗舰GTX580,总算达到设计图要求,达到了16组,512SP,顺次阉割出了GTX570,GTX560。
3.
Kepler架构最大变化在于, 对每个SM阵列,将SP数量扩大到6倍,达到192SP。谓之曰SMX阵列。
每个SMX阵列,包含192个CUDA核心,单次并行吞吐量是2048个线程。
- Kepler 1.0架构,官方设计是15组SM,2880SP,然而旗舰GTX580最后只弄出了8组,1536SP,顺次阉割出GTX570、GTX560。
- Kepler 2.0架构,旗舰GTX680,总算达到设计图要求,达到了15组,2880SP,顺次阉割出了GTX670,GTX660。
特别版,GTX Titan Z,直接把两块GK110并在一起,合出了30组,5760SP,同时支持双精度浮点计算。
其阉割掉双精度之后,就是GTX690。
值得一提的是,GTX游戏卡直接把双精度阉割掉了,因为只有Tesla做科学计算的时候,才会用双精度浮点运算。
4.
Maxwell架构是老黄的无奈之举。因为台积电把20nm工艺让给了ARM系(Apple和高通)。
还是基于28nm的Maxwell,继续在SM上大刀阔斧闹改革,将192SP降低为128SP,谓之曰SMM阵列。
Maxwell 最新架构,官方设计是16组SM,2048SP,为旗舰GTX980,顺次阉割出GTX970,GTX960。
特别版,GTX TitanX,24组SM,3072SP,较之TitanZ,阉掉了双精度浮点数支持。
TitanX是老黄在GTC 2015向DL界主推的一块民用卡,因为DL无需高精度浮点,用Tesla太奢侈。
#include "cuda_runtime.h"#include "device_launch_parameters.h"#include<stdio.h>#include<stdlib.h>#include<string.h>
int main(){ int deviceCount; cudaGetDeviceCount(&deviceCount);
int dev; for (dev = 0; dev < deviceCount; dev++) { cudaDeviceProp deviceProp; cudaGetDeviceProperties(&deviceProp, dev); if (dev == 0) { if (/*deviceProp.major==9999 && */deviceProp.minor = 9999&&deviceProp.major==9999) printf("\n"); } printf("\nDevice%d:\"%s\"\n", dev, deviceProp.name); printf("Total amount of global memory %u bytes\n", deviceProp.totalGlobalMem); printf("Number of mltiprocessors %d\n", deviceProp.multiProcessorCount); printf("Total amount of constant memory: %u bytes\n", deviceProp.totalConstMem); printf("Total amount of shared memory per block %u bytes\n", deviceProp.sharedMemPerBlock); printf("Total number of registers available per block: %d\n", deviceProp.regsPerBlock); printf("Warp size %d\n", deviceProp.warpSize); printf("Maximum number of threada per block: %d\n", deviceProp.maxThreadsPerBlock); printf("Maximum sizes of each dimension of a block: %d x %d x %d\n", deviceProp.maxThreadsDim[0], deviceProp.maxThreadsDim[1], deviceProp.maxThreadsDim[2]); printf("Maximum size of each dimension of a grid: %d x %d x %d\n", deviceProp.maxGridSize[0], deviceProp.maxGridSize[1], deviceProp.maxGridSize[2]); printf("Maximum memory pitch : %u bytes\n", deviceProp.memPitch); printf("Texture alignmemt %u bytes\n", deviceProp.texturePitchAlignment); printf("Clock rate %.2f GHz\n", deviceProp.clockRate*1e-6f); } printf("\nTest PASSED\n"); getchar();}
Comment:
deviceProp.name为GPU名字,如果没有GPU则会输出 Device Emulation
deviceProp.totalGlobalMem返回的是全局储存器的大小,对大数据或一些大模型计算时显存大小必须大于数据大小,如图返回的是2GB的存储大小,
deviceProp.multiProcessorCount返回的是设备中流多处理器(SM)的个数,流处理器(SP)的个数SM数×每个SM包含的SP数,其中帕斯卡为每个SM,64个SP,麦克斯韦为128个,开普勒为192个,费米为32个,
deviceProp.totalConstMem返回的是常数储存器的大小,如同为64kB
deviceProp.sharedMemPerBlock返回共享储存器的大小,共享存储器速度比全局储存器快,
deviceProp.regsPerBlock返回寄存器的数目;
deviceProp.warpSize返回线程束中线程多少;
deviceProp.maxThreadsPerBlock返回一个block中最多可以有的线程数;
deviceProp.maxThreadsDim[]返回block内3维度中各维度的最大值
deviceProp.maxGridSize[]返回Grid内三维度中各维度的最大值;
deviceProp.memPitch返回对显存访问时对齐时的pitch的最大值;
deviceProp.texturePitchAlignment返回对纹理单元访问时对其参数的最大值;
deviceProp.clockRate返回显存的频率;
附另一个可能是不错的链接,关于cuda programming:http://blog.csdn.net/augusdi/article/details/12833235
只看一些基础概念:
2.1 线程网格(Grid)、线程块(Block)、线程(Thread)、线程束(Warp)
2.1.1 内核函数
内核函数是并行计算中最基本的单元函数,其特点是:
统一的处理逻辑代码,分布并行掌控不同区域的数据,以此达到多区域数据联动并行执行。
NVIDIA为了CPU在逻辑上能调度GPU计算的函数,规定了统一的格式。
以__global__限定符为始,声明:__global__ void helloworld()。
__global__意思为,GPU执行,CPU调用
调用时,需要分配 <<<线程块,块内线程数>>>。
如执行helloword,使用1个线程块,块内使用256个线程,则
helloworld<<<1,256>>>
2.1.2 线程网格(Grid)
线程网格在编程时并不存在,它只是抽象上的并行网格体系。
不同种类的内核函数,每种内核函数调度数个的线程块,这数个线程块逻辑上被判为一个Grid。
2.1.3 线程块(Block)
线程块是一个3D结构,强调3D坐标系时,需要以dim3类型声明三维大小。
dim3是个结构体, 成员x、y、z,代表方向轴尺度。
如helloworld<<<dim3(1,1,1), 256>>>。
当然,大部分操作基本使用的是1D坐标系,线程块默认全部扩展到X轴上。
一般写成helloworld<<<1,256>>>。
通常在内核函数内,需要获取线程块编号,以便对数据集的不同区域处理,四大重要属性:
?dim3 gridDim(不是指有多少Grid,而是指一个Grid有多少Block)
?dim3 blockDim(不是指有多少Block,而是指一个Block有多少Thread)
?dim3 blockIdx
?dim3 threadIdx
对于1D坐标系,有int tid= (blockDim.x*blockIdx.x) + threadIdx.x;
tid指明当前线程的编号,是内核函数里最基本的控制变量。
int step=(blockDim.x * gridDim.x);
由于CUDA限制每个Block的线程数(2.0以上通常使用1024,以下通常使用512)
所以在常规元素分解模型中,通常把每个Block的线程数设置为常量(固定不动)
这时,有两个策略:
① 其一,不固定Block:
这种方法最为常用,由于CUDA对每个任务而言,对Block数量的限制很松,
如图:
这时候,可以采取为每个线程分配一个元素的方法,用
BLOCKS=(N+THREADS−1)/THREADSBLOCKS=(N+THREADS−1)/THREADS
算出一个动态的Block数量的需求,这时候,for(i=tid;i<N;i+=step)等效于for(i=tid;i<N;i+=1)
因为这个循环根本不会执行第二次。
① 其二,固定Block数量:
这时,这时候,为了跑完全部的N个元素,有些Thread会启动人工循环。
i+=step会将元素坐标继续跳转,因为N必然大于step,你不能用+1来取剩余的元素吧?
这两种方法本质上是等效的,由于在物理执行时,同时并行线程最多大概是3072,
几百万、甚至几千万的Block会被CUDA扔到等待队列里,由CUDA自己安排自动循环,
有时甚至比你的人工循环更高效,所以,通常用①的方法,为了保持写法一致,step也会作为默认跳转量。
2.1.4 线程(Thread)
CUDA逻辑体系里最基本的执行单位,等效于CPU的线程。
内核函数一旦被指明了线程块大小,线程大小后,每个线程就分配到了一个内核函数的副本。
区别这些的线程的唯一方法就是线程编号tid,通过tid,让不同线程窥视数据集的不同部分。
用相同的逻辑代码,执行数据空间的不同子集。
2.1.5 线程束(Warp)
线程束对用户透明,它是NVIDIA强行规定的。目前显卡都固定为32。
逻辑上,线程束将32个线程编为一组。
一般微机系统,如8086,它的访存模式是串行的。每一个总线周期,吞一个字节进来。
NVIDIA的GPU在一个总线周期内,能够最大吞32*4=128字节。
前提是当个线程束内的线程,逐序访问显存,这特别需要设计数据存储形式。
使用线程束的目的是掩盖单个总线周期过长的问题,通常要跑500~600个T周期。
一般来说,一个CUDA程序必然少不了以下三步:
?cudaMalloc:创建新的动态显存堆
?cudaMemcpy:将主机(Host)内存复制到设备(Device)显存
?显存处理完之后,cudaMemcpy:设备(Device)显存复制回主机(Host)内存,释放显存cudaFree
其中第三步最容易遗忘。要知道,CPU最后是无法使用显存中的数据的。
一个例子:
/* GPU版HelloWorld,主要目的是演示CUDA基本程序框架: *? 将HelloWorld复制进显存 *? 让GPU完成strcpy函数 *? 将显存中的HelloWorld转回内存,并且打印 *//****kernel.cu****/#include "cuda_runtime.h"#include "device_launch_parameters.h"__global__ void cudaStrcpy(char *des, char *src) /*内核函数*/{ while ((*src) != ‘\0‘) *des++ = *src++; *des = ‘\0‘;}/****gpu_helloworld.cu****/#include "device.cu"#include "kernel.cu"#include "cstring"void helloworld(char *str1, char *str2){ InitCUDA(); char *dev_str1=0, *dev_str2=0; int size = strlen(str1) + 1; cudaMalloc((void**)&dev_str1, size); /*cuda系函数必须放在cu文件里*/ cudaMalloc((void**)&dev_str2, size); cudaMemcpy(dev_str1, str1, size,cudaMemcpyHostToDevice); cudaStrcpy<<<1,1>>>(dev_str2, dev_str1); /*单线程块、单线程*/ cudaMemcpy(str2, dev_str1, size, cudaMemcpyDeviceToHost);}/****main.cpp****/#include "cstdio"#include "cstring"extern void helloworld(char *str1, char *str2);int main(){ char src[] = "HelloWorld with CUDA"; char *des = new char[strlen(src)+1]; helloworld(src, des); printf("%s\n", des);}
另一个例子:
/* 向量加法是CUDA 7.0在VS中提供的样例模板,演示了并行算法的经典trick:循环消除。 *利用单个线程块中,多个线程并发执行,来消除循环。 *时间复杂度估计,不能简单从O(n)迁移到O(1),因为GPU同时并行量存在限制。 *即便是Kepler架构中拥有192SP的SM阵列,理论同时并行量也不过是2048。 *//****kernel.cu****/__global__ void kernel_plus(int *a, int *b, int *c){ int x = threadIdx.x; c[x] = a[x] + b[x];}/****gpu_vectoradd.cu****/void vectorAdd(int *a, int *b, int *c,int size){ if (!InitCUDA()) return; int *dev_a = 0, *dev_b = 0, *dev_c = 0; cudaMalloc((void**)&dev_a, size*sizeof(int)); cudaMalloc((void**)&dev_b, size*sizeof(int)); cudaMalloc((void**)&dev_c, size*sizeof(int)); cudaMemcpy(dev_a, a, size*sizeof(int), cudaMemcpyHostToDevice); cudaMemcpy(dev_b, b, size*sizeof(int), cudaMemcpyHostToDevice); kernel_plus << <1, size >> >(dev_a, dev_b, dev_c); cudaMemcpy(c, dev_c, size*sizeof(int), cudaMemcpyDeviceToHost);}/****main.cpp****/extern void vectorAdd(int *a, int *b, int *c, int size);int main(){ int a[5] = { 1, 2, 3, 4, 5 }, b[5] = { 10, 20, 30, 40, 50 }, c[5] = { 0 }; vectorAdd(a, b, c,5); printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n", c[0], c[1], c[2], c[3], c[4]);}
Ref: http://blog.csdn.net/qiexingqieying/article/details/51734347
首先对这些框架进行总览。
库名称 | 开发语言 | 速度 | 灵活性 | 文档 | 适合模型 | 平台 | 上手难易 |
Caffe | c++/cuda | 快 | 一般 | 全面 | CNN | 所有系统 | 中等 |
TensorFlow | c++/cuda/Python | 中等 | 好 | 中等 | CNN/RNN | Linux, OSX | 难 |
MXNet | c++/cuda | 快 | 好 | 全面 | CNN | 所有系统 | 中等 |
Torch | c/lua/cuda | 快 | 好 | 全面 | CNN/RNN | Linux, OSX | 中等 |
Theano | python/c++/cuda | 中等 | 好 | 中等 | CNN/RNN | Linux, OSX | 易 |
接下来将对这些框架进行分别介绍。
Caffe
TensorFlow
MXNet
[CUDA] CUDA to DL