CPU |
您所在的位置：网站首页 › 如何调用cuda › CPU |
CPU
|
这几天收拾准备准备GUP“库达”计算，开学搞搞计算机图形图像处理。安装问题在另一篇博文里有记载。在学习库达前，首先来阐明一些计算机体系结构概念知识，方便入手: CPU体系架构什么是CPU？ 上述的内部工作流程可以简化成一个流水线 极大地减少了时钟周期，但增加了一些芯片面积。 打个比方，工厂进行加工要走一系列流程，如果它的每个细节处理都安排员工来专门做这事，就做的越来越快，减少了制造这个产品的总时间，但是同时工厂也得雇佣更多的员工增加了工厂的规模和开支。 在流水线处理时，要注意依赖关系，并不是流水线越长越好 旁路：相互依赖的指令操作时，可以开一个（旁路）后门不进行上述的写回等工序直接引用到下一个指令， 停滞：流水线的停滞，就拿工厂来说，中间的一个步骤罢工，后面的步骤只能等待。 分支预测不能消极的等待，所以我们做了分支预测，根据停滞前的那条指令来猜测下一条指令。 猜测还是要有一定的策略基础并非完全玄学，考试遇到不会的题，有的人猜对那可能是玄学，但有的人还是根据自己平时的积累进行排除选项提高预测概率。 具体方法给社会贡献了多少博士与论文啊（嘎嘎嘎） 超标量superscalar提升IPC（ CPU 每一时钟周期内所执行的指令多少）,增大流水线宽度。也就如工厂一道工序进行加工的产品数目增加 但同时用的资源也多了，旁路网路增多，寄存器，存储器增多。工序占得地方大了，运载的车也就多了 乱序执行OoO（OUT of Order）我们写的一个程序逻辑上认为计算机也会按照我们的步骤来进行，但是为了充分挖掘计算机内部的吞吐率，CPU进行重排指令。 存储器层次计算机的主要工作占绝大数是进行访问存储器——数据的搬来搬去 。但存储器的访问速度又很慢，所以进行改良： 缓存：利用将数据放在尽可能接近的位置来供下一次更快的使用。（举个难忘的例子：我大学体育选的是乒乓球，到了期末考试，要进行一分钟55次攻防球来回，我准备了俩球，把另一个球放在了裤兜里，重要的是我还把拉链拉上了，再次取 的时候我花了好几秒，导致那次没有达成目标，第二次我直接找旁边的一女生借，虽然事先没有说明，但不得不说快多了！不出意外，有点懵逼）大概就是这么个意思。 有了前面一些简单概念的铺垫，接下来开始正式的cuda介绍。其实，cuda 的产生就是为了给CPU 分担一些计算的任务等，加快程序数据的吞吐率： CUDA编程
CUDA程序构架分为两部分：Host和Device。一般而言，Host指的是CPU，Device指的是GPU。在CUDA程序构架中，主程序还是由CPU来执行，而当遇到数据并行处理的部分，CUDA 就会将程序编译成GPU能执行的程序，并传送到GPU。而这个程序在CUDA里称做核（kernel）。CUDA允许程序员定义称为核的C语言函数，从而扩展了C语言，在调用此类函数时，它将由N个不同的CUDA线程并行执行N次，这与普通的C语言函数只执行一次的方式不同。执行核的每个线程都会被分配一个独特的线程ID，可通过内置的threadIdx变量在内核中访问此ID。在 CUDA 程序中，主程序在调用任何GPU内核之前，必须对核进行执行配置，即确定线程块数和每个线程块中的线程数以及共享内存大小。 1. 1核函数 定义：在GPU进行的函数通常称为核函数一般通过__global__修饰（在核函数里，都用双下划线来修饰），调用通过，第一个参数代表block线程块数目，第二个参数代表线程块内含有的线程数目thread。根据2可以看出它的操作单位是总的线程格的block。调用时必须声明核函数的执行参数。正式编程前必须使用cudaMalloc（）函数先为wrap分配适当的内存空间。分配前要精打细算，因为总的内存空间一定，一个wrap分配不当的空间会导致其他线程内存不够用导致越界报错甚至死机。（字面上 还是蛮难的）。任何程序入门编程： #include #include int main() { printf("Hello world\n"); }乍一看，除了头文件多了一个以外，其他毫无区别。常说cuda编程用的是c语言，难道真的一成不变？ 答：cuda编程面向GPU，但如果程序里并没有用到GPU处理，那就真的是一模一样。但如果上述程序你真的不想让显卡就这么闲着，你可以这样： // 首先，修改文件后缀为.cu #include #include //第一步：核函数申明 __global__void kernel(void)//注意：这里的返回值只能为void { } int main(void) { //核函数的调用在主函数 kernel ()//启用一个block的一个thread printf("Hello,world\n"); return 0; } 1.2 dim3结构类型 dim3是基于unit3定义的矢量类型，表示3个数据成员unsigned int x;unsigned int y;unsigned int z;可用来表示一维、二维、三维线程块。相关内置变量 2.1 gridDim,线程格的维度定义，即gridDim.x,gridDim.y,gridDim.z。 2.2 blockDim，线程块的维度，即同上。 2.3blockIdx,线程块ID的索引，及同上。 2.4threadIdx,线程ID的索引，即同上。一维的block，包含的线程ID，threadId=threadIdx.x，但对于二维三维，存在以下计算公式：threadID=threadIdx+threadIdx.yblockDim.x;threadID=threadIdx.x+threadIdx.yblockDim.x+threadIdx.zblockDim.yblockDim.z。计算线程索引偏移量为启动线程的总数。如stride=blockDim.x*gridDim.x;threadId+=stride 1.3函数修饰符kernel核函数，在主机端调用kernel在设备端创建多个线程。 1.4常用的内存函数cudaMalloc() 用处：用来分配全局存储器（提示：少用，容易沾满） 函数原型：cudaError_t cudaMalloc(void **devPtr,size_t size) 将cudaMalloc（）分配的指针传递给设备和主机上。 cudaMemcpy() 用处：进行CPU与GPU的数据拷贝 cudaMemcpy（Md(gpu),M(原地址i)cudaMemcpyHostToDevice） cudaFree（） 函数原型：cudaError_t cuda Free用处：释放内存与c函数free（）一样。 1.5 GPU内存类型 全局内存global memory，通常意义上的设备内存。 共享内存 shared memory，GPU的每个线程块位于一个相同的核处理器SM，之间共享内存。形式：__shared__float a[i] 常量内存 constant memory,常量内存代替全局内存能有效的减少内存带宽。要求：当我们要将数据拷贝到常量内存时，要使用cudaMemcpyToSymbol（），使用cudaMemcpy仍会拷贝到全局。 纹理内存 Texture memory从硬件角度来说，纹理内存实际上是存储在全局内存上；唯一的区别是，当某一个变量绑定纹理内存后，在程序运行过程中会将部分信息存储在纹理缓存中，以减少线程块对全局内存的读取，进而提高程序的运行速度。使用：有两种方法：纹理参考与纹理对象。一维纹理内存 用texture+类型声明。 cudaBindTexture（）绑定到纹理内存中。 tex1Dfetch()读取纹理内存的数据。 二维纹理内存 texture texIn。 cudaBindTexture2D（）绑定到纹理内存中。 通过tex2D()来读取纹理内存中的数据。 先写这里，补作业了！emmm，我太懒了！ 固定内存将数据保存在主机内存里，确保了该内存始终驻留在物理内存。 好处：如果知道主机内存的物理地址，可以直接通过“直接内存访问”DMA技术来在GPU和主机之间复制数据。由于DMA介入并不需要访问CPU，大大提高了访问速度。这里简单介绍一下DMA 的工作原理：CPU将DMA初始化后，DMA控制器获得权限，CPU即进入后台，由DMA控制器输出命令，直接控制I/O接口进行DMA传输。详细请戳——>百度百科DMA形式：通过cudaHostAlloc（）函数来分配，通过cudaFreeHost（）释放。 1.6 原子性 atomic functions,操作的执行过程分解为更小的部分。函数的调用：atomicAdd(addr，y)生成一个院子的操作序列，这个操作序列包括读取addr(相应地址)出的值,将y增加到该值。结果保存到addr地址。 1.7 流 粒度是计算机领域中粒度指系统内存扩展增量的最小值。CUDA 的线程组成为Grid __syncthreads(); }有的程序执行到是，有的程序执行到否，导致线程并没有达到同步更不可能达到同步，程序就会死锁。打个比方：你约好和班里的几名同学一起吃饭，不等到互相别吃饭，你在楼的西侧等他，他在楼的东侧等你，都没有等到对方，也就意味着在那一直等 ， 死锁。 1.9 加速比对一个程序进行加速的时候，很多时候要预先估算出该程序使用GPU加速后的加速比，尤其做项目的时候！（从大佬那里听来的~） 如何估算程序的GPU加速比？ Amdahl阿姆达尔定律通俗来讲，N个处理器：（（1/（P（并行）/N+S串行））P,S表示并行与串行的计算所占比例P+S=1使尽可能的系统并行化，由公式可得，程序中可并行代码的比例决定你增加处理器（总核心数）所能带来的速度提升的上限。 另外，GPU加速比针对的是每一个步骤。即一个程序M由P、A、B若干步骤组成，对P步骤进行加速，则GPU加速比=P步骤加速前的时间/P加速后的时间，而不是S程序加速前的时间/S加速后的时间这个非常重要，因为很多人不懂GPU加速，他们往往只关注整个程序的加速效果。而对于一个完整的程序，很多都包括一些非常耗时的操作（比如读取图片等）。这种情况下就算你对程序中的某个步骤使用GPU加速到极致，整个程序可能只感觉快了一点点，这是非常不公平的。 估算加速比通常情况下有几种惯用的手段： CPU来实现大的for循环资源的管理（内存、显存的申请与释放等）核函数级别的优化（1）原程序CPU完成，程序并行化程度高，且依赖性低 这种为最理想的情况，常见于简单的图像算法。通俗来讲即for循环，循环的数据之间没有依赖性，在这种情况下，加速比容易破10+，图片越大，加速比也能达到更高。 （2）原程序GPU上完成，资源管理差 程序已经由GPU核函数或者调用GPU API实现，但是资源（显存和句柄的申请和释放等）管理比较差。比如之前在流那节讲到的高效运用多流这种情况下，看资源申请所占的时间比例，不过一般加速比能到几倍就很不错了。 （3）原程序GPU上完成，但核函数实现较差 如果只剩下对核函数进行优化，那么一般可得到的加速比就更低了，大概在2X. 1.10 线程调度面临的现状是我们从软件上启动的线程数目远远大于硬件上的可用的执行单元数目。 可以参考NVIDIA公司对显卡的介绍以及cuda开发手册查看显卡的SM，sp,block数。
通俗来说就是将一组很大的数据通过算法化简成一个个小块，
好了，框架就这基本要素。 就这？就这？之后会出一个CUDA项目实战，再偶尔看看c++，就就…开学了吧。 |
工作流程简化为：取指——>译码——>执行——>访存——>写回
这里可以简单理解为：一张桌子有8张椅子，32个人要同步吃饭。就保证8为周期，进行4次，每一人恰一口。 wrap：针对的是一个block里的一组线程，32个线程为一个单位，有序号排列0-63，两组wrap，如果一个块大小不是32的倍数，最后一个wrap会被填充。 特征：一个SM在调度过程中只允许一个wrap来占用。 注意：wrap里的线程需要天然同步，wrap内部线程沿不同分支执行时，没办法为每一个ALU都设置一个逻辑复杂的机构，整个系统的速度会严重下降！ 利用线程之间互相独立的代码来实现延迟掩藏。
线程之间交叉进行加和，每一轮需要的线程数目减半，但是由于wrap分割的限制（线程同步 ），每一轮闲置下来的线程并不能真正的释放自由。这里改变一下线程偏移量顺序：
右面改变了线程偏移量，由1,2，4——>4,2,1.但从根本上提高了加速比。这里回忆wrap，一个wrap管理8个线程同步，右面的一半线程会提早全部结束，就可以进行内存释放等其他操作，减少了资源占用所以，要避免分支发散。 （我感兴趣的是它在医学CT等图像处理方面应用的，特意跑去知网看了看，也终于第一次看完了几篇完整的论文，里面涵盖的知识点多且片面，大都是对基于系统的介绍，怪不得老师会说这方面的论文层出不穷，但国产一个操作系统的真正实现还只是星星点点）
【本文地址】
今日新闻 |
推荐新闻 |