OpenHero 开勇's profile开勇 OpenHeroPhotosBlogListsMore Tools Help

Blog


    April 30

    CUDA硬件实现分析(一)------安营扎寨-----GPU的革命

     CUDA硬件实现分析(一)------安营扎寨-----GPU的革命

    CUDA硬件实现分析(一)------安营扎寨

    ------GPU的革命

    序言:有个不会写计算机程序的朋友看了blog,问我,这个GPU也能当故事写吗?我觉得或许GPU真的算是一场革命吧,他的发展或许在酝酿中,不过到08年底,09年初,一定会有一场轰轰烈烈的竞争。那个时候或许从OS层面都会给人带来震撼。如果把CPU的多core看成由几个特种兵组成的,每个特种兵都手里面都拿着8杆枪(SSE)。那么GPU可以看成农民起义……一上来就是成百上千的人,虽然单兵作战能力比不上CPU的单个core,但是毕竟人数众多。就现在GPU的性能,在并行运算上如果不考虑double硬件的成本,已经早早超过CPU的并行运算能力。这或许就是一场革命,这次革命不知是简简单单的GPUCPU的转变,而是并行算法和串行算法的竞争。并行算法虽然研究到现在已经有很多年,但是真正的实际运用,离我们普通大众还是差很远。但是GPU,并行计算的出现,一下子把我们和并行计算的距离拉近了好多。现在在学校里面学习计算机的时候都是从串行算法开始,养成了很多固定的串行思维。当遇到问题并行划分的时候,就还带着串行的思想,那就不好了:)

    正文:前面我们已经说到线程的一些概念,但是这些概念都是软环节的。我们常常会听到某某单位说他们的软硬件配置如何如何的好。软件再好,每个士兵都是可造之才,但是如果硬件条件跟不上,也没他们的勇武之地。就像国内做过一些跳槽原因的统计,很多人已经跳槽都是为了高工资,但实际统计结果表明,很多人都觉得在以前那个公司里面学不到东西,或者说得不到发挥自己长处的地方。那就得看公司有没有这样的机会让你发挥你的能力了。看到这里,扯远了~很多人都看得不耐烦了……书接上回 CUDA 线程执行模型分析(二)大军未动粮草先行------GPU的革命。已经讲到CUDA在线程模型是一个什么样子,经过几天的吸收,也应该在脑子里面有一些印象了。但是你会问一个问题,我们是不是可以开无数个线程来执行啦?或许招兵的人都想找到很多很多的人,但是你得也考虑一下你的粮食有多少,军营有多大。在这么我们得讨论一下现在的支持CUDAnvidia的显卡现在的硬件情况。

    <!--[if !vml]--><!--[endif]-->空洞的讲解或许还是没怎么又说服力,下面以G80为例子。
    1. G80
    里面有16Multiprocessor.
    2.
    每个Multiprocessor都有一组(G80里面是8个)32位的Processor(每个Processor都是SIMD架构,什么叫SIMD架构:军训的时候,大家都到了食堂,不是像在学校里面,每个人那自己的碗筷就三三俩俩的去吃饭,那可要讲究纪律,啥叫纪律,一群人站在桌子面前,连长没发话,谁也不敢坐下来……连长一声令下:“坐下”。所有的人才按照敢坐下来,也是同时坐下来- -!要是谁没有同步坐下来,那就惨了- -!再来一次,一定要是同步坐下去的,都能听到声的,咵!恨不得把板凳给坐碎了------还是部队的东西结实,兄弟们怎么坐到军训结束只有坐坏屁股的,没有听说凳子坐坏的- -!所以啊,爱惜公物就是爱惜自己。记住了吧,这就是SIMDSingle Instruction Multiple Data )还有共用的Instruction Unit(这玩儿就不用翻译是啥了吧~看SIMD,自己理解去)。在G80里面有两个Instruction Unit模块。
    3.
    每一个时钟周期内,按照warp(这玩儿咋翻译啦?就理解为运行的时候,一个block里面一起运行的thread,例如block里面有512thread,但是每次只有32thread在运行,那么这32thread就是一个运行的warp- -! 还好不是rap- -!俺就真没法解释了)
    4.
    每一个warp里面包含的thread数量是有限的,现在的规定是32个。将来不知道会不会有变化?不知道,这个只有CUDA开发人员知道了。

    我们还是按照我们的既定方式学习吧~看图说话- -!接下来又是一张图:<!--[if !vml]--><!--[endif]-->
    刚才我们已经说到了硬件上的处理器的模型,现在我们来看看内存吧。前面的章节我们已经讲到了内存的registerlocalsharedconstanttextureglobal,现在看到对应的硬件的位置了吧?哦,对了,有人会问,localglobal跑那里去了,你看看,是不是多了一个Device Memory嘛。LocalGlobal 都是在Device Memory上的。只是在线程模型运行的时候,为了方便说明线程的模型,把它单独划分出来好说明。其实ConstantTexture都是在Device Memory里面的,但是人家是只读的Cache,所以比Globallocal要快一些。不过Cacheshared Memory比起来,又显得慢一些了。

    硬件的架构差不多就是这么多。如果还不是很清楚,我们再讲讲军营的生活。还是讲吃食堂吃饭吧(回想起大学的时候和几个好朋友一起成立了一个编程爱好者协会……结果后来变成了好吃协会- -!)。由于我们现在的连队人数众多,一个grid里面就可以有65535thread,所以一下子让这么多人一起吃饭,是不可能的。我们有好多个grid怎么办?所以每次Device(就是显卡)只处理一个grid。但是每个grid里面又有那么多人。一来是咱们地方小,不能容纳那么多人一起吃饭,而来是也没必要再建立那么多食堂。大家可以轮流吃,早吃的早走,晚来的也晚不了多久。

    这次是Grid1先到达食堂,然后他把士兵分成了几个block,(我们上次已经讲过的,blockgrid的关系,应该还记得把)但是现在桌子还是少,为了大家不乱起来,每个block的人就指定到某一个multiprocessor(一个周期只能容纳32个人一起用餐)哪里。

    这个block1的人就围着multiprocessor1这张桌子坐,但是这张桌子也不是太大,所以block里面的人又分成32个人一组(warp)来吃饭。那如果这个block没有32个人,就占不完这张桌子,所以一个multiprocessor又可以让别的blockthread一起来用餐。

    因为一次就把所有的registershared memory分配给了所有的block的所有的thread。所以怎么样合理的少用registershared memory就可以同时让更多的thread得到处理。

    讲了这么多,怎么来看自己的显卡那片军营如何拉?下面给出两种方法,一种是用一个小软件,GPU-Z的工具,这个工具可以得到很多显卡的参数,这个可以上网搜索到。

    <!--[if !vml]--><!--[endif]-->

    第二种方法是用CUDA的命令,这个只是针对支持CUDA的显卡的,下面给一段代码:

    /********************************************************************

    *  InitCUDA.cu

    *  This is a init CUDA of the CUDA program.

    *********************************************************************/

    #include <stdio.h>

    #include <stdlib.h>

    #include <cuda_runtime.h>

     

    /************************************************************************/

    /* Init CUDA                                                            */

    /************************************************************************/

    bool InitCUDA(void)

    {

      int count = 0;

      int i = 0;

     

      cudaGetDeviceCount(&count);

      if(count == 0) {

            fprintf(stderr, "There is no device.\n");

            return false;

      }

     

      for(i = 0; i < count; i++) {

            cudaDeviceProp prop;

            if(cudaGetDeviceProperties(&prop, i) == cudaSuccess)

            {                

                  printf("name:%s\n",                                         prop.name);

                  printf("totalGlobalMem:%u\n",                   prop.totalGlobalMem);

                  printf("sharedMemPerBlock:%u\n",                prop.sharedMemPerBlock);

                  printf("regsPerBlock:%d\n",                           prop.regsPerBlock);

                  printf("warpSize:%d\n",                               prop.warpSize);

                  printf("memPitch:%u\n",                               prop.memPitch);

                  printf("maxThreadsPerBlock:%d\n",               prop.maxThreadsPerBlock);

                  printf("maxThreadsDim:x %d, y %d, z %d\n",      prop.maxThreadsDim[0],prop.maxThreadsDim[1],prop.maxThreadsDim[2]);

                  printf("maxGridSize:x %d, y %d, z %d\n",      prop.maxGridSize[0],prop.maxGridSize[0],prop.maxGridSize[0]);

                  printf("totalConstMem:%u\n",                    prop.totalConstMem);

                  printf("major:%d\n",                                  prop.major);

                  printf("minor:%d\n",                                  prop.minor);

                  printf("clockRate:%d\n",                              prop.clockRate);

                  printf("textureAlignment:%u\n",                       prop.textureAlignment);

     

                  if(prop.major >= 1) {

                        break;

                  }

     

            }

      }

      if(i == count) {

            fprintf(stderr, "There is no device supporting CUDA 1.x.\n");

            return false;

      }

      cudaSetDevice(i);

           

      printf("CUDA initialized.\n");

      return true;

    }

    April 28

    April 28

    心理面有点很无助地感觉~。心理面有揪心的痛……
    好想回到母亲的身边,好想,好想……妈妈受的苦,已经够多了,够多了……为什么现在还要受这样的痛苦。
    母亲用一次次面对挫折的行动教会了我什么叫做坚持,什么叫做永不放弃。为了我,你已经付出很多了~
    现在还要受苦,做儿子的心理好痛苦,这种痛苦,揪心的痛。好想现在就飞到你的身边。

    真想骂,那CNN的是什么样的治安!在市区中心还能被抢劫!现在连罪犯的影子都没找到!
    CNN的是什么医院,手术做了以后还会做第二次!那是医疗事故!

    我现在又能做什么啦?
    慈母手中線. 遊子身上衣. 臨行密密縫. 意恐遲遲歸. 誰言寸草心. 報得三春暉。
    妈妈,孩子心理也能感受到你的痛~更痛~
    妈妈现在还不让我回家~很多事情都瞒着我,还不让我回家~当又一次预感到不对的时候,打电话回家询问,才知道~

    从重庆的小山沟到香港大都市,一步一步,突然发现自己还是那么的无助,连妈妈都不能照顾……
    可是我现在能做的又是什么啦……只能远远的祝福母亲早日康复,早日健康起来。

    CUDA 线程执行模型分析(二)大军未动粮草先行------GPU的革命

     CUDA 线程执行模型分析(二)大军未动粮草先行------GPU的革命

    CUDA 线程执行模型分析(二)大军未动粮草先行

    ------GPU的革命

    序:今天或许是比较不顺心的一天,从早上第一个电话,到下午的一些列的事情,有些许的失落。有的时候真的很想把工作和生活完全分开,但是谁又能真正的分得那么开,人非草木!很多的时候都想给人生下一些定义,添加一些注释。但是生活本来就是不需要添加注释的自解释的代码。用0来解释?还是用1来解释?0,天地之始,1,万物之源。谁又能说清楚,是先有0,还是先有1,他们本就是同体……要想成事,就应该拿得起,放得下。感叹人生的同时,人生的旅程是不会停止的……手下还有招来的那么多将士,都还等着啦!

    正文:书接上回--CUDA 线程执行模型分析(一)招兵------ GPU的革命》,经过几天的征兵,将士也招了差不多五六千人,五六千人,就是五六千张嘴,得吃饭。这和我们的CUDA内存模型又有什么相识啦?或许有很多人都参加过军训,知道军队里面吃饭吗?到吃饭的时间,各处训练的队伍都拉到食堂前面。先到的就先进去吃,后到的就在外面拉歌,等待下一波,这里我们暂时还不讲解怎么安排每一对吃饭。我们先讲讲每一个Grid手下的兄弟们吃饭都需要那些家伙。下面来看看站长图:

    <!--[if !vml]-->
    这里我们有一个block的士兵(16个人)在吃饭,每个人有一双筷子(Registers),一个碗(local memory),大家公用一张桌子(Block 里面的shared memory)。记得军训那年,在河北宣化,八个人一张桌子,一起抢桌子上的吃的时候,那个景象。训练了一天,一个个都是如狼似虎的,嘴巴里面咬着,筷子里面夹着,碗里面盛着,研究还看着桌子上的那几根鱼骨头(鱼肉早就进肚子里面了- -!)那个时候就看谁的筷子夹得多,碗装得多了。不过再多,筷子就那么长,碗就那么大~咱们再来看看block的情况:

    如图:……不需要注释了吧……

    餐厅里面肯定不止是一张桌子了,我已经不记得当时军训的时候,那个军营的食堂有多少个桌子,不过我们吃饭的时候可以装下几个连一起吃饭……想想那吃饭的阵势……

    哦,对了,我还记得在食堂里面有几个大桶,那玩儿应该不能称之为桶了,太大了 - -!可以分为三类,一类是装菜的,一类是装饭的,一类是装水果的。这几个桶都是大家可以共用的,记得有好吃的菜,大家都最先吃完那个菜,然后再跑到中间这个几个桶的地方找这个菜,装完一份跑到桌子那边一起分享……很是怀念那段日子。 Long long ago了……

    我能回想起来的那个时候的饭堂,如图(一个大的Grid吃饭图……):

    <!--[if !vml]--><!--[endif]-->


    CUDA的内存模型也就这个样子了~

    这些都是我们承诺给士兵的,每个人都要吃上饭。但是实际上我记得食堂再大,也不能一下子装下两个团的人吃饭,那个时候都是每个团里面的几个连队先吃饭,然后后到的就在外面拉歌,等里面的吃好了,后面的再进去吃……不是每个食堂都像北大的“万人坑”那样能装1w- -实地考察,好像要是真装了一w人,大家也就不用坐着吃饭了~- -

    对应手册上的第二章,运行模型就讲完了,接下来就要把队伍拉出来练习练习了……

    对了,我们的SIMD还没讲啦~队伍人数那么多,现在就那么几杆枪~怎么练习,还没讲啦……接下来的章节更精彩……

    PS: 题外话,或许是小的时候喜欢画图吧~现在还记得小学的美术老师给我们上的第一堂美术课,在黑板上画的正方形。在那个时候好像脑子里面就可以看穿物体一样,看到物体的这一面,就可以在大脑里面勾画出他的另一面~记得大学的时候参加定向越野,看到地图的时候,第一感觉,图形都好熟悉,然后再和实际的物体一对照,都是那么清晰,看到一栋建筑物,一对照地图,就可以知道他的后面是什么样,仿佛整个地图就在脑子里面构建起了地形图~呵呵很是奇妙……不知道大家有没有这样感觉的时候……

    April 25

    CUDA 线程执行模型分析(一)招兵---GPU的革命

    CUDA 线程执行模型分析(一)招兵------ GPU的革命

    新一篇: CUDA 线程执行模型分析(二)大军未动粮草先行------GPU的革命

     

    CUDA 线程执行模型分析(一)招兵
    ------ GPU的革命
     序:或许看到下面的内容的时候,你会觉得和传统的讲解线程,和一些讲解计算机的书的内容不是很相同。我倒觉得有关计算机,编程这些方面的内容,并不都是深奥难懂的,再深奥难懂的事情,其实本质上也是很简单的。一直以为计算机编程就像小时候搭建积木一样,只要知道游戏规则,怎么玩就看你自己了。或许是从小学那会,就喜欢在做数学题的时候用一些简便方法来解题,养成了一些习惯,喜欢把复杂的问题都会尝试用最简单的额方法来解决,而不喜欢把简单的问题弄得很复杂。不再多说了,有的朋友已经看得不耐烦了……ps:再罗嗦一句,如果下面看不明白的,就当小说看了,要是觉得不像小说,那就当故事看,要是觉得故事不完整,写得太乱,那就当笑话看,在各位学习工作之余能博得大家一笑,也倍儿感荣幸……ps2:想好再说……突然想到了,确实是了一段时间再想到的,既然叫GPU革命,那就得招集队伍啊,下面我就开始招兵了。
    正题:
    要真正进入CUDA并行化开发,就必须先了解CUDA的运行模型,才能在这个基础上做并行程序的开发。
    CUDA在执行的时候是让host里面的一个一个的kernel按照线程网格(Grid)的概念在显卡硬件(GPU)上执行。每一个线程网格又可以包含多个线程块(block),每一个线程块中又可以包含多个线程(thread)。
    在这里我们可以拿古时候的军队作为一个例子来理解这里的程序执行模型。每一个线程,就相当于我们的每一个士兵,在没有当兵之前,大家都不知道自己做什么。当要执行某一个大的军事任务的时候,大将军发布命令,大家来要把对面的敌人部队的n个敌人消灭了。然后把队伍分成M个部分,每一个部分完成自己的工作,有的是做侦查的工作,有的是做诱敌的工作,有的是做伏击的工作,有的是做后备的工作,有的是做后勤的工作……反正把一个大任务按照不同的类别,不同的流程不同,分别由M个部分来完成。
    这里我们可以把大将军看着是Host,它把这次军事行动分解成一个一个的kernelkernel_1kernel_2……kernel_M,每一个kernel就交给每一个Grid(副将?千户?就看管的人多人少了,如果GPU硬件支持少一点,那就是千户;要是GPU硬件高级一些,管理的人多一些,那就副将?戚家军也不过四五千人,咱也不能太贪心,一下子就想统军百万,再说了,敢问世上韩信一样的将才又哪有那么多啦?)来完成。当要执行这些任务的时候,每一个Grid又把任务分成一部分一部分的,毕竟人太多,他一个官不过来,他只要管理几个团队中间的高级军官就可以了。Grid又把任务划分为一个个的Block(百户?),这里每一个Grid管理的Block也是有限的,(人就那么多……想管多少得看硬件的支持)。毕竟显卡上的GPU硬件还是很少,Thread(线程)相对于真正的军队来说人还是少了很多。所以到Block这个层的时候,就直接管理每一个Thread(士兵)。
    由于古代通信不是很方便(从GPU的发展史来看,如果按照中国的历史,现在的GPU也就还处在战国时代吧……),所以每一个Block(百户)内部的Thread(士兵)才能方便的通信,按照既定的规则进行同步;而各个block之间就没那么方便了,大家不能互相通讯。不过同一个(千户)Grid管理的block之间是共享同一个任务分配的资源的。每一个Grid都可以从大将军那里分配到一些任务,和一些粮食,同一个Gridblock都可以分到这个Grid分配到的粮食。而每一个(千户)Grid本身的任务就不一样,所以Grid除了知道自己做的事情外,其他的Grid他都不会知道了。----这差不多就是一个运行模型。下面让我们来看看在GPU中东图例说明:
    看到这张图,我们可以对应来讲解我们的Thread部队。一个大将军Host,分配了任务中的两个任务(Kernel1, Kernel2)给了千户(Grid1Grid2)来完成。千户Grid1里面把自己的队伍分成了6个百户Block,然后每一个百户又把任务分配给了自己的士兵(Thread)来具体完成。这里得说明的是,由于千户拿到的任务Kernel是定了的,所以到每个士兵(Thread)也就那里就只会埋头做同样的事情(就像戚继光招的兵:在胡宗宪的幕僚郑若曾所著的《江南经略》中,有着这样一份详细的招生简章,如果不服气,大可以去对照一下:凡选入军中之人,以下几等人不可用,在市井里混过的人不能用,喜欢花拳绣腿的人不能用,年纪过四十的人不能用,在政府机关干过的人不能用。以上尚在其次,更神奇的要求还在下面:喜欢吹牛、高谈阔论的人不能用,胆子小的人不能用,长得白的人不能用,为保证队伍的心理健康,性格偏激(偏见执拗)的人也不能用。……概括起来戚继光要找的是这样一群人四肢发达,头脑简单,为人老实,遵纪守法服从政府,敢打硬仗,敢冲锋不怕死,具备二愣子性格的肌肉男----《明朝那些事儿》)。
     
    为了方便统一管理,大家都去掉了自己的名字,按照Grid1Blockxy),Threadxy)这样的编号来称呼每一个Thread士兵。如果你要找到某一个Thread,你就跑到军营里面大叫:喂,Grid1手下的Block1管理的三排第二个Thread12)出来。对于每个士兵自己来说,他要知道自己的位置,就得知道自己的长官都是谁。Thread12)要知道自己再整个Grid手下算第几个兵(钢七年第……个兵),当Grid1叫到他的号了,他得马上回答:我在Block11)的编号是:   
    unsigned int xIndex = blockDim.x * blockIdx.x + threadIdx.x;
     unsigned int yIndex = blockDim.y * blockIdx.y + threadIdx.y;
    如果要得到线性的编号,我们可以自己算一下:这个士兵是在Grid部队下的第xIndex行,yIndex列站着(这里我们必须注意:blockDim.x (这里为5),blockDim.y(这里为3)不是block的坐标,这里是block的size,切记!)。如果从第一个士兵哪里算是线性编号0,那他的线性编号就是
    unsigned int index = xIndex(6) + size_x * yIndex(5); 这里的size_x就是一行一共有多少个士兵(Thread),例如上图,这里一行有3个block每一个block里面的每一行有5个Thread,所以size_x就应该为3×5=15,一个Grid的一行有15个士兵,那刚才叫道的那个人的线性编号就应该是……还要我算吗?如果不得81,自己再算一下……(编号是从0,开始的)计算过程:index = 6+5×15=81
    说了那么多,咱也不能光说不练假把式。下面给一个简单的Thread 测试的Demo。本来打算把整个代码都copy过来,但是考虑到又会被别人copy出去,这样copycopy去,在编程中很容易出现错误,所以作者也不提倡在做编程中直接copy代码,这样很危险的,很多自己都不知道的bug就隐藏在copy的代码当中……so,截图……
    作者这里做了一个简单的测试,测试了512个线程。这里只有一个grid,从这一个grid里面也只分了一个bock
    dim3    grid(size_x / BLOCK_DIM, 1);--》    dim3 grid(1, 1);一个Grid里面一个blcok
    dim3    block(BLOCK_DIM, 1, 1);--》dim3 block(512, 1, 1);一个Block里面分配512个Thread;
    这里的每一个任务kernel就是:
    __global__ static void ThreadDemo1(unsigned int* ret)
    {
     unsigned int xIndex = blockDim.x * blockIdx.x + threadIdx.x;
     unsigned int yIndex = blockDim.y * blockIdx.y + threadIdx.y;
     
     if(xIndex < size_x && yIndex < size_y)
     {
            unsigned int index = xIndex + size_x * yIndex;
            ret[index] = xIndex;
            ret[index + size_x*size_y] = yIndex;
     }
    }
    计算自己的线性id 然后把自己的坐标写入到线性id对应的数组里面。Ps:说明一下,这个记录id坐标的数组ret[],ret的前一半记录的是线程的x坐标,后一般是记录的y坐标。PS2:题外话,cu是C的扩展,这里的const定义的常量的用法在ANSIC C里面是行不通的,但是在C++中是可用的。
     
    每个任务kernel都说好了,然后就是host下达命令
     ThreadDemo1<<<grid,block>>>(ret);
    来运行程序。所有的士兵都开始工作了,把自己的坐标x,y写入到ret数组里面。
    下面是得到的结果:
    (0,0) (1,0) (2,0) (3,0) (4,0) (5,0) (6,0) (7,0) (8,0) (9,0) (10,0) (11,0) (12,0) (13,0) (14,0) (15,0) (16,0) (17,0) (18,0) (19,0) (20,0) (21,0) (22,0) (23,0) (24,0) (25,0) (26,0) (27,0) (28,0) (29,0) (30,0) (31,0) (32,0) (33,0) (34,0) (35,0) (36,0) (37,0) (38,0) (39,0) (40,0) (41,0) (42,0) (43,0) (44,0) (45,0) (46,0) (47,0) (48,0) (49,0) (50,0) (51,0) (52,0) (53,0) (54,0) (55,0) (56,0) (57,0) (58,0) (59,0) (60,0) (61,0) (62,0) (63,0) (64,0) (65,0) (66,0) (67,0) (68,0) (69,0) (70,0) (71,0) (72,0) (73,0) (74,0) (75,0) (76,0) (77,0) (78,0) (79,0) (80,0) (81,0) (82,0) (83,0) (84,0) (85,0) (86,0) (87,0) (88,0) (89,0) (90,0) (91,0) (92,0) (93,0) (94,0) (95,0) (96,0) (97,0) (98,0) (99,0) (100,0) (101,0) (102,0) (103,0) (104,0) (105,0) (106,0) (107,0) (108,0) (109,0) (110,0) (111,0) (112,0) (113,0) (114,0) (115,0) (116,0) (117,0) (118,0) (119,0) (120,0) (121,0) (122,0) (123,0) (124,0) (125,0) (126,0) (127,0) (128,0) (129,0) (130,0) (131,0) (132,0) (133,0) (134,0) (135,0) (136,0) (137,0) (138,0) (139,0) (140,0) (141,0) (142,0) (143,0) (144,0) (145,0) (146,0) (147,0) (148,0) (149,0) (150,0) (151,0) (152,0) (153,0) (154,0) (155,0) (156,0) (157,0) (158,0) (159,0) (160,0) (161,0) (162,0) (163,0) (164,0) (165,0) (166,0) (167,0) (168,0) (169,0) (170,0) (171,0) (172,0) (173,0) (174,0) (175,0) (176,0) (177,0) (178,0) (179,0) (180,0) (181,0) (182,0) (183,0) (184,0) (185,0) (186,0) (187,0) (188,0) (189,0) (190,0) (191,0) (192,0) (193,0) (194,0) (195,0) (196,0) (197,0) (198,0) (199,0) (200,0) (201,0) (202,0) (203,0) (204,0) (205,0) (206,0) (207,0) (208,0) (209,0) (210,0) (211,0) (212,0) (213,0) (214,0) (215,0) (216,0) (217,0) (218,0) (219,0) (220,0) (221,0) (222,0) (223,0) (224,0) (225,0) (226,0) (227,0) (228,0) (229,0) (230,0) (231,0) (232,0) (233,0) (234,0) (235,0) (236,0) (237,0) (238,0) (239,0) (240,0) (241,0) (242,0) (243,0) (244,0) (245,0) (246,0) (247,0) (248,0) (249,0) (250,0) (251,0) (252,0) (253,0) (254,0) (255,0) (256,0) (257,0) (258,0) (259,0) (260,0) (261,0) (262,0) (263,0) (264,0) (265,0) (266,0) (267,0) (268,0) (269,0) (270,0) (271,0) (272,0) (273,0) (274,0) (275,0) (276,0) (277,0) (278,0) (279,0) (280,0) (281,0) (282,0) (283,0) (284,0) (285,0) (286,0) (287,0) (288,0) (289,0) (290,0) (291,0) (292,0) (293,0) (294,0) (295,0) (296,0) (297,0) (298,0) (299,0) (300,0) (301,0) (302,0) (303,0) (304,0) (305,0) (306,0) (307,0) (308,0) (309,0) (310,0) (311,0) (312,0) (313,0) (314,0) (315,0) (316,0) (317,0) (318,0) (319,0) (320,0) (321,0) (322,0) (323,0) (324,0) (325,0) (326,0) (327,0) (328,0) (329,0) (330,0) (331,0) (332,0) (333,0) (334,0) (335,0) (336,0) (337,0) (338,0) (339,0) (340,0) (341,0) (342,0) (343,0) (344,0) (345,0) (346,0) (347,0) (348,0) (349,0) (350,0) (351,0) (352,0) (353,0) (354,0) (355,0) (356,0) (357,0) (358,0) (359,0) (360,0) (361,0) (362,0) (363,0) (364,0) (365,0) (366,0) (367,0) (368,0) (369,0) (370,0) (371,0) (372,0) (373,0) (374,0) (375,0) (376,0) (377,0) (378,0) (379,0) (380,0) (381,0) (382,0) (383,0) (384,0) (385,0) (386,0) (387,0) (388,0) (389,0) (390,0) (391,0) (392,0) (393,0) (394,0) (395,0) (396,0) (397,0) (398,0) (399,0) (400,0) (401,0) (402,0) (403,0) (404,0) (405,0) (406,0) (407,0) (408,0) (409,0) (410,0) (411,0) (412,0) (413,0) (414,0) (415,0) (416,0) (417,0) (418,0) (419,0) (420,0) (421,0) (422,0) (423,0) (424,0) (425,0) (426,0) (427,0) (428,0) (429,0) (430,0) (431,0) (432,0) (433,0) (434,0) (435,0) (436,0) (437,0) (438,0) (439,0) (440,0) (441,0) (442,0) (443,0) (444,0) (445,0) (446,0) (447,0) (448,0) (449,0) (450,0) (451,0) (452,0) (453,0) (454,0) (455,0) (456,0) (457,0) (458,0) (459,0) (460,0) (461,0) (462,0) (463,0) (464,0) (465,0) (466,0) (467,0) (468,0) (469,0) (470,0) (471,0) (472,0) (473,0) (474,0) (475,0) (476,0) (477,0) (478,0) (479,0) (480,0) (481,0) (482,0) (483,0) (484,0) (485,0) (486,0) (487,0) (488,0) (489,0) (490,0) (491,0) (492,0) (493,0) (494,0) (495,0) (496,0) (497,0) (498,0) (499,0) (500,0) (501,0) (502,0) (503,0) (504,0) (505,0) (506,0) (507,0) (508,0) (509,0) (510,0) (511,0)
    一目了然,正是我们定义的512个线程。
    现在大家应该知道CUDA是怎么运行了?0 0 要是还不清楚,跟贴问我吧 - - !
     
    好了,现在招兵的任务已经完成了,兵是招来了,下面就改训练了,不过天时已晚,大家还是吃饱喝足,好好的休息一下吧~
    明天继续GPU的革命。
    且听:《沙家浜》里胡傳魁已經出場了:十幾個人來七八條槍,不是很威風嗎?----我可要说十几个人怎么来分这七八条枪。
    士兵们执行任务了,是一哄而上,还是……且听SIMD- -!
     
    PS:上面在输出结果的时候用了一个小技巧,或许有的人知道,就不用看了,有不知道的,接着看下去,或许可以方便你以后的调试:)
    再输出的时候,在系统debugging 信息里面的command Arguments里面用一个管道输出到一个1.txt文件里面。管道用法大家自己可以到网上查~我就不用讲解了。也很简单,就是在你的哦xx.exe 后面加一个 “>” 管道,对了还有 “<” 管道,呵呵:) 一个是输出,一个是输入。
    April 24

    Visual Assist 添加支持*.cu文件

    Visual Assist 添加支持*.cu文件

     
    Visual Assist 添加支持*.cu文件
    1.       打开注册表,在HKEY_LOCAL_MACHINE\SOFTWARE\Microsoft\VisualStudio\8.0\Languages\File Extensions\ 下面添加子键 .cu 然后copy .cpp的键值到.cu。这样才能表示cu也是VS下的VC的工程文件。


    2.       打开注册表
    HKEY_CURRENT_USER\Software\Whole Tomato\Visual Assist X\VANet8 ExtSource键添加键值.cu
    3.       打开Visual Assist属性,在projects 的C/C++ Directories custom下面添加CUDA的头文件目录,这样才能在Visual Assist 生成规则的时候找到CUDA自身的特殊定义才能生成Visual Assist的关键字,如__global__.
     

    快速开发 CUDA windows 程序

      书接上回CUDA windows 开发 
    http://blog.csdn.net/OpenHero/archive/2008/04/10/2278164.aspx 

    已经写到建立一个简单的CDUA程序,但是上面需要设置的参数还是很多,这样对于初学者来说增加了难度。在这里作者开发了一个针对VS2005CUDA wizard CUDA VS2005 Wizardhttp://blog.csdn.net/OpenHero/archive/2008/04/18/2305856.aspx 这样就可以很方便的在VS2005的环境中进行CUDA的开发。
    下面是快速入门的流程:
    1。安装CUDA VS2005 Wizard》就可以可以在VS的工程里面看到一个CUDAWin32App的工程向导。
    CUDA new project
    点击这个就可以创建一个简单的CUDA工程,里面已经就有一个简单的example,可以进行CUDA开发了。
    2.我这里创建的是一个测试Thread demo,所以生成的cu文件名字就是threaddemo1CUDA threaddemo
    可以看到这里已经由VS wizard向导自动生成了一些简单的hello cuda的程序。
    3.我们再来看工程的环境配置,就可以发现Debug Release EmuDebug EmuRelease已经在系统设置里面了,已经由Wizard自动生成了工程的各个环节配置。
    configuration
    4.工程环境变量的查看和配置,右键点击工程,选择属性或者 “properties" 就可以看到project的环境配置里面就多了一个CUDA选项,这里就可以对CUDA进行环境配置。
    properties
     http://p.blog.csdn.net/images/p_blog_csdn_net/OpenHero/5.JPG
    5.编译程序
    http://p.blog.csdn.net/images/p_blog_csdn_net/OpenHero/6.JPG
    6.运行输出CUDA initialized. Hello CUDAtime
    http://p.blog.csdn.net/images/p_blog_csdn_net/OpenHero/7.JPG
    7.调试,需要把cu文件关联到VCproject里面,不然调试的时候不能断点跟踪进去,所以必须在系统的VC++project 环境中加入*.cu文件,VS才能把你的程序和源文件关联起来,进行调试。
    如果没找到关联cu,在调试的时候就会发现找不到源文件。
    cu set
    8.简单的快速开发就讲到这里了。我的字体这里显示高亮度,或许有些朋友会问,怎么才能让CU文件显示高亮度,可以查看我的另一篇翻译,《syntax highlighting when editing your .cu files in Visual Studio
    9.再下一节中,准备对CUDA的线程模型做一个简单的实用讲解;欢迎再来……hoho >_<!
    ps:你的留言是对我的最大鼓励,Open Heart Bravely fly!