子曰：工欲善其事，必先利其器。我们要把显卡作为通用并行处理器来做并行算法处理，就得知道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);
这个函数必须要在使用 __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); } |
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]; } } } |
分配 width×height 32位float的CUDA array例子:
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(); cudaArray* cuArray; cudaMallocArray(&cuArray, &channelDesc, width, height); |
下面是一些例子，内存之间的拷贝，回想一下有几种内存～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); |
__constant__ float constData[256]; float data[256]; cudaMemcpyToSymbol(constData, data, sizeof(data)); |
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); |
float data[256]; int size = sizeof(data); CUdeviceptr devPtr; cuMemAlloc(&devPtr, size); cuMemcpyHtoD(devPtr, data, size); |