并行课件备注

5-10不需别人用pthread_join清理门户，自己了断了。由于线程也占用资源，如果你不设置成为detach状态，那么当你的线程推出后，你必须执行

pthread_join调用才能释放这些被占用的资源，如果设置成detach状态，线程再推出后将自动释放自己占用的资源

这些占用的资源不包括使用malloc分配的内存和ipc资源

5-39其实函数的执行过程非常简单，在第一个线程执行到pthread_cond_wait(&cond,&mut)时，此时如果X<=Y，则此函数就将mut互斥量解锁，再将cond条件变量加锁，此时第一个线程挂起（不占用任何CPU周期）。

而在第二个线程中，本来因为mut被第一个线程锁住而阻塞，此时因为mut已经释放，所以可以获得锁mut，并且进行修改X和Y的值，在修改之后，一个IF语句判定是不是X>Y，如果是，则此时pthread_cond_signal()函数会唤醒第一个线程，并在下一句中释放互斥量mut。然后第一个线程开始从pthread_cond_wait()执行，首先要再次锁mut，如果锁成功，再进行条件的判断（至于为什么用WHILE，即在被唤醒之后还要再判断，后面有原因分析），如果满足条件，则被唤醒进行处理，最后释放互斥量mut。

至于为什么在被唤醒之后还要再次进行条件判断（即为什么要使用while循环来判断条件），是因为可能有“惊群效应”。有人觉得此处既然是被唤醒的，肯定是满足条件了，其实不然。如果是多个线程都在等待这个条件，而同时只能有一个线程进行处理，此时就必须要再次条件判断，以使只有一个线程进入临界区处理。6-20Sectionsaredistributedamongthethreadsintheparallelteam.Eachsectionisexecutedonlyonceandeachthreadmayexecutezeroormoresections.It’snotpossibletodeterminewhetherornotasectionwillbeexecutedbeforeanother.Therefore,theoutputofonesectionshouldnotserveastheinputtoanother.Instead,thesectionthatgeneratesoutputshouldbemovedbeforethesectionsconstruct.6-22DataScopeAttributesAlldataclausesapplytoparallelregionsandworksharingconstructsexcept“shared,”whichonlyappliestoparallelregions.6-23PrivateCauseFor-loopiterationvariableisPRIVATEbydefault.6-29AtomicConstructSinceindex[i]canbethesamefordifferentIvalues,theupdatetoxmustbeprotected.Useofacriticalsectionwouldserializeupdatestox.Atomicprotectsindividualelementsofxarray,sothatifmultiple,concurrentinstancesofindex[i]aredifferent,updatescanstillbedoneinparallel.7-4We’vealreadydefinedspeeduptobethesequentialexecutiontimedividedbytheparallelexecutiontime.Wejustpluginthenumeratorandthedenominatorfromthepreviousslide.Sincethevalueinthedenominatorisalowerbound,thequotient(i.e.,thespeedup)isanupperbound.7-2首先我们介绍一下为什么要使用gpu进行计算，gpu计算比传统的cpu计算好在哪里。然后是gpu的基本架构，然后介绍一下有哪些常用的调用gpu的方法。后两部分使我们的重点，首先介绍一下cuda的编程模型，这一章有4个小节，讲到这的时候再细说。然后就是怎么用cuda编程，我们在这里介绍一些基本的编程方法，还有一些比较高级的用法，如果有时间的话就将，没时间的话，大家看一下文档。7-3三个整数数组，我们想要计算A[i]+B[i]然后将结果存入C[i]中，我们首先看一下传统的编程方法7-4这是一个传统的做法，需要迭代N次，时间复杂度O(N)。在这个循环中，我们发现，A[i]+B[i]和A[i-1]+B[i-1]是没有关系的。这就意味着这些运算是可以并行处理的。这时候我们就可以用多线程。7-6一些大型的应用程序，例如计算化学，用来计算分子间是如何相互作用的，还有天气和气候的模拟程序，等等一些其它的大型程序，它们要处理的数据量都是百万，上亿级别的数据。需要数千个线程同时执行，这已经远远超过了cpu的负载能力。而且操作系统也不支持。我们需要一个专门的硬件来处理这种大规模的数据。而gpu正是一个非常好的选择。因为gpu本来是为游戏设计的，而游戏中最多的操作就是图像的矩阵运算，而且是大规模的矩阵运算，和我们要处理的数据非常类似，因此gpu也就有了另外一个功能，用于高性能计算。英伟达也开发了cuda用于gpu编程。下面介绍一下gpu的基本架构。7-8一个cpu可以有2个，4个，8个，甚至更多个核，cpu可以做任何类型的计算，串行的，并行的，各种IO操作，能够进行指令预取，指令流水，分支预测，乱序执行等等一些其它的功能，在同样的芯片面积上，gpu去掉了一些cpu中的功能，或者简化了一些cpu中的功能，取而代之的是更多的计算核心，一个gpu一般都含有几千个核，能够同时运行10000个线程，因此，GPU是专门用来处理大规模的计算密集型程序。这些程序能够高度的并行化。运行在gpu上会获得非常高的加速比。7-9这个是英伟达特斯拉显卡核心GP100的架构图，这个核心由6个GPC组成，这些概念不重要，了解一下就行。每个GPC包含5个TPC，每个TPC包含2个SM单元，这里面最重要的就是SM概念。整个gp100核心包含了60个SM单元，这里展示的是一个完整的gp100核心，总共包含60个SM单元，不同的产品可能会有不同个数的SM单元。对于gp100核心的产品，最多是60个SM单元。下面详细介绍一下SM单元。7-10每一个sm单元被分成了两个处理器块，每个处理器块包含了32个单精度CUDA核（图中绿色方块）、一个指令缓存、一个warp调度（后面会介绍warp的概念）、两个指令分发单元8个ld/st单元，用来计算访存地址，8个sfu单元，specialfunctionunit，特殊函数单元，例如cos，sin，平方根等等。合起来一个SM单元里总共包含64个单精度浮点cuda核心，因此，也就有32个双精度浮点cuda核心。整个gp100核心总过有3840个单精度和1920个双精度cuda核心。一个gp100核心，它的双精度浮点的峰值运算速度为5300GFLOPs。以上是gp100核心的一个简单介绍，更详细的架构信息可以再官方的白皮书中找到，在百度上搜pascal-architecture-whitepaper7-11下面我们介绍一下三种不同的调用gpu的方法7-12这是三种常用的调用gpu的方法，cudaoptimizedlibraries,openacc,programminglanguages，下面我们详细的介绍一下7-13为了让gpu能够被广泛的应用，英伟达用cuda重新编写了一些常用的库函数，例如BLAS(BasicLinerAlgebraSubprograms基本线性代数程序,FFTW(FastFourierTransformintheWest)快速傅里叶变换程序，等等一些其它的数学函数库，虽然英伟达重写了这些函数库，但是他们没有修改这些函数的调用接口，因此我们不需要要修改源程序，只需要在编译程序时告诉编译器使用cuda数学函数库，而不是普通的数学函数库。这样我们编译出来的程序就能调用gpu了。这种方法是最简单的，但同时也是限制最多的，因为很多时候我们要做的计算不是一个数学函数就能解决的。所以这种方法非常有局限性，但是如果你的程序中需要用到数学函数，可以考虑使用cuda的数学函数库。这个不是我们的重点。7-14openacc是一个基于预编译指令的编程模型，我们需要在代码中插入一些openacc的预编译指令来指导编译器进行并行化处理，因此，我们需要专门的编译器，openacc来编译程序。看一下左边的例子，源程序就是一个简单的for循环结构。我们如果想让这个循环并行执行呢，就需要插入图中的那些预编译指令，openacc会自动的对这些结构进行并行化处理。这种方法比较实用于那些已经完成的代码，如果用cuda将这些代码重写一遍，会非常的耗时耗力，那么通过openacc这种自动化的并行处理，能够省去很多的人力物力。但是这种方法只适用于一些简单的结构，对于复杂的结构，例如dowhile循环，openacc就无能为力了，只能靠人来进行手工并行化了。这个也不是我们的重点。7-15cuda支持的编程语言很多，第一个是cuda对Fortran语言的扩展，虽然Fortran语言很古老，比你们熟悉的C语言还要早10几年，但是在科学计算领域，fortran一直都被广泛地应用着。第二个是对脚本语言python的扩展，在科学计算领域应用的比较少，第三个是opencl，opencl是为异构编程而设计的，什么是异构编程呢，假如我们有多个计算设备，例如cpu，英伟达的gpu，或者是amd的gpu，如何将这些设备统一起来呢，opencl便提供了一个标准，这个标准规定了软硬件api的规范，但它不提供具体的实现，每个厂商根据这个规范来编写具体的实现代码。程序员只需要调用统一的api就行，不同的设备会调用不同的实现来完成。第四个是cuda对C和C++的扩展。我们后面只讲如何利用c和c++进行cuda编程。其它的有兴趣的可以自己去看。7-17在这一章，我们会介绍cudacc++的编程模型。首先我们介绍一下cuda编程时需要用到哪些软硬件环境，然后在介绍cuda程序的执行过程，cpu和gpu之间是如何协同工作的。然后我们详细讲一下cuda中线程的概念，以及和线程相关的一些其它重要的概念，这些都是和后面的编程紧密相关的。在这一章的最后，我们介绍一下cuda的内存模型，访存的速度一直都是程序的瓶颈，要想写出高效的程序，就一定要了解设备的存储结构，充分的利用cache，提高程序的效率。7-18要想利用cuda进行编程，首先，你要有一块英伟达的显卡，AMD的不行，cuda不支持。这块显卡可以是专门的gpu计算显卡，例如我们前面介绍的TeslaPascalGP100，不过这种专门计算的显卡一般都非常贵，普通用户是买不起的，普通用户买来也没什么用，这种gpu计算卡一般都是给高性能计算的集群使用的。但这并不意味着我们不能用cuda编程了，普通用户可以利用英伟达的游戏显卡来进行cuda编程，例如我们笔记本里面的GT系列的显卡，还有更好的GTX系列的显卡都可以用来进行gpu编程，但是一些太老的显卡呢，就不行了。如果你想看一下你的显卡是否支持cuda编程，可以上这个网站，然后点击cuda-enabledgeforceproducts，如果在里面找到了你的显卡型号，那么你的显卡就可以用来cuda编程。7-19利用cuda编程，我们需要安装cudatoolkit，这个软件包会为我们安装一些cuda程序编译，运行会用到的库，它还会更新显卡驱动，以便能够运行cuda程序。如果你的操作系统是windows的话，要想编译cuda程序，你可能还要安装visualstudio，vs是对cuda支持最好的ide。如果你不想安装vs，你也可以直接从命令行调用cuda的编译器nvcc。Nvcc.exe可能没在你默认的系统路径里面，你需要自己找到nvcc.exe的位置，一般在它的安装目录下都能找到。如果是在linux平台的话，安装完cudatoolkit之后，可以在终端里直接敲命令nvcc来编译cuda程序。7-21一个cuda程序主要是由两部分组成，串行部分和并行部分。串行部分在cpu上执行，并行部分在gpu上执行。大多数时候，我们称cpu为host，gpu为device。串行部分主要进行逻辑控制，例如if语句，switch语句，还有输入输出，包括读取数据，向屏幕显示信息。并行部分最主要的工作就是是计算。例如矩阵相乘。Cuda程序在运行的过程中，普通指令就在cpu上执行，当遇到cuda指令时，便将代码和数据发送到gpu上，然后调用gpu进行进行计算。Gpu计算完成后，返回到cpu中，cpu继续执行。这个过程和普通的函数调用没有什么区别，只不过被调用的函数是在gpu上执行的。这是一个大致的cuda程序的执行过程，下面我们看一下具体是如何工作的。7-22当cpu要调用gpu进行计算时，首先要将程序和数据拷贝到gpu的内存中。因为Gpu在进行计算时，需要从自己的内存中读数据。因为gpu是通过pcie接口和cpu相连的，如果gpu直接从cpu的内存中读数据，需要不断的通过pcie来进行数据传输。这个过程消耗的时间要远大于我们直接将数据拷贝到gpu内存的时间。而且也会造成gpu要等待数据这样的时间浪费。因此，我们需要现将数据通过pcie拷贝到gpu的内存中，然后进行计算。7-23准备好数据之后，cpu发送指令，命令gpu开始进行计算，计算时，gpu就不需要和cpu进行通信了。直接在自己的内存中读写数据。7-24Gpu计算完之后，便将结果从自己的内存中拷贝到cpu内存中。具体的执行过程会在后面体现出来。下面我们进入cuda编程中最重要的一部分，cudathreads。7-25与传统的cpu线程相比，cuda线程更加的轻量化，能够快速的创建几千个线程，并且能够快速地进行上下文切换，因此，尽量让cuda线程做计算工作，让cpu做逻辑工作。7-26在介绍cudathreads之前，先普及以下cudakernel地概念.程序中可以并行执行地部分成为cudakernel，一般情况下，cudakernel都是一个函数，函数里面时可以并行执行地代码。Cpu调用这个函数，函数里面地cudaapi便在gpu上执行。Gpu上所有地线程都执行相同地代码，但是可以选择不同地路径，比如遇到分支语句，可能偶数线程和奇数线程执行地时不同地代码。和cpu线程一样，每个cuda线程都有一个唯一地标识。右图是一个简单的cudakernel例子。我们有四个线程，threadIdx.x表示每个线程地id，首先在输入数组中取出数据，0号线程取得第0号数据，1号线程取得第1号数据，依次类推。每个线程调用func函数来处理数据，然后将结果存到输出数组的相应位置。这是一个简单的例子，我们只用到了四个进程，真正运行在gpu上的程序都会用到几百个，上千个线程。为了高效地管理这些线程，Cuda采用了层次化的管理方式。7-27首先将线程分为warp，每32个线程分为一个warp，一个warp是cuda任务调度，程序执行的最小单元。AwarpinCUDA,then,isagroupof32threads,whichistheminimumsizeofthedataprocessedinSIMDfashionbyaCUDAmultiprocessor.但是在编程时，我们不需要考虑warp。我们需要考虑的是blocks。一个或者多个warps构成了一个block7-28程序执行时，至少使用一个warp，既32个线程。这是cuda里最小的调度和执行单位。如果你的程序连32线程都不到，那么cpu完全能满足你的需求，不需要将程序迁移到gpu上，这样反而会使程序运行变慢。在一个warp里面的线程的线程号是连续递增的。但是我们编程时并不会直接操作warp，我们操作的是block。7-29一个block内的线程可以通过共享内存来进行数据交换和同步操作，也可以通过调用同步api来进行同步。相比之下，通过共享内存来同步会更快一些。由于硬件资源的限制呢，一个线程块最多可以包含1024个线程每个线程块内的新城id都是从0开始的。并且在这个线程块内是唯一的7-30一个或多个block构成了一个grid。具体包含多少个block，是在编程时指定的，每个block里包含多少个线程也是在编程时指定的，但必须是32的整数倍。因为block是由warp组成的。当我们提交一个kernel到gpu上运行时，需要指定一个grid里包含多少个block，一个block里包含多少个thread。然后整个kernel被当作一个grid载入到gpu上运行。好了，到现在我们已经了解了cudakernel，线程，线程块，grid，这些概念了，那么接下来我们看一个更加具体的kernel的执行过程。7-31当我们向gpu提交了一个kernel时，这个kernel被看做是一个grid，假设我们指定Grid里面有8个block，如果你的GPU含有两个SM单元，向SM单元分配任务时，都是直接将一个block分配给一个SM单元。那么在这个例子里面，每个SM单元会被分配4个block，每个SM单元中的block顺序执行。而不同SM单元中的block并行执行。Block0和block1就是并行执行的。7-32如果我们有4个SM单元，每个SM单元执行两个block，每个SM单元内的block顺序执行。我们知道block是由warp组成的，那么这些warp是如何执行的呢？7-33先回忆一下SM单元的结构。一个SM单元里包含了两个warp调度器，每个warp调度器又有两个指令分发单元，每个warp调度器还对应着32个单精度的浮点cuda核。当一个SM单元被分配了一个block执行时，7-34当一个block在SM单元上运行时，它首先被分解成多个warp，然后warp调度器来决定哪个warp可以执行。一个SM单元每次可以同时执行两个warp，当一个warp被调度运行时，warp中的所有线程都会执行同一个指令，如果执行到一条分支指令的话，分支指令的所有路径都会被顺序执行，不符合当前路径的线程会处于阻塞状态。我们看一下下面的例子，7-36如果我们的kernel中有一段代码，用来区别偶数进程和奇数进程。那么这段代码会在block中差生两个执行路径，一个用来执行奇数进程，一个用来执行偶数进程，如下图所示7-37假设warp里面有8个线程，这8个线程执行上面的代码，从开始执行到分支语句之前，这8个进程每次都是执行相同的指令，但是当执行到分支语句的时候，就出现问题了，这个if语句一共有两个分支，它首先执行第一个分支，既偶数进程分支，那么这8个进程中一共有4个进程符合次分支，因此，这4个进程首先执行，而另外4个进程则等待。当这4个进程执行结束后。再执行第二个分支，既奇数进程分支，这时，另外4个处于等待的进程开始执行，而首先执行的那4个进程则进入等待状态。当奇数进程执行完后，这8个进程再同时执行后面的指令。这中情况会造成50%的性能损失。那么如何避免呢。7-38这段代码依然会在block产生两个分支，但是与前一个不同的是，同一个warp内的线程，它们会选择相同的路径，这样在一个warp内，就不产生branchdivergence了。当然了，实际情况可能不会这么简单，我们可能还需要改变原有的数据结构和算法，这些都需要根据实际情况来定。但是我们遵循的一个原则是，让warp内的线程做同样的工作。以上便是cuda线程的所有内容了。下面我们介绍cuda的内存模型7-39与传统的cpu线程相比，cuda线程更加的轻量化，能够快速的创建几千个线程，并且能够快速地进行上下文切换，因此，尽量让cuda线程做计算工作，让cpu做逻辑工作。7-40GPU的全局内存是GPU的主要的存储器，之所以是全局的，主要是因为GPU与CPU都可以对它进行写操作。任何设备都可以通过PCI-E总线对其进行访问。全局内存的功能类似于C语言程序中的堆。cudaMalloc()hastwoparameters:AddressofapointertotheallocatedobjectSizeofallocatedobjectintermsofbytesTheaddressofthepointervariableshouldbecastto(void**)becausethefunctionexpectsagenericpointer;thememoryallocationfunctionisagenericfunctionthatisnotrestrictedtoanyparticulartypeofobjectscudaFree()hasoneparameter:Pointertofreedobject7-41Oncethehostcodehasallocateddevicememoryforthedataobjects,itcanrequestthatdatabetransferredfromhosttodevice.ThisisaccomplishedbycallingoneoftheCUDAAPIfunctions,cudaMemory().PleasenotecudaMemcpycurrentlycannotbeusedtocopybetweendifferentGPU’sinmultiGPUsystems右图是CPU和GPU之间传输关系图，可以看出来，CPU和GPU之前传输速度相对很差。GPU和GPUmemory传输速度要快得多，所以，对于编程来说，要时刻考虑减少CPU和GPU之间的数据传输。7-42常量内存其实是全局内存的一种虚拟地址形式，并没有特殊保留的常量内存块。常量内存有几个特性，第一个是高速缓存，第二个时只读，第三个是它支持将单个值广播到线程束中的每个线程。常量内存的大小比较小，一般被限制为64KB。常量内存的声明方式有两种，一种是在编译时声明，需要用到“__constant__”关键字；另一种是在运行时通过主机端定义为只读内存，使用cudaMemcpyToSymbol函数。7-43textureMemory驻留在deviceMemory中，并且使用一个只读cache（per-SM）。textureMemory实际上也是globalMemory在一块，但是他有自己专有的只读cache。这个cache在浮点运算很有用。textureMemory是针对2D或3D空间局部性的优化策略，所以thread要获取2D或3D数据就可以使用textureMemory来达到很高的性能.Globalmemory没有Cache，访问速度很慢，Sharedmemory访问速度很快，但是容量很小，对于较大的数组，将其绑定至texturememory往往是个不错的选择。Texturememory可以cache，而且容量很大。7-44Sharedmemory可以用于block内线程之间的数据共享。Sharedmemory实际上是可受用户控制的一级缓存。每个SM中的一级缓存与Sharedmemory共享一个64KB的内存段。其访问速度仅次于registers，延迟较低。需要注意的是SM=Streamingmultiprocessor而不是SharedMemory。Sharedmemoryisanefficientmeansforthreadstocooperatebysharingtheirinputdataandintermediateresults.Canallocatesharedmemorystatically(sizeknownatcompiletime)ordynamically(sizenotknownuntilruntime)each“__’’consistsoftwo“_’’characters.Onecanalsoaddanoptional“__device__”infrontof“__shared__”inthedeclarationtoachievethesameeffect.7-45这里是计算强度的计算方法，所谓的计算强度，就是浮点计算次数/IO次数，就是平均每个数据所参与的浮点计算操作的次数。当计算强度>1时，说明每个数据参与超过一个浮点计算。这种情况下，就需要尽量使用sharedmemoryload，来减少访存延迟。7-46这两种内存属于不能操作的内存，是由一套自动机制来达到很好的性能。7-47这里我们对CUDA内存模型进行总结。包括每一类内存的位置、访问权限、变量的生存周期等。其中registers和localmemory是由编译器控制和分配（Non-programmable），而shared，global，constant和texturememory可受程序员控制(Programmable)。Registers和local是On-chip内存，Loal,global,constant,和texturememory是设备内存。7-48下面我们开始将cuda编程的基本语法，其实和c语言的语法一样，只是增加了一些cuda的api，7-49在这一章，我们只介绍一些基本的cuda编程的api，以及一些特殊内存的使用方法。7-50我们看第一个例子，数组求和。有两个数组，a，b，将他们对应的元素相加，然后将结果存入c中。右图所示的是一个传统的单线程的编程方法。这是一个简单串行的例子，可以看到程序中做主要工作的是add函数，并行化的工作也就集中在了add函数上。我们首先考虑一个双核的cpu，如何让两个核同时工作呢。7-51其中一个方法是，一个核处理奇数索引的数据，另一个核处理偶数索引的数据。这两个代码只是两个例子，实际写多线程代码时是不一样。如果我们需要处理的数据量太大了，需要用到几千个核，这时，我们需要将这个代码改成cuda程序了。7-52我们首先看一下main函数要如何修改。我们先说一下大致过程，cudaapi的具体信息后面再介绍。首先是定义三个数组，然后时三个指针，每个指针都以dev开头，很明显这三个指针是要在gpu上用到的。不是一定要以dev开头，只要符合c语言命名规则的都可以，以dev开头是为了和cpu上的变量区分开。下一条语句，很明显是cudaapi了，一般情况下cudaapi都会以cuda开头。这个api是用来在gpu上开辟内存空间的。A,b,c这三个数组不仅在cpu上会用到，在gpu上也会用到，因此我们也需要在gpu上开辟空间，要注意的是于c语言的malloc不同，cuda的malloc函数需要多传递一个指针的指针最为参数，具体语言后面会说到。为b，c数组开辟空间的语句直接省略了。下一条语句是初始化，初始化只需要做一次就行，既可以在cpu上完成初始化，也可以在gpu上完成初始化。如果在gpu上初始化会更快一些，但是这不是我们关注的地方。所以我们在cpu上完成初始化。在cpu上完成初始化之后，我们需要将数据拷贝到gpu中，cudamemcpy就是用来再cpu和gpu之间传递数据的，最后一个参数是用来指定传输方向的，从cpu到gpu或者从gpu到cpu。下一条语句就是告诉gpu要再gpu上执行哪个函数。我们这个例子中是add函数。我们前面讲过一个kernel被当作一个grid来执行。现在呢，add函数就是一个kernel，尖括号里的内容就是grid的大小，第一个参数N指定了一个grid有多少个block，第二个参数1指定了每个block里有几个线程。虽然我们指定了每个block里有一个线程，但是由于warp是执行和调度的最小单元，因此，每个SM单元还是会开启32个线程，只不过，32个线程只有一个做真正的任务，其余的一直处于等待状态。这里我们指定每个block里有一个线程只是一个例子，你们写真正的cuda代码时千万不要这样写，尽量让每个block里的线程数能被32整除。在下一条语句是将结果从gpu拷贝到cpu中。最后一条语句是释放gpu上的空间，b，c也要释放，我这里注释掉了。下面我们看一下这几个cudaapi的具体使用方法。7-53C语言中的malloc函数会返回一个指针，但是因为cuda的api全都会返回一个错误码，所以只能传递一个指针的指针，用来存储开辟的内存空间的地址。如果只有__device__修饰符，表明了这个函数只能在gpu上调用，并且只能运行在gpu上，如果只有一个__host__修饰符，表明这个函数只能在cpu上调用和执行。7-54在gpu上运行add函数时需要加几个尖括号，这几个尖括号被称为executionconfiguration。Dg，db，都是3维数据结构，dg指明了grid的三个维度的大小，db指明了block三个维度的大小，Ns指明了每个block使用的共享内存是多少，这是一个可选参数，默认是0。S指明了与这个kernel相关的cudastream是什么，默认是0，cudastream是比较高级的内容，我们不会讲到，感兴趣的可以自己在网上查一下。我们重点看一下dim3这个类型，我们在调用add函数时，只传递了两个整形变量，N和1.这时候，编译器会自动将这两个整型变量转换成两个dim3类型变量，分别代表grid和block的维度，grid的x维赋予N，其它两个维度都是1，block的x维度赋予1，其它两个维度都是1.如果我们不想让其它两个维度是一，可以用下面的代码。7-55声明两个dim3的变量，然后对它们的各个维度分别赋值。7-56Cudamemcpy只有一个host修饰符，意味着这个cuda函数只能在cpu端调用。它不仅能在cpu和gpu之间传递数据，还能在cpu内存上进行数据复制，也能在gpu内存上进行数据复制。只要指定传输的方向即可。下面我们看一下add函数是如何定义的。7-57我们看到add函数定义的时候多了一个__global__修饰符，global修饰符表明了这个函数要在gpu上执行，但是可以在cpu上调用。一般kernel函数都要加上__global__修饰符。使用global修饰符对函数也有一些要求。首先，global函数的返回值必须是void，其次，在调用global函数的地方必须指明它的executionconfiguration。最后，global函数的调用是异步的，意思是，一旦我们将这个函数交给gpu执行后，调用语句马上返回，不会等到gpu执行完才返回。我们看一下add函数内部。7-58Add函数里用到了一个变量blockIdx.x，这是由cuda定义的全局变量，指明了当前线程所属的block的索引，除了blockidx.x外，还有blockidx.y和blockidx.z，分别指明了x,y,z这三个维度的索引。除了blockidx还有threadidx，指明了当前thread在block内的索引。同样也有三个，分别是threadidx.x,threadidx.y,threadidx.z。在我们这个例子里面呢，grid只有一维，其它两个维度的大小都是一。而每个block里只有一个线程，所以blockidx.x的索引就相当于是线程的索引了。Add函数很简单，对于小于N的线程做加法操作。为什么写小于N呢，虽然我们可以肯定tid不会超过N，但是有可能由于我们的疏忽tid超过N了，这会在gpu上造成数组访问越界。如果你确定不会超过N，也可以不加这个判断。这一节到这里就介绍完了，下面我们讲一下如何使用gpu的共享内存和同步操作。7-62下面我们详细的讲一下这个过程，假设我们由N个线程，两个数组x，y的大小都是N*M，这两个数组都是一维数组。这N个线程每次可以同时处理N组数据，我们一共有N*M组数据，总共需要M次迭代。第一次迭代，这N个线程分别处理的是x1*y1,x2*y2一直到xN*yN。第二次迭代就是从xN+1*yN+1开始，一直到x2N*y2N。每个进程迭代M次，迭代结束后，再将每个进程的累加和加起来，得到最后的结果。下面我们看一下具体的代码怎么写。7-63首先看一下main函数的结构，这是前半部分的代码，一直到调用点乘函数。N代表了数据规模，blockspergrid表示grid里由多少个block，threadperblock表示每个block里由多少个thread。Main函数的结构呢和我们上一个数组相加的结构是一样的，不同的是这个例子中的C数组的大小并不是N，而是blockspergrid，后面会讲到为什么是这个，我们先看一下dot函数的结构7-64前面讲过，cuda会为每个block来分配一块共享内存，block里的所有线程共享这块内存，因此，在dot函数的开头，我们为block里的每一个线程申请了一个共享内存，共享内存的总大小为threadsperblock。下一步是要计算线程的全局id，threadidx表示的是当前线程在block内的索引，我们取数据时需要用到当前thread在整个grid中的索引是多少，blockdim.x表示block的x维度的大小。整个计算过程就相当于将二维数组的坐标转换成一维数组一样，我们这里就不再说了。下一个是cacheindex，表示当前线程在共享内存中的位置，因为每个block有一块共享内存，所以我们只需要threadidx来索引即可。下一步便是遍历这个线程需要处理的数据，进行累加操作。以0号线程为例，第一次处理a[0],b[0]，第二次处理a[n],b[n],间隔是总的进程数，代码里面总的进程数就是用blockdim.x*griddim.x来表示的，因为grid和block都是一维的。如果它们是二维或者三维的，还需要加上其它维度的大小。最后每个进程将累加和存入相应的共享内存中。下一步就是要对共享内存中的数据进行累加，计算出这个blcok的累加和是多少。因为只有当所有进程的数据都写入共享内存之后，我们才能开始累加操作。所以在此处要有一个同步操作。只有当所有进程都到达这个同步点的时候，所有的进程再同时执行下面的指令。执行到这里的时候，是一个多线程的环境，我们可以充分的利用多线程来累加共享内存中的数据。下面我们看一下累加操作。7-65这是共享内存cache中的数据，假设cache一共有8个数据需要累加，也就是说我们由8个线程，每个线程对着一个数据，累加的过程就是前一半的线程将自己的数据和后一半的数据相加。现在是8个数据，8个线程，第一次是迭代是前4个线程自己的4个数据和后4个数据相加，结果存入到前4个数据中。然后我们只需要累加4个数据了，不断的重复次过程，直到结果都存在了第一个数据中。这里我们要注意需要同步一下。7-66Dot函数的最后一步操作，将cache[0]的数据存入数c中，7-67然后将数据传输到cpu端，在cpu上进行最后的累加操作。以上便是共享内存的使用。7-71Image中的每一个方格都是一个像素，我们假设每个像素都能发射出一个平行的光，当一束光遇到场景中的物体时，会产生一个或多个焦点，我们找到离image最远的那个交点。7-72为了方便，我们场景中的物体全是球形物体7-73Hit函数用来计算这个表面与ox，oy这条射线有没有相交，如果相交的话，交点的深度是多少，既交点到image的距离是多少。下面我们看一下没有使用contantmemory的主程序7-74Cudaevent函数主要是用来计时的，cudaeventrecord函数里的0代表是是哪个cudastream，cudastream我们就不再讲了。还有cpubitmap结构，这是一个像素矩阵，我们直到这是一个二维矩阵就够了，这个不是cuda定义的。从代码中可以看到，我们在cpu和gpu端分别定义了20个球体。For循环是对这20个球体进行初始化。7-75然后将初始化后的球体拷贝到gpu端。下面就是计算每个像素点与每个物体的交点了，kernel函数我们就不再展示了。我们看红色框里的内容也是用来计时的，它和前面的计时的相对应，我们把这个用法记住就行了。这是普通的写法，我们知道contantmemory是用来存储常量的。那么程序里面哪些是不会改变的呢，球体这个数组，一旦我们初始化之后，球体数组就不再改变了。下面我们看一下怎么修改主程序，kernel函数不需要修改。7-76左边的是修改后的代码，右边是原来的代码，我们来看一下不同之处，我们在开头多了一个constant的声明，这表明这个球体数组存储在了constantmemory上。并且constantmemory必须声明维全局变量。Constantmemory的初始化还和以前一样，当然了，为了更快的初始化，我们也可以在gpu上完成初始化。将constantmemory拷贝到gpu上时，要用特殊的拷贝函数，用法和传统的cudamemcpy一样，只不过cudamemcpytosymbol不需要指明拷贝的方向，因为只能从cpu端拷贝到gpu端。其它的没有任何区别了，以上便是contantmemory的用法了。7-80我们用一个heatingmodel来说明如何利用texturememory7-81不同的颜色