注册 登录  
 加关注
   显示下一条  |  关闭
温馨提示!由于新浪微博认证机制调整,您的新浪微博帐号绑定已过期,请重新绑定!立即重新绑定新浪微博》  |  关闭

My Unix World

不要迷恋Unix,Unix只是计算世界很小的一部分!

 
 
 

日志

 
 

【Copy】CUDA编程接口(一)----一十八般武器  

2008-12-14 22:33:29|  分类: L-X |  标签: |举报 |字号 订阅

  下载LOFTER 我的照片书  |
2008-05-25 06:00 作者: OpenHero 开勇 出处: 天极网 责任编辑:>dizzarz
"); document.write("");
");

  子曰:工欲善其事,必先利其器。我们要把显卡作为通用并行处理器来做并行算法处理,就得知道CUDA给 我提供了什么样的接口,就得了解CUDA作为通用高性能计算平台上的一十八般武器。(如果你想自己开发驱动,自己写开发库- -那我不得不佩服你很有时间,想必也不会有很多人想自己在去实现一个CUDA吧,呵呵,虽然实现一个也不是太难)。前面我们讲到了一些简单的CUDA的C 语言扩展的规则,下面就具体来讲解CUDA给我听哦买提供了多少方便的API函数。在开发CUDA的时候,CDUA也给我们提供了一套完整的API函数。 从一开始就在想,怎么把这些枯燥的API函数,或者这CUDA的一十八般武器说得清楚。如果按照中文的翻译的第四章那么讲解,或许晕的人更多,只知道这些 是武器,而不知道什么武器是用来干嘛的。从序言看到,十八般武艺所列兵器大同小异,形式和内容却十分丰富。有长器械,短器械;软器械、双器械;有带钩的、 带刺的、带尖的、带刀的;有明的、暗的;有攻的、防的;有打的、杀的击的;也有射的、挡的。我们来看CUDA的时候,一看到这么多的API函数,先来给他 分一些类,然后才好徐徐道来。

  一、 API总结:

  1.通用的一些接口,前一章节也有提高过:数学函数,时间函数,同步函数,原子操作;

  2.控制Device的函数;就是得到设备信息,管理设备信息的函数。设置那块显卡工作,得到那块显卡的性能。这里有分为driver级别的API和runtime级别的API;有人会问什么是driver级别和runtime级别请看图:

  

  这个图我们在前一章已经看到过了。不会不清楚吧~- -!driver级别的API就是提供驱动级别的API,就像写驱动一样的感觉。Runtime级别的API就是封装了一些Driver级别的API,按 照一些常规的方法封装了一些底层的API。其实这里就像我们平时生活中一样,最开始对汽车不熟悉的时候,买一辆车回来能开就ok了;能用熟悉 Runtime级别的API就行了。慢慢的,感觉汽车自带的音箱不 好,自己就开始买一些原始设备回来改装车;慢慢的感觉整车都有点不爽了,然后慢慢的发现想修改发动机,修改外形……就开始改装车了,这样的工作,就得从 Driver级别开始做了,玩得更高级一些的就自己设计图纸,自己来用一些零部件来组装车了。这就是Driver API和Runtime API的关系。Runtime的在开始的用起来一般都比较方便,慢慢的发现如果高层(high-level)的Runtime API用起来不方便,就用底层(low-level)的Driver API来自己做改装的车……

  3.内存管 理,host 的内存,device的内存,global的内存,constant第,shared的,这里会分出来一章单独讲Texture(纹理)内存的使用,说实 在的Texture也是内存~非要搞那么神秘,没办法,也只好拿出来单讲……PS:内存管理也分为Runtime级别的API接口,和Device级别的 API接口。

  4.程序运行控制:像Stream,Event,Context, Module, Execution control这样的咱都把归类到运行管理的。这里也得分清楚有Runtime级别的,也有Driver级别的。

  5.好了,这里就剩下OpenGL 和Direct3D的接口函数了,其实把,这也是为了方便做图来用的,主要是OpenGL和DX都已经成了图形显示方面的标准,so~显卡也得照顾这两个 东东了,要不然显卡自己画……hoho要是真自己再来实现OpenGL或者DX,CUDA就真的会头大了,hoho~~还好就借用现有的图形显示的程序来 做就行了。牛顿人家都说是站在“牛头人”(巨人)的肩膀上才能看得更远……咱也不要非自己费那么大的经去做一些无用功。想想吃不饱的时候没看到多少人减肥 的,倒是现在吃得好了,减肥的人多了……长胖了去健身房减肥,始终感觉有点怪怪的……拿钱去做无用功,(有用功就是减肥)不过现在也有人在做实验,把多余 的这些减肥的人的能量转化为电力……扯远了……提一下:他也有两个层次的API,有Runtime层次的,也有Driver层次的。

  好了,差不多就这5个部分的API了。下面我们就来个个讲解CUDA的十八般武器。

  二、 API讲解

  想想前面五个部分,下面一部分一部分的讲解:

  1. 一些通用的函数:

  数学函数,前面章节已经讲了,提一下CUDA 2.0里面添加了一两个新的函数。

  时间函数,clock ()这个自己去查C的函数库;还有CUDA提供的几个时间函数,用起来也没什么难的,只是说一些,CUDA提供的几个时间函数在计算Device上的运行时间的时候,和CPU上的时间函数比起来还不是那么的完美的准确,所以在做时间数据的时候,最好多取几次求个平均。

  同步函数前面章节讲到了, __syncthreads()函数,原子操作函数也就和以前的原子操作函数一个道理,也不用多加解释。不过其实我都觉得这__syncthreads()函数都应该归到程序运行控制部分。

  2.Device管理

  Runtime 和Driver层面的API都提供了设备管理的函数,其实两个层面的API提供的功能都差不多,可以在API的说明中查到他的区别。这里需要指出来的是虽然Host主机上的多线程程序是应该可以同时访问同一块显卡(Device) 的,毕竟显卡就是按照PCIE标准插在PCI插槽上的标准设备(也有AGP接口的),这就是一个多线程访问硬件的问题了。So~本来是应该可以的。但是由 于CUDA的设计原因,这里的host上的当个多线程的线程每个线程都要执行 CUDA kernel的时候,就必须执行在多个Device上面。保证每个线程访问的Device不是同一个。多个线程线程A不能分享线程B在Device上创建 的资源。

  Runtime API:cudaGetDeviceCount() 和 cudaGetDeviceProperties() 提供了遍历硬件设备,得到某个设备性能参数的功能。

  


int deviceCount;
  cudaGetDeviceCount(&deviceCount);
  int device;
  for (device = 0; device < deviceCount; ++device) {
  cudaDeviceProp deviceProp;
  cudaGetDeviceProperties(&deviceProp, device);
  }
  cudaSetDevice() 设置某一块Device作为这个主机host上的某一个运行线程的设备:

  cudaSetDevice(device);

  这个函数必须要在使用 __global__ 的函数或者Runtime

  的其他的API调用之前才能生效。 如果没有调用cudaSetDevice(),device 0 就会被设置为默认的设备,接下里的如果还有cudaSetDevice()函数也不会有效果。

  Driver API:

  cuDeviceGetCount()和cuDeviceGet() 看名字就知道干嘛的~(英语不好的这应该能看明白吧- -!不要被我这个那 国家四级都没过的人BS你哈~!~)

  


int deviceCount;
  cuDeviceGetCount(&deviceCount);
  int device;
  for (int device = 0; device < deviceCount; ++device) {
  CUdevice cuDevice;
  cuDeviceGet(&cuDevice, device);
  int major, minor;
  cuDeviceComputeCapability(&major, &minor, cuDevice);
  }
  3.内存管理:

  Device 上的内存可以被分配成线性的,也可以分配为CUDA的数组形式的。CUDA的内存可以为1维,2维,还有3维(2.0版本)。内存的类型有 unsigned8,16或者32位的int,16位(只有driver API可以做到)float,32位的float。这里分配的内存也只能通过kernel里面的函数通过处理纹理的方法来处理。这个地方也是GPU的历史 原因了,以前都是处理图像的,所以这里叫纹理。……叫啥都是别人取得名字 - -!在计算机里面不就是内存嘛 - -!

  Host 的runtime的运行库也提供按照page-locked的内存管理的函数,page-locked的内存要比pageable方式快很多。好的东西往 往比较少~page-locked就是很稀少的。如果通过减少分配pageable的内存来分配多的page-locked内存,系统需要的分页内存就少 了,这也就会让系统的性能降低了。所以在处理这块的时候要合理。

  Runtime API:

  使用 cudaMalloc() 或者 cudaMallocPitch() 来分配线性内存,通过cudaFree()释放内存.

  下面是分配一个大小为256 float数组的方法:

  float* devPtr;

  cudaMalloc((void**)&devPtr, 256 * sizeof(float));

  在使用2D数组的时候最好用cudaMallocPitch()来分配,在guide的第五章在讲到内存之间的调度的时候,就会看到他的好处。下面是一个分配一个大小为width×height 2D float数组的例子:

 


 // host code
  float* devPtr;
  int pitch;
  cudaMallocPitch((void**)&devPtr, &pitch,
  width * sizeof(float), height);
  myKernel<<<100, 512>>>(devPtr, pitch);
  // device code
  __global__ void myKernel(float* devPtr, int pitch)
  {
  for (int r = 0; r < height; ++r) {
  float* row = (float*)((char*)devPtr + r * pitch);
  for (int c = 0; c < width; ++c) {
  float element = row[c];
  }
  }
  }
   CUDA 的数组方式,需要用 cudaMallocArray()和cudaFreeArray(). cudaMallocArray()又需要cudaCreateChannelDesc()来管理,这个其实可以在第guide的第五章里面可以看到,我 们后面也会详细的介绍内存的调度和管理,和传统的GPU的内存方式不一样的地方.

  分配 width×height 32位float的CUDA array例子:

 


 cudaChannelFormatDesc channelDesc =
  cudaCreateChannelDesc();
  cudaArray* cuArray;
  cudaMallocArray(&cuArray, &channelDesc, width, height);
   cudaGetSymbolAddress用来在全局中定位一个数组的位置,然后cudaGetSymbolSize ()来回忆他分配的时候大小----如果有全局编程或者多线程编程经验的,或者用过几个函数同时处理一个数据的经验,都会为了把数据的独立性弄出来,不能 让数据和函数耦合太大,一般都不会让函数直接牵扯上数据,只是在函数处理的时候重新定位数据----这地方有点绕~~

  下面是一些例子,内存之间的拷贝,回想一下有几种内存~linear的有两个函数可以分配的,还有CUDA array的内存:

 


 cudaMemcpy2DToArray(cuArray, 0, 0, devPtr, pitch,
  width * sizeof(float), height,
  cudaMemcpyDeviceToDevice);
  The following code sample copies some host memory array to device memory:
  float data[256];
  int size = sizeof(data);
  float* devPtr;
  cudaMalloc((void**)&devPtr, size);
  cudaMemcpy(devPtr, data, size, cudaMemcpyHostToDevice);
  从host上面拷贝内存到device的constant上面:

 


 __constant__ float constData[256];
  float data[256];
  cudaMemcpyToSymbol(constData, data, sizeof(data));
  Driver API:

  cuMemAllocPitch()被推荐来作为2D的数组分配函数,会在内存对齐方面做一些check~然后保证在处理数据(像 cuMemcpy2D()这样的拷贝函数) 的时候达到最优的速度。下面就是一个float的width×height的2D的数组,并在循环里面处理的例子: ----(这个地方扩展提两个东东:一个是内存对齐,做程序优化的时候,或者处理网络问题的时候,都会遇到这些问题,内存对齐是一个普遍存在的问题,像 SSE这样的编程的时候也是要求内存对齐的,看以后要是有机会以单独讲解内存对齐的问题:)第二个是CUDA的内存访问的问题,就是很多朋友都会在写 kernel的时候,搞不明白里面的threadid,block id和传进来的内存的关系,这个地方必须要搞清楚的;内存和线程是CUDA编程必须搞明白的两个概念,不然到时候就会很混乱。其实看看前面的章节,应该能 看明白的,不明白就用手画图~要是感觉还是有点不太清楚,也不要担心,接下来的章节会单独把一些难懂的问题更详细的讲解。)

 


 // host code
  CUdeviceptr devPtr;
  int pitch;
  cuMemAllocPitch(&devPtr, &pitch,
  width * sizeof(float), height, 4);
  cuModuleGetFunction(&cuFunction, cuModule, “myKernel”);
  cuFuncSetBlockShape(cuFunction, 512, 1, 1);
  cuParamSeti(cuFunction, 0, devPtr);
  cuParamSetSize(cuFunction, sizeof(devPtr));
  cuLaunchGrid(cuFunction, 100, 1);
  // device code
  __global__ void myKernel(float* devPtr)
  {
  for (int r = 0; r < height; ++r) {
  float* row = (float*)((char*)devPtr + r * pitch);
  for (int c = 0; c < width; ++c) {
  float element = row[c];
  }
  }
  }
  cuArrayCreate()和cuArrayDestroy()来创建和释放CUDA array类型的数组。下面是一个 width×height 32-bit float类型的CUDA array 分配的例子:
  CUDA_ARRAY_DESCRIPTOR desc;
  desc.Format = CU_AD_FORMAT_FLOAT;
  desc.NumChannels = 1;
  desc.Width = width;
  desc.Height = height;
  CUarray cuArray;
  cuArrayCreate(&cuArray, &desc);
  这个也是几个内存之间拷贝的例子:

 


 CUDA_MEMCPY2D copyParam;
  memset(?Param, 0, sizeof(copyParam));
  copyParam.dstMemoryType = CU_MEMORYTYPE_ARRAY;
  copyParam.dstArray = cuArray;
  copyParam.srcMemoryType = CU_MEMORYTYPE_DEVICE;
  copyParam.srcDevice = devPtr;
  copyParam.srcPitch = pitch;
  copyParam.WidthInBytes = width * sizeof(float);
  copyParam.Height = height;
  cuMemcpy2D(?Param);
  拷贝host上面的内存到device上面:

  


float data[256];
  int size = sizeof(data);
  CUdeviceptr devPtr;
  cuMemAlloc(&devPtr, size);
  cuMemcpyHtoD(devPtr, data, size);
   本来想一次把下面的4, 5点也讲了~但是如果一下讲出来~lz帖子太长了~~这就不好了~~呵呵,实在的,看第四章的中文翻译的时候,就看了前面几个就不想看了~帖子不能老长老 长的又不吸引人- -!hoho~所以后面的4,5就在下贴里面发了~~看了这么多也比较累的~好好的休闲一下~你还会发现出了4,5部分,还少了一个部分,那就是CUDA 自己的函数,怎么定义,有device的,有global的,这个又有怎么区分啦~且听下回讲解。
  评论这张
 
阅读(290)| 评论(0)
推荐 转载

历史上的今天

评论

<#--最新日志,群博日志--> <#--推荐日志--> <#--引用记录--> <#--博主推荐--> <#--随机阅读--> <#--首页推荐--> <#--历史上的今天--> <#--被推荐日志--> <#--上一篇,下一篇--> <#-- 热度 --> <#-- 网易新闻广告 --> <#--右边模块结构--> <#--评论模块结构--> <#--引用模块结构--> <#--博主发起的投票-->
 
 
 
 
 
 
 
 
 
 
 
 
 
 

页脚

网易公司版权所有 ©1997-2017