CUDA优化策略

CUDA程序优化

CUDA程序优化应该考虑的点：

精度：只在关键步骤使用双精度，其他部分仍然使用单精度浮点以获得指令吞吐量和精度的平衡；

           延迟：需要首先缓冲一部分数据，缓冲的大小应该可以保证每个内核程序处理的一批数据能够让GPU慢负荷工作；

           计算量：计算量太小的程序使用CUDA很不合算；当需要计算的问题的计算密集度很低的时候，执行计算的时间远远比IO花费的时间短，整个程序的瓶颈出现在PCI-E带宽上。

优秀的CUDA程序特征：

在给定的数据规模下，选用算法的计算复杂度不明显高于最优的算法；

Active warp的数量能够让SM满载，并且active block的数量大于2，能够有效地隐藏访存延迟；

                当瓶颈出现在运算指令时，指令流的效率已经过了充分优化；

         当瓶颈出现在访问IO时，程序已经选用了恰当的存储器来储存数据，并使用了适当的存储器访问方式，以获得最大带宽；

CUDA的编写与优化需要解决的问题：

         确定任务中的串行和并行的部分，选择合适的算法；

               按照算法确定数据和任务的划分方式，将每个需要实现的步骤映射为一个满足CUDA两层并行模型的内核函数，让每个SM上至少有6个活动warp和至少2个活动block；

        在精度不组或者发生一处时必须使用双精度浮点或者更长的整数类型；

        优化显存访问：合并采用相同block和grid的kernel；尽力避免将线程私有变量分配到local memory；

        优化指令流：在误差可接受的情况下，使用CUDA算术指令集中的快速指令；避免多余的同步；在只需要少量线程进行操作的情况下，使用类似“if threaded<N”的方式，避免多个线程同时运行占用更长时间或者产生错误结果；

        资源均衡：调整每个线程处理的数据量，shared memory和register和使用量；通过调整block大小，修改算法和指令以及动态分配shared memory，都可以提高shred的使用效率；register的多少是由内核程序中使用寄存器最多的时刻的用量决定的，因此减小register的使用相对困难；

        节约register方法：使用shared memory存储变量；使用括号明确地表示每个变量的生存周期；使用占用寄存器较小的等效指令代替原有指令；

         与主机通信优化：尽量减少CPU与GPU间的传输：使用cudaMallocHost分配主机端存储器，可以获得更大带宽；一次缓存较多的数据后再一次传输，可以获得较高的贷款；需要将结果显示到屏幕的时候，直接使用与图形学API互操作的功能；使用流和异步处理隐藏与主机的通信时间；使用zero-memory技术和Write-Combined memory提高可用带宽；

测量程序运行时间：

        CUDA内核程序的运行时间：可以在设备端测量，也可以在主机端测量；

        CUDA API的运行时间：只能在主机端测量；使用CUDA runtime API时，会在第一次调用runtime API函数时启动CUDA环境，计时的时候应该避免将这一部分计入，因此在正式测试之前应当首先及你选哪个一侧包含数据输入输出地就爱上你，使得GPU从平时的节能模式进入工作状态，使得测试结果更加可靠；

设备端测量时间：

      调用clock（）函数：返回的是GPU的时钟周期，需要除以GPU的运行频率才能得到以秒为单位的时间；

使用CUDA API事件管理功能；

主机端测量时间：使用c标准库中的clock_t()函数测试，由于其精度很低，因此应该运行多次然后求平均运行时间；注意异步函数（比如内核函数和带有asyn后缀的存储器拷贝函数），在GPU上执行完成之前，CPU线程已经得到了它的返回值；从主机测量一系列CUDA调用需要的时间的时候，要首先调用cudaThreadSynchronize()函数等，使得GPU线程执行完毕后，进入CPU线程，从而得到正确的执行效果；在一串流中的第一个流（ID为0的流）的行为总是同步的，因此使用这些函数对0号流进行测时，得到的记过是可靠的。

任务划分原则：

        在两次主机—设备通信之间进行尽量多的计算；考虑使用流运算隐藏主机—设备通信时间，通过Pinned memory、zero—copy、write—combined memory等手段提高实际传输带宽；

        尽量使得每个block中线程数量是32的整数倍，最好保持在64~256之间，并根据任务的具体情况确定每个维度上的大小，以减少计算访存地址时的整数除法和求模运算；

              对一个block的任务进行划分后，再按照block的维度和尺寸要求对grid进行划分：每个block 的访存均匀分布在显存的各个分区中；block 间的负载可以存在一定程度的不均衡；

Grid和Block的维度设计：

        首先考虑block的尺寸，grid的尺寸一般越大越好；

        每个SM中至少要有6个active warp用于隐藏流水线延迟，并且拥有至少2个active block；

计算每个SM上active warp和active block的数量：

确定每个SM使用的资源数量：使用nvcc的—keep编译选项，或者在.cu编译规则（cuda build rule）中选择保留中间文件，得到.cubin文件，用写字板打开后可以看到imem 和reg分别代表内核函数中每个线程使用的local memory和register数量；

根据硬件确定SM上的可用资源：可以用SDK中的deviceQuery 获得每个SM中的资源；根据内核不同，SM上的warp总数上限，block总数上限，寄存器数量，shared memory数量都不同；

每个block中的线程数量不能超过512；

计算每个block使用的资源，并确定active block和active warp数量：

     e.g. 每个block中有64个线程，每个block使用256 Byte shared memory，8个寄存器文件，

那么：每个人block使用的shared memory： 256 Byte；

      每个block使用的寄存器文件数量： 8*64 = 512；

      每个block中使用的warp数量：64/32 = 2;

                 如果在G80/92 GPU中运行这个内核程序：

                        由shared memory数量限制的active block数量： 16384、256  = 64；

                                           由寄存器数量限制active block数量：8192/512 = 16;

                        由warp数量限制的active block数量 24/2 = 12;

                        每个SM中的最大active block数量：8；

这些计算可以由NVIDIA在CUDA SDK中提供的 CUDA occupancy calculator完成；

   Block 的维度和每个维度上的尺寸的主要作用是避免做整数除法和求模运算，对执行单元效率没有什么显著影响；

计算grid中各个维度上block的数量：grid在x轴上的block数量 = （问题在x轴上的尺寸+每个block在x轴上的尺寸-1）/每个block在x轴上的尺寸；

 存储器访问优化：    

 主机—设备通信优化：

      目前一条PCI—E 2.0*16通道的理论带宽是每向8GB/s, 远小于显存和GPU片内存储器带宽；

      Pinned memory：强制让操作系统在物理内存中完成内存申请和释放工作，不用参会页交换，因此速度比pageable memory快；

                         声明这些内存会占用操作系统的可用内存，可能会影响到操作系统运行需要的物理内存；

                         需要合理规划CPU和GPU各自使用的内存，使整个系统达到最优；

异步执行：

      内核启动和显存内的数据拷贝（Device to Device）总是异步的；

      内存和显存间的数据拷贝函数有异步和同步两个版本：

              同步(顺序执行)： cudaMemcpy(a_d,a_h,size,cudaMemcpyHostToDevice);

                               cpuFunction();

              异步(同时执行)： cudaMemcpyAsync（…………）；

                                               cpuFunction()；

      属于同一个流中的内核启动总是同步的；

      如果几次内核启动属于不同的流，那么他们的执行可能是乱序的；

利用异步提高计算效率：

      使用流和异步是CPU和GPU同事进行运算；

      利用不同流之间的异步执行，使流之间的传输和运算能够同时执行，更好地利用GPU资源；

全局存储器访问优化：

       需要考虑half-warp访问的对齐问题，不同的硬件要求不同；（存疑？？？？？？？？）

       采用合并访问；

       尽量避免间隔访问：比如按列访问矩阵，可以借助shared memory来实现这一点；

Shared memory访问优化：

       共享存储器被组织为16个可以被同时访问的存贮器模块，称为bank；

Bank组织方式：宽度32bit，相邻的32bit字被组织在相邻的bank中，每个bank在每个时钟周期可以提供32bit的带宽；

一个warp被分为两个half-warp进行访问；

避免bank conflict：在SDK中，使用宽度为17或则会threadDim.x+1的行来避免bank conflict；（存疑？？？？？？？？）

Shared memory采用了广播机制：在相应一个对同一个地址的读请求时，一个32bit字可以被读取并同时广播给不同的线程；

当一个half-warp中有多个线程读取同一个32bit字地址中的数据时，可以减少bank conflict的数量；

如果half-warp中的线程全都读取同一地址中的数据时，此时完全不会发生bank conflict；

如果half-warp内有多个线程要对同一地址进行读写操作，此时会产生不确定结果，这种情况应该使用shared memory的原子操作；

共享存储器保存着加载kernel时传递过来的参数，以及kernel执行配置参数，如果参数列表很长，应该将其中一部分参数放入constant memory；

使用纹理存储器：

        主要用于存放图像和查找表：不用严格遵守合并访问条件，就能达到较高带宽；

                                  对于少量数据的随机访问，效率不会太差；

                                  可以使用线性滤波和自动类型转换等功能调用硬件的不可编程计算资源，不必占用可编程计算单元；

使用常数存储器：

        主要用于存放指令中的常数；速度低于shared memory；

指令流优化：

增大吞吐量手段：

避免使用地吞吐量指令；

       优化每种类型的存储器，有效利用带宽；

       允许线程调度单元精良用多的数学计算来覆盖访存延迟，需要有教导的算术密度；

吞吐量：每个多处理器在一个时钟周期下执行的操作数目；

算术指令：尽量使用单精度浮点单元进行运算，在计算能力小于等于1.2的设备中，每个双精度的变量将会转换成单精度格式，双精度运算也会转为单精度算术运算；

          单精度浮点基本算术运算：加，乘，乘加运算的吞吐量是每个时钟周期8个操作；

          求导数运算：每个时钟周期2个操作；

          单精度除法：每个时钟周期0.88个操作；

          单精度浮点倒数平方根：2；

          平方根：1；

          对数：2；

          正弦余弦：参数较大的时候，采用归约操作将x的绝对值减小；有快路径和慢路径（大参数）；

          整数算术运算：整数加法（8），乘（2）；除法和取模开销特别大，尽量地避免或者用位运算代替；

          比较，min，max：（8）；

          位运算（8）；

          类型转换（8）；

控制流指令： If, switch, do, for, while 可能引起一个warp线程跳转到不同的分支，严重影响指令吞吐量；

访存指令：包括任何读写memory的指令；

对于local memory只有在register不够用或者编译器无法解析的时候才会发生；

将较大的数据（float，double）拆分成每个线程32bit，或者将多个[u]char,[u]short合并成每个线程32bit的形式访问；

在访问local/global memory时候，会有额外的400~600个时钟周期的访问延迟；

同步指令：_syncthreads()的吞吐量是每时钟周期8个操作

1. memory coalescing，保证内存融合。因为global memory在CC为1.x上是按照half wrap进行访问读写的，而在2.x上是按照wrap进行访问读写的。在显存中，有多个存储器控制器，负责对显存的读写，因此，一定要注意存储器控制器的负载均衡问题。每一个存储器控制器所控制的那片显存中的地址空间称为一个分区。连续的256Byte数据位于同一个分区，相邻的另一组256Byte数据位于另一个分区。访问global memory就是要让所有的分区同时工作。合并访问就是要求同一half-wrap中的thread按照一定byte长度访问对齐的段。在1.0和1.1上，half-wrap中的第k个thread必须访问段里的第k个字，并且half-wrap访问的首地址必须是字长的16倍，这是因为1.0和1.1按照half-wrap进行访问global memory，如果访问的是32bit字，比如说一个float，那么half-wrap总共访问就需要16个float长，因此，每个half-wrap的访问首地址必须是字长的16倍。1.0和1.x只支持对32bit、64bit和128bit的合并访问，如果不能合并访问，就会串行16次。1.2和1.3改进了1.0和1.1的访问要求，引进了断长的概念，与1.0和1.1上的端对齐长度概念不同，支持8bit-段长32Byte、16bit-段长64Byte、32bit-64bit-128bit-段长128Byte的合并访问。对1.2和1.3而言，只要half-wrap访问的数据在同一段中，就是合并访问，不再像1.0和1.1那样，非要按照顺序一次访问才算合并访问。如果访问的数据首地址没有按照段长对齐，那么half-wrap的数据访问会分两次进行访问，多访问的数据会被丢弃掉。所以，下面的情况就很容易理解：对1.0和1.1，如果thread的ID与访问的数据地址不是顺序对应的，而是存在交叉访问，即：没有与段对齐，那么，就会16次串行访问，而对1.2和1.3来讲，会判断这half-wrap所访问的数据是不是在同一个128Byte的段上，如果是，则一次访问即可，否则，如果half-wrap访问地址连续，但横跨两个128Byte，则会产生两次 传输，一个64Byte，一个32Byte。当然，有时还要考虑wrap的ID的奇偶性。1.2和1.3放宽了对合并访问的条件，最快的情况下的带宽是最好的情况下的带宽的1/2，然而，如果half-wrap中的连续thread访问的显存地址相互间有一定的间隔时，性能就会灰常差。比如，half-wrap按列访问矩阵元素，如果thread的id访问2*id的地址空间数据，那么，半个wrap访问的数据刚好是128Byte，一次访问可以搞定，但是，有一半数据会丢失，所以，也表示浪费了带宽，这一点一定要注意。如果不是2倍，而是3倍、4倍，那么，有效带宽继续下降。在程序优化时，可以使用share memory来避免间隔访问显存。

　　2. bank conflict，bank冲突。先说一下，share memory在没有bank conflict情况下，访问速度是global和local的100倍呢，你懂的。类似global memory的分区，share memory进行了bank划分。如果half-wrap内的很多thread同时要求访问同一个bank，那么就是bank conflict，这时，硬件就会将这些访问请求划分为独立的请求，然后再执行访问。但是，如果half-wrap内所有thread都访问同一个bank，那么会产生一次broadcast广播，只需要一次就可以相应所有访问的请求。每个bank宽度长度为32bit。对于1.x来讲，一个SM中的share memory被划分为16个bank，而2.x是32个bank。1.x的bank conflict和2.x的bank conflict是不一样的。对1.x来讲，多个thread访问同一个bank，就会出现bank conflict，half-wrap内所有thread访问同一个bank除外。但是，对2.x来讲，多个thread访问同一个bank已经不再是bank conflict了。比如：
　　__shared__ char Sdata[32];
　　char data = Sdata[BaseIndex+tid];
　　在1.x上属于bank conflict，因为，0~3thread访问同一个bank，4~7访问同一个bank，类推，这种情况属于4-way bank conflict。但是，对于2.x来讲，这种情况已经不是bank conflict了，以为2.x采用了broadcast机制，牛吧，哈哈。 这里要多看看矩阵乘积和矩阵转置例子中的share memory的使用，如何保证memory coalescing和避免bank conflict的。

　　3. texture memory是有cache的，但是，如果同一个wrap内的thread的访问地址很近的话，那么性能更高。

　　4.以下是要注意的：
　　(1)在2.x的CC上，L1 cache比texture cache具有更高的数据带宽。所以，看着使用哈。
　　(2)对global memory的访问，1.0和1.1的设备，容易造成memory uncoalescing，而1.2和1.3的设备，容易造成bandwidth waste。 而对2.x的设备而言，相比1.2和1.3，除了多了L1 cache，没有其他的特别之处。
　　(3）采用-maxrregcount=N阻止complier分配过多的register。
　　(4)occupancy是每个multiprocessor中active wrap的数目与可能active wrap的最大数目的比值。higher occupancy并不意味着higher performance，因为毕竟有一个点，超过这个点，再高的occupancy也不再提高性能了。

　　 5.影响occupancy的一个因素，就是register的使用量。比如，对于1.0和1.1的device来讲，每个multiprocessor最多有8192个register，而最多的simultaneous thread个数为768个，那么对于一个multiprocessor，如果occupancy达到100%的话，每个thread最多可以分配10个register。另外，如果在1.0和1.1上，一个kernel里面的一个block有128个thread，每个thread使用register个数为12，那么，occupancy为83%，这是因为一个block有128个thread，则，由于multiprocessor里面最大的simultaneous thread为768，根据这个数目计算，最多同时有6个active block，但是6个active block，就会导致总共thread个数为128*6*12个，严重超过了8192，所以不能为6，得为5，因为128*5<768, and 128*5*12<8192, 5是满足要求的最大的整数。如果一个kernel里面的一个block有256个thread，同样一个thread用12个register，那么occupancy为66%，因为active block为2。可以在编译选项里面加入--ptxas-options=-v查看kernel中每个thread使用register的数量。同时，NV提供了CUDA_Ocuppancy_calculator.xls作为occupancy计算的辅助工具。顺便说一下，对于1.2和1.3的device来讲，每个multiprocessor最多的simultaneous thread个数为1024个。

　　6. 为了隐藏由于register dependent寄存器依赖造成的访问延迟latency，最小要保证25%的occupancy，也就是说，对于1.x的device来讲，一个multiprocessor最少得发起192个thread。对于1.0和1.1来讲， occupancy为192/768=25%，达到要求，但是对于1.2和1.3而言，192/1024=18.75%，不过，也只能这样。对于2.x系列的device来讲，由于是dual-issue，一个multiprocessor最多发起simultaneous thread个数为1536个，所以，一个multiprocessor最少同时发起384个thread时，occupancy为384/1536=25%，又达到了25%。

　　7. 对于block和thread的分配问题，有这么一个技巧，每个block里面的thread个数最好是32的倍数，因为，这样可以让计算效率更高，促进memory coalescing。其实，每个grid里面block的dimension维度和size数量，以及每个block里面的thread的dimension维度和size数量，都是很重要的。维度呢，采用合适的维度，可以更方便的将并行问题映射到CUDA架构上，但是，对性能不会有太大改进。所以，size才是最重要的，记住叻! 其实，访问延迟latency和occupancy占有率，都依赖于每个multiprocessor中的active wrap的数量，而active wrap的数量，又依赖于register和share memory的使用情况。首先，grid中block的数目要大于multiprocessor的数目，以保证每个multiprocessor里面最少有一个block在执行，而且，最好有几个active block，使得blocks不要等着__syncthreads()，而是占用了hardware。其次，block里面的thread的数目也很重要。对于1.0和1.1的设备来讲，如果一个kernel里面block的大小为512个thread，那么，occupancy为512/768=66%，并且一个multiprocessor中只有一个active block，然而，如果block里面的thread为256个thread，那么，768/256=3，是整数，因此，occupancy为100%，一个multiprocessor里面有3个active block。但是，记住了，higher occupancy don't mean better performance更高的占有率并不意味着更好的性能。还是刚才那个例子，100%的occupancy并不比66%的occupancy的性能高很多，因为，更低的occupancy使得thread可以有更多的register可以使用，而不至于不够用的register分配到local memory中，降低了变量存取访问速度。一般来讲啊，只要occupancy达到了50%，再通过提高occupancy来提高性能的可能性不是很大，不如去考虑如何register和share memory的使用。保证memory coalescing和防止bank conflict。记住如下几点：
　　(1)block里面thread个数最好为wrap大小的倍数，即：32的倍数。使得计算效率更高，保证memory coalescing。
　　(2)如果multiprocessor中有多个active block时，每个block里面的thread个数最好为64的倍数。
　　(3)当选择不同的block大小时，可以先确定block里面thread个数为128到256之间，然后再调整grid中block大小。
　　(4)如果是让问延迟latency造成程序性能下降时，考虑在一个block里面采用小block划分，不要在一个multiprocessor中分配一个很大的block，尽量分配好几个比较小的block，特别是程序中使用了__syncthreads()，这个函数是保证block里面所有wrap到这里集合，所以，block里面的thread越少越好，最好是一个wrap或者两个wrap，这样就可以减少__syncthreads()造成的访问延迟。
　　(5)如果如果一个block里面分配的register超过了multiprocessor的最大极限时，kernel的launch就会fail。

　　8. share memory的使用量也是影响occupancy的一个重要因子。thread与share memory的元素之间，没有必要是一对一的。一个线程可以一起负责处理share memory数组中的第一个、第二个以及第三个元素，都ok的。第一个thread处理share memory中的第一个元素，第二个thread负责处理第二个元素，类推如下，这种情况不是必须的，有时也没必要这么做。在代码里面，采用一个thread负责处理share memory数组中的多个元素的方法，是非常好的策略。这是因为如果share memory里面各个元素要进行相同的操作的话，比如乘以2，那么，这些操作可以被负责处理多个元素的一个thread一次搞定，分摊了thread处理share memory元素数量的成本费用。

　　9. 当上面那些high level级别的优化策略都检查使用过以后，就可以考虑low level级别的优化：instruction optimization指令集优化。这个也可以很好的提高性能的。指令集的优化，可以稍微总结如下：
　　(1)尽量使用shift operation位移运算来取代expensive昂贵的division除法和modulo取余运算，这里说的都是integer运算，float不行的。如果n是2幂数，(i/n)=(i>>log2(n)), (i%n)=(i&(n-1)). 其实，这只是一个量的问题，对于1.x的设备而言，如果一个kernel里面使用了十多个tens of这样的指令，就要考虑用位移运算来取代了;对于2.x的设备而言，如果一个kernel里面使用了20个这样的指令，也要考虑使用位移运算来取代除法和取余运算。其实，compiler有时会自动做这些转换的，如果n是2的幂数。
　　(2)reciprocal square root，对于平方根倒数1.0f/sqrtf(x)，编译器会采用rsqrtf(x)来取代，因为硬件做了很多优化。当然，对于double型的平方根倒数，就采用rsqrt(x)啦。呵呵，记住了。
　　(3)编译器有时会做一些指令转化。在要做计算的单精度浮点型常数后面，一定要加入f，否则，会被当做双精度浮点运算来计算，对于2.x以上的设备来讲，这一点很重要，记好了。
　　(4)如果追求速度speed，而不是精度precision，那么尽量使用fast math library。比如，__sinf(x)、__expf(x)比sinf(x)和expf(x)有更快的速度，但是，精度却差一些。如果是__sinf(x-2)则比sinf(x-2)的速度要快一个数量级，因为x-2运算用到了local memory，造成太高的访问延迟。当然，在compiler option中使用-use_fast_math可以让compiler强制将sinf(x)和expf(x)转化为__sinf(x)和__expf(x)进行计算。对于transcendental function超越函数，作用对象是单精度浮点型数据时，经常这么用，其他类型的数据，性能提升不大。
　　(5)对于2和10为底做指数运算，一定要采用exp2()或者expf2()以及exp10()或者expf10()，不要采用pow()和powf()，因为后者会消耗更多的register和instruction指令。 另外，exp2()、expf2()、exp10()、expf10()的性能和exp()以及expf()性能差不太多，当然比pow()和powf()要快10多倍呢。加好了哈。
　　(6)减少global memory的使用，尽量将global memory数据加载到share memory，再做访问。因为访问uncached的显存数据，需要400~600个clock cycle的内存延迟。

　　10. 下一个就是control flow了。一定要避免在同一个wrap里面发生different execution path。尽量减少if、swith、do、for、while等造成同一个wrap里面的thread产生diverge。因为，一旦有divergence，不同的execution path将会顺序的串行的执行一遍，严重影响了并行性。但是：
switch（threadIdx.x）
{
case 0：
break;
case 1:
break;
...
case 31:
break;
}
　　上面这个例子，则不会发生divergence，因为控制条件刚好和wrap里面的thread相对应。
　　其实，有时，compiler会采用branch predication分支预测来打开loop循环或者优化if和switch语句， 这时，wrap就不会出现divergence了。在写code时，我们也可以自己采用#pragma uroll来打开loop循环。在使用branch predication时，所有指令都将会执行，其实，只有预测正确的真正的执行了，而预测错误的，其实就是thread，不会去读取该instruction的地址和数据，也根本不会写结果。其实，编译器做分制预测，是有条件的，只有分支条件下的指令instruction的个数小于等于某个阈值的时候，才会做分支预测branch predication。如果编译器觉得可能会产生多个divergent wrap，那么阈值为7，否则为4。(这里很不理解7和4是怎么来的)。

　　11. 在loop循环的counter，尽量用signed integer，不要用unsigned integer。比如：for(i = 0; i < n; i++) {out[i] = in[offset+stride*i];} 这里呢，stride*i可以会超过32位integer的范围，如果i被声明为unsigned，那么stride*i这个溢出语句就会阻止编译器做一些优化，比如strength reduction。相反，如果声明为signed，也没有溢出语句时，编译器会对很多地方做优化。所以，loop counter尽量设置为int，而不是unsigned int。

　　12. 在1.3及其以上的device上，才支持double-precision floating-point values，即：64位双精度浮点运算。当使用double时，在编译器选项里面添加：-arch=sm_13

　　13. 还有一点需要注意，如果A、B、C都是float，那么A+(B+C)并不一定等于(A+B)+C。

　　14. 先看下面两个语句：float a; a = a * 1.02;
　　对于1.2及其以下的device来讲，或者1.3及其以上device，但是没有打开支持double运算的选项，那么，由于不支持double，所以，1.02*a这个乘积是一个float;
　　对于1.3及其以上的device来讲，如果打开了支持double运算的选项，那么，a*1.02是一个double，而将乘积赋值给a，这个结果是float，所以，是先做了从float到double的promotion扩展，然后做了从double到float的truncation截取。
　
　15. 多GPU编程。如果有p个GPU同时并行，那么，程序中就需要p个CPU threads。这些thread可以用OpenMP(小规模)或者MPI(大规模)进行管理。GPU之间的数据拷贝，必须通过CPU实现。对于OpenMP，是这样的：一个CPU thread将数据从对应的GPU中拷贝到host端的share memory region中，然后另一个CPU thread将数据从host端的share memory region拷贝到对应的GPU中。也就是说：OpenMP是通过share memory进行数据拷贝的。而对于MPI而言，数据是通过message passing进行传递的。一个CPU thread使用cudaMemcpy将数据从device拷贝到host，然后通过MPI_Sendrecv()，另一个CPU thread就使用cudaMemcpy将数据从host端拷贝到呃device端。编译选项，记着采用nvcc -Xcompiler /openmp或者nvcc -Xcompiler mpicc。