OpenHero 开勇's profile开勇 OpenHeroPhotosBlogListsMore  |  Tools Help |
|
May 12 CUDA编程接口（一）------一十八般武器------GPU的革命CUDA编程接口（一）------一十八般武器------GPU的革命
CUDA编程接口（一）------一十八般武器
------GPU的革命
序言：所谓一十八般武器，不同的年代又有不同的说法，最早的汉武年间的：矛、镗、刀、戈、槊、鞭、锏、剑、锤、抓、戟、弓、钺、斧、牌、棍、枪、叉。到三国的：九长：刀、矛、戟、槊、镗、钺、棍、枪、叉；九短：斧、戈、牌、箭、鞭、剑、锏、锤、抓。再到明清的：弓、弩、枪、刀、剑、矛、盾、斧、钺、戟、黄、锏、挝、殳（棍）、叉、耙头、锦绳套索、白打（拳术）。《水浒传》里的：矛、锤、弓、弩、铳、鞭、锏、剑、链、挝、斧、钺、戈、戟、牌、棒、枪、扒。今天的武术届又有：刀、枪、剑、戟、斧、钺、钩、叉、鞭、锏、锤、抓、镗、棍、槊、棒、拐、流星。400多种古代冷兵器时代的武器，常用的也只有这么多种。也就像我们的API一样，API有无数多个，你自己都可以给自己造几个API出来，常用的，或者就那么多个。要打天下也不能扛着锄头，竹竿干吧。秦国之所以能统一六国，在武器上的统一，提供同一个的型号的武器装备（看秦的历史，就可以发现所有的兵器都是同一型号生产，弓弩上的器件可以互换，从兵马俑坑中找到的剑戟，箭头的尺寸误差很小，都可以互换），也是他能战胜其他六国的很好的基础。
正文：
子曰：工欲善其事，必先利其器。我们要把显卡作为通用并行处理器来做并行算法处理，就得知道CUDA给我提供了什么样的接口，就得了解CUDA作为通用高性能计算平台上的一十八般武器。（如果你想自己开发驱动，自己写开发库- -那我不得不佩服你很有时间，想必也不会有很多人想自己在去实现一个CUDA吧，呵呵，虽然实现一个也不是太难）。书接上回《 CUDA硬件实现分析（二）------规行矩步------GPU的革命》前面我们讲到了一些简单的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讲解
记得小时候练过几天棍法，现在好像起式都记不起来啦，哈哈，只记得棍要齐眉，剑齐耳朵，不知道错没，hoho。大学的时候一直想给武术队的张锐学九节鞭，不过一直太忙，除了和他练习过几次散打对抗，看他刀+鞭，双刀，枪，剑，九节鞭练习过几次以外，到大学毕业，也都没机会学九节鞭，～现在也就留着那一根九节鞭在身边。看到它就能想起一些朋友吧：）大学毕业也就各奔东西，都不知道他现在在干嘛……
又走神了……想想前面五个部分，下面一部分一部分的讲解：
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<float>();
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));
-#%#$^$^*^&系统int中断：(@$:突然发现已经是早上8点多了 - -！又一夜！@#￥%……！#……
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的，这个又有怎么区分啦～且听下回讲解。Hoho
ps：熬夜不好～熬夜很伤身……
有的时候，我们经常会用旧的东东来和新的东东比较，就像C和C++都不知道争论多少年了。其实很多时候我到觉得是没必要的争论，除非你是做C或者C++本身开发的。就像新的硬件不停的变化，以前的概念或者今天就不能用了，有的时候，我们关心架构在这个之上的程序开发就够了，没有太多的必要去问茴香豆的茴有几种写法。有的时候看到论坛里面的争论的时候，很多都不太清楚问题的也参加到争论里面，感觉就是明白事的没有几个，倒是起哄的不少 - -！人家说当局者迷，旁观者清，我看现在很多时候倒是当局者清，旁观者瞎起哄- -！呵呵，好像合乎了现在很多选秀节目的心理哈，起哄的人越多，人家的节目越红，哈哈～----扯远了……
Ps2：在学习API的时候，最好的方法就是多实验，多尝试API的性能，就像练习武术中的器械一样，经常用才会精通的，才能在这个基础上想出新的招式。就像有人在CSDN论坛问，怎么才能弄好ACM比赛，看到像刘汝佳写的《算法艺术与信息学竞赛》的书，都是一些方法，没有代码，就觉得很不解，就像问用什么代码来实现。我的回答就是：阅尽天下A×，心中自然无码。多做代码的练习，找一本数据结构的书，对照上面的代码都自己实现一遍，找一个代码练习的书，自己都重新写一遍。要不然就把像MSDN这样的介绍API的资料上的Demo能实现多少都去实现以下，呵呵。其实学习怎么编码都是一些基础工作。最重要的到最后都是算法的实现。其实到最后你会发现写程序就是数组的处理，就是字符串的处理……仅此而已～。所以不要小看了C语言那些书上的小程序例子，不要小看了输出（**）星星这样的例子，实际程序中很多时候都是处理这些星星～hoho～不要一上来就想去写什么游戏，写什么网络软件～先把字符串处理好了，就nb了，真的～你看ACM的题目，topcoder的题目，几乎都是字符串处理- -！哈哈。多实践，多去做一些demo例子。到真正用的时候就会觉得手到擒来，怎么用怎么顺手了。
还是那句话：约尽天下A×，心中自然无码。
再多ps一句：上次和刘汝佳聊天，他准备在新版的书中用java来添加一些代码实现讲解，哈哈～hoho May 05 CUDA硬件实现分析（二）------规行矩步------GPU的革命CUDA硬件实现分析（二）------规行矩步------GPU的革命
新一篇: ubuntu 8.04 install on xp， 3d desktop cool！
CUDA硬件实现分析（二）------规行矩步
------GPU的革命
序言：换位思考。当今的生活，节奏快，任务忙。慢慢的忽略了身边的很多事，很多人。再加上接受“高等”教育的人越来越多，“有自己思想的”人越来越多，慢慢的都习惯从自己的思维角度来思考问题，尤其是读工科的学生更是喜欢按照自己的角度来思考问题。慢慢的忽略了换位思考。有很多朋友说学工科的人都喜欢走极端。或许这个就像金庸小说里面少林高僧给两个偷学少林武功的人的建议。在忙碌的生活和紧张的工作中，找个时间，能让自己停下来，想想做过的事情，让自己忙碌的脚步，休息一会儿。往往在team开发中，遇到问题的时候，就需要沟通和交流，但是沟通和交流的基础就是换位思考。在一个平等的环境中的沟通和交流才能算真正意义上的思想的交流。其实学工科的时候有一个小窍门，那就是找规则。有既定的规则，那就是定理和定义。如果你能找到新规则，那就是新发现，可以写paper。我们遇到新东西的时候，也最好在自己的既有思维中找到影子，找到相同的规则。这样就可以很好的学习新东西。不过往往学工科的思维比较有规则性，在加上平时看的工科的书都是规则性太强，长此以往很容易形成偏执的性格。平时就需要多看一些能扩展思维的书籍，或许能消减一些戾气吧。
正文：前面已经讲解了很多概念上的东西，其实CUDA的最重要的两个东西，就是线程和内存。只要掌握了这两个东西，CUDA的东西也就很简单了。它的编写语言是C扩展的，所以，就当C语言用就行了，只是主要它的特殊的几个标志就ok了。前面讲解了线程和内存的模型，大概，应该，似乎，可以在你的脑海里面有一个概念了吧。只要有这个概念，我的文章的目的就达到了。前面的《CUDA硬件实现分析（一）------安营扎寨-----GPU的革命》已经讲解了线程在CUDA的具体运行过程。下面我们一起来看看内存在CUDA的硬件实现中的一些规定。这也比较合理吧，大军安营扎寨了，就应该颁布规则制度，只有了解CUDA的规则制度，才能真正的把各个线程都管理好。才能在这个平台上让程序高效的运行。
这里我们先明确几个
一．Threads，Warps, Blocks
1. 一个warp最多有32个threads。只有在总线程少于32的时候，才可能在一个warp里面少于32个线程。 2. 每一个block最多有16个warp。就是说一个一个block里面最多有512个thread。 3. 每一个Block在同一个SM上执行，也就是同一个block的warp都在同一个SM上运行。 4. G80有16个SM。 5. 所以最少16个blocks才能占全所有的SM。 6. 如果资源（看看前面讲解的线程都要从device哪里分什么资源）够线程分，一个SM上面可以跑多余一个block的线程。就是同时可以跑2个，3……个block的线程。 二．访问速度
Register—HW 一个时间周期 Shared Memory------HW一个时钟周期 Local Memory --- DRAM，no cache，慢 Global Memory --- DRAM， no cache，慢 Constant Memory --- DRAM， cached， 1……10s……100s个周期，这个和cache的locality有关。 Texture Memory --- DRAM，cached， 1……10s……100s个周期，这个和cache的locality有关。 Instruction Memory（不可见）--- DRAM，cached 三．CUDA程序架构
如图 四．语言扩展
单从学习语言来说，我到觉得应该精学一种，然后其他的就触类旁通了。在学习新语言的时候，要想快速的入门，也有诀窍。1.变量的定义方式。2.函数的定义方式。3.逻辑控制方式（if，loop……）。只要把这3个东西弄明白了，管他啥新的语言，20分钟就可以入门……然后入门都可以了，那要慢慢的更入的研究，那就得看你对这门语言的了解了。其实万变不离其中。其实从计算机编程语言的角度出发，就是定义一些数据，然后对数据进行操作，so……学习语言就从这个角度入手，那就很简单了。像java或者C#等一些语言比C语言多的新的特性不外乎就是方便你开发而已。 所以我们这里再来说说CUDA的语言，不外乎就是扩展了C语言，为了方便在GPU显卡上运行，规定一个特定环境。就定义一些特定的变量，说明他们是在GPU上的。这里有说明内存，和函数，是在GPU上的……so，这样一来，CUDA就扩展了C语言的变量分配定义和函数定义。
 上面这张图是来之Fall 2007 syllabus，上面已经说得很清楚各个变量定义的时候的位置和生存周期。其实就是在C语言的常规变量的时候，定义了变量的位置而已。其中有一个约束限制，就是指针变量，在kernel里面的指针变量，只能指向从global上面分配的内存。
五．内建的变量
所谓内建，就是CUDA自己在kernel里面定义的一些变量。就像我们以前计算线程id的时候，就利用了他的自己的变量，dim3 gridDim；dim3 blockDim；dim3 blockIdx；dim3 threadIdx；注意gridDim，他的gridDim.z在现在的CUDA1.1版本中没定义。 内建变量，[u]char[1..4], [u]short[1..4], [u]int[1..4], [u]long[1..4], float[1..4] 就是构建了有4个变量的struct； uint4 param; ---》等价位一个struct里面有4个int。 int y = param.y; dim3 就是unit3这样的struct。 六．通用的数学函数
· pow, sqrt, cbrt, hypot
· exp, exp2, expm1
· log, log2, log10, log1p
· sin, cos, tan, asin, acos, atan, atan2
· sinh, cosh, tanh, asinh, acosh, atanh
· ceil, floor, trunc, round
但是这里要指出，有几个函数是不精确的，但是可以很快的运行：
– __pow
– __log, __log2, __log10
– __exp
– __sin, __cos, __tan
七．在host部分的运行库（CUDA Runtime ）
1. 提供Device管理的api（有多个显卡的时候怎能来设置，这里我们还没讲当多显卡并行运行库）
2. 初始化调用的Runtime函数
3. 每一个host的线程（thread）只能调用一个device的函数在一个device上运行。就是同时不能有几个host（主机）线程在调用同一个device上的运行函数。
八．内存管理函数
我们讲了那么多的内存，下面就来看看到底是什么样的函数： – cudaMalloc（），cudaFree（）；内存分配和释放（device上的）
– 内存copy：cudaMemcpy(), cudaMemcpy2D(), cudaMemcpyToSymbol(), cudaMemcpyFromSymbol()；
九．线程同步函数
void __syncthreads();
同步同一个block里面的线程，让block里面的线程都允许到这一点的时候，就等待同一个block里面的其他线程，就像军训的时候，大家吃饭吃完了还不能一个个走，必须得一个桌子的人都吃完了才能走。还得列队一起走，呵呵，这就是同步。所以最好保证每一个kernel里面的处理都是很快的，这样才不会让其他thread等待太久，不然会挨骂的 - -！hoho
哥们来点实际的吧- - 肯定很多人都这么在吼了。嘿嘿，下面让我们看一段代码，example里面的transpose：
#define BLOCK_DIM 16
// This kernel is optimized to ensure all global reads and writes are coalesced,
// and to avoid bank conflicts in shared memory. This kernel is up to 11x faster
// than the naive kernel below. Note that the shared memory array is sized to
// (BLOCK_DIM+1)*BLOCK_DIM. This pads each row of the 2D block in shared memory
// so that bank conflicts do not occur when threads address the array column-wise.
__global__ void transpose(float *odata, float *idata, int width, int height)
{
__shared__ float block[BLOCK_DIM][BLOCK_DIM+1]; //(1)
// read the matrix tile into shared memory
unsigned int xIndex = blockIdx.x * BLOCK_DIM + threadIdx.x;
unsigned int yIndex = blockIdx.y * BLOCK_DIM + threadIdx.y;
if((xIndex < width) && (yIndex < height)) // (2)
{
unsigned int index_in = yIndex * width + xIndex;
block[threadIdx.y][threadIdx.x] = idata[index_in];
}
__syncthreads(); //(3)
// write the transposed matrix tile to global memory
xIndex = blockIdx.y * BLOCK_DIM + threadIdx.x;
yIndex = blockIdx.x * BLOCK_DIM + threadIdx.y;
if((xIndex < height) && (yIndex < width))
{
unsigned int index_out = yIndex * height + xIndex;
odata[index_out] = block[threadIdx.x][threadIdx.y];
}
}
是不是在代码面前都感觉是那么的亲切啊～嘿嘿，有时候看算法的书看久了以后，看到代码实现，那才叫亲切啊>_<!
下面对加红的3处代码分别讲解：
（1） shared__ float block[BLOCK_DIM][BLOCK_DIM+1];
为啥要用shared啊？知道为啥不，快啊～shared在block里面，是大家公用的，不用再从global哪里调换。还记得军训的时候食堂吃饭的时候那个内存讲解吗？桌子上的菜，肯定比你从中间几个桶里面去拿菜要快啊～ （2） if((xIndex < width) && (yIndex < height))
由于CUDA在生成线程的时候，是一块一块的生成的，所以有的时候，x和y有可能越界。数据都是均匀的分配到每一个block里面去的，很可能就会出现有的blcok里面分配到的数据不一定完整，要是想不明白的（先画图想想），等后面的章节，举例子说明。 （3） __syncthreads();
同步，就是等每一个thread都运行到这里的时候，再接下来运行。这里也避免了数据访问的冲突。 （4） 还有一个得讲讲，为啥下面的还重新计算一下x和y，因为线程被挂起来了以后，重新进入warp的时候，分配的id就不一定是一样的了。这里就有点像铁打的营盘流水的兵，每一波人都可以分配到同一张桌子吃饭，但是不一定是上一波人了～。碗筷还是那些碗筷，位置还是那个位置，坐的人不一样了。
上面的都是C语言扩展库的东西，就是说，这些是扩展C语言的东东，只要是基于C语言的规则都可以运行的，不管是host还是device。
好了，今天要讲的部分应该差不多了，要是再写多一些，或者就有会像手册第4章那样，看了一半很多人就不想看了～- -！太长，lz 帖子太长，顶lz- -！我可不想看到这样的留言，hoho～
从小到大，我们听过好多规则，但是好多时候听进去的有多少？很多玩游戏的人，游戏的规则都不清楚- -ps，俺现在还不清楚StarCraft里面的某些兵的战斗力是多少 - -！就玩- -！小小的bs一下自己。只是为了消遣罢了·
我始终觉得自己不适合玩游戏，我也知道自己玩游戏不会完成专业玩家，so～还是老老实实做自己的工作，做一些能养活自己的事情，或许能让自己更充实一些。在游戏里面充实以后，或许在现实中就会空虚了- -！到时候就不是游戏人生，那就是人生被游戏了。呵呵。想起高中的数学老师，虽然在很多journal上发表了文章，也自己出版了数学的几本书了。但是和别人打麻将的时候总是输多赢少。问他为啥数学这么厉害，玩麻将不行～他说，我玩麻将的时候只是消遣，根本就没用脑袋想麻将这回事。
或许我们做事情的时候，就应该专注做自己的一些方面吧，只要坚持下来做了，做好了，做精了，那就够了。小时候就从长辈哪里听到，半灌水叮当响这样的词，或许我们更应该低调做人，高调做事吧。
Ps：说得都有点多了- -就当消遣。俺熬夜写这个也不容易啊，哈哈～>_<!lz，又没人逼你写。
|
|
|
