计算机系统结构-附录 第1-10章 计算机系统评价指标-专用加速器

大学计算机基础第1章计算机系统结构概述 图F.12所示。图F.12SAN结构示意图SAN具有如下优点：=1\*GB2⑴高性能、高速存取：目前Infiniband可提供100Gbps的带宽；=2\*GB2⑵高可用性：具有更多的存储空间进行冗余；=3\*GB2⑶集中存储和管理：通过整合各种不同的存储设备形成一个统一的存储池；=4\*GB2⑷高可扩展性：服务器和存储设备相分离，两者的扩展可以独立进行；=5\*GB2⑸支持更远的距离，另外通过光纤通道网卡、集线器、交换机等互联设备，用户可根据需要可以灵活地放置服务器和存储设备。在数据库等其它面向事务处理的应用场合，SAN的优势更加明显。SAN也有弱点：=1\*GB2⑴设备的互操作性较差：依赖于专用网络协议。=2\*GB2⑵多主机共享SAN需要额外协调软件。（3）iSCSIiSCSI在议TCP/IP封装SCSI存储协议，是一种IP的SAN。iSCSI采用Initiator/Target(启动设备/目标设备)的实现模式：Initiator发出iSCSI命令，而Target接收到该命令后，负责解释和执行。下面以iSCSI为例，分析实际的数据存取路径。Linux的SCSI驱动程序的实现分为三个模块：SCSI上层，SCSI中层，SCSI底层。可以用iSCSI模块替换三层模块中的任意一层来对SCSI命令进行封装并通过TCP/IP网络传输。可以通过替换Linux的SCSI驱动的中层模块来对iSCSI协议进行解释和封装。如图F.13所示，（1）当用户向一个iSCSItarget设备发出访问请求时，应用程序通过系统调用访问文件系统，（2）文件系统首先解释这个I/O命令，分析I/O命令所在的设备和地址。然后它把这个请求发送给SCSI上层驱动程序，（3）SCSI上层驱动程序把相应的I/O命令转化为SCSI命令，并发给iSCSIInitiator，（4）iSCSIInitiator把SCSI命令封装打包，通过网络发送到iSCSITarget模块，（5）当iSCSITarget模块收到iSCSI命令包后，它把iSCSI命令包中的SCSI命令发送到SCSI的底层驱动程序，（6）由SCSI底层驱动程序完成I/O请求。把数据按原路返回给用户。图F.13访问iSCSITarget设备的I/O流程（4）存储集群（StorageCluster）存储集群把多个存储节点、元数据服务器和客户服务器通过高速网络连接起来，通过元数据服务器对外提供统一命名空间。客户端从元数据服务器获取所请求文件的存储节点信息，然后直接存取相应的存储节点。这种三方合作I/O模式，使得存储集群具有良好的扩展性，能够扩展包含更多的存储节点和元数据服务器。存储集群通常使用分布式文件系统来统一管理所有存储节点（客户端、元数据服务器和存储节点）。1998年左右GoogleFileSystem（GFS）统一管理成百上千普通服务器（早期就使用普通桌面计算机），爬取和存储Internet上的网站数据，并建立检索，实现了搜索引擎。注意到存储节点也具有计算能力，能够执行本地处理，这也就是后来Hadoop分布式文件系统提供Mapreduce计算模式。同时，Lustre分布式文件系统也被开发出来，用于存储超级计算机生产数据。现在，存储集群已经广泛应用到各种数据中心，在第八章中将进一步讨论。图F.14Lustre文件系统结构图盘古存储系统是阿里云的超大规模分布式存储系统，它的一个存储集群部署了超过100000台服务器，总存储容量超过10EB。它不仅为阿里云服务，还为阿里巴巴集团和蚂蚁金服的其它独立业务，如对象存储、云计算等提供高性价比、统一的存储服务。盘古作为阿里巴巴集团的统一存储平台，必须在满足服务质量（QualityofService,QoS）的同时，最大程度的降低总成本并支持多个计算集群。盘古系统也是三方体系结构，其框架图如图2-1所示。它由应用服务器（ApplicationServers，ASs），元数据服务器（MetadataServers，MSs）和数据存储服务器（ChunkServers，CSs）组成。应用服务器，也就是存储客户端，为用户提供在线服务。通常，盘古将ASs和CSs一起划分为多个业务集群（BusinessZones），其中每个业务集群由一组ASs和CSs组成。每个业务集群提供了专一的在线服务（例如，对象存储服务），避免了多业务并置而导致潜在的硬件资源争用。类似于GFS，在盘古中，当客户端需要获取或写入数据时，ASs首先向MSs查询要访问数据所在的位置，然后直接与CSs交换数据。盘古中的每个数据文件都被划分为一系列的Chunk，分布在多个CSs上。Chunk是最小的数据管理单元，其最大大小为64MB，每个Chunk被分配了一个全局唯一的ChunkID。每个MS都有一个全局映射表，用于维护文件到Chunk以及Chunk到CSs的映射关系。MSs采用ParallelRaft协议来实现完全分布式的元数据管理，具有较高的负载均衡、可靠性和可用性。为了方便访问开源生态系统中的应用程序，盘古还兼容于HDFS。因此，它可以无缝访问海量数据并与阿里云其它存储产品进行交互。图F.15阿里盘古存储系统框架图F.6排队理论（增加）本节构建数学模型来粗略计算请求-服务模型系统的响应时间和吞吐率，这种数学模型称之为排队论。通过模型量化分析，能更好的理解和设计服务系统。本节以I/O请求处理作为说明，其分析方法和结论能够应用到其他请求-服务系统中。F.3.1排队模型图F.16将I/O系统视为黑匣子服务系统首先采用黑盒方法建模方式。如图F.16所示，I/O请求不断到达I/O服务系统，当I/O设备处理完请求之后，请求就“离开”。请求到达往往是离散和随机的，某一时刻请求达到和请求离开不一定相等。但架构师通常对系统的长期、稳定状态感兴趣，而不是瞬时情况。如果系统处于稳定状态，那么在一段时间内进入系统的任务数必须等于离开系统的任务数。这种流量平衡状态是稳定状态的前提。如果系统已经被观察或测量了足够长的时间并且平均等待时间稳定，那么可以认为服务系统达到稳定状态。这样引出Little定律，建立平均任务数、新任务的平均到达率和执行任务的平均时间之间的关系。平均任务数=任务到达率×平均响应时间Little定律适用于任何处于平衡状态的系统，只要系统内部不会创建或者删除任务，并且有不受响应时间影响的稳定任务达到。注意到达率和响应时间必须使用相同的时间单位。假设系统处于平衡状态，因此请求输入速率必然等于请求输出速率。观察这个系统Timeobserve分钟。在观察过程中，记录了每项任务的服务时间，然后把所有任务服务时间加起来。Timeobserve期间完成的任务数是Numbertask，所有任务在系统中花费的累计时间是Timeaccumulated。请注意，任务可以在执行时间上重叠，因此Timeaccumulated大于Timeobserve。该时间段内：平均任务数=Timeaccumulated/Timeobserve平均响应时间=Timeaccumulated/Numbertask任务到达率=Numbertask/Timeobserve将上面的三个定义代入这个公式，并交换结果右边的两项，就得到Little定律：系统中的平均任务数=到达率×平均响应时间。考虑到服务系统（服务器）包括一个先进先出的请求队列和服务装置。后者是具体执行请求服务的设备。对于这个服务器，给出下列定义：■服务时间Timesever——服务任务的平均时间；平均服务费率为1/服务时间，也用符号μ表示。■排队时间Timequeue——队列中的平均排队时间。■响应时间Timesystem——系统中的每个任务的平均响应时间，即排队时间和服务时间的总和。■任务到达率Arriverate——任务到达的平均数量/秒，也由符号λ表示。■服务任务数量Lengthserver——正在服务中的平均任务数。■队列长度Lengthqueue——队列的平均长度。■系统中任务数Lengthsystem——系统中的平均任务数，它是队列长度和服务长度之和。Timequeue关注任务排队的时间，而Timesystem是从进入服务装置到完成的全部时间。因此是Timesystem=Timequeue+Timesever。使用Little定律得到服务中的平均任务数（Lengthserver）等于到达率×服务时间。服务器利用率是正在服务平均任务数除以服务率。对于单个服务器，服务速率为1/服务时间。因此，服务器利用率（每个服务器的平均数任务）：服务器利用率=Arrivalrate×Timeserver服务利用率必须在0到1之间；否则服务装置将会过载，导致有任务不能被完成，这违反了系统处于平衡状态的假设。请注意，这个公式只是对Little定律的重述。利用率也称为服务强度，在许多场合中用符号ρ表示。例题假设I/O系统具有单个磁盘，每秒钟服务50个I/O请求。假设磁盘服务一个I/O请求的平均时间是10毫秒。此I/O系统的利用率是多少？解答使用上面的等式，将10毫秒表示为0.01秒。服务器利用率=到达率×服务时间=50×0.01=0.5I/O系统利用率为0.5队列把任务提交给服务器的策略称为队列规则。最简单和最常见的规则是先进先出(FIFO)。据此，可以将队列中的等待时间与队列中的平均任务数关联起来：Timequeue=Lengthqueue×Timeserver+平均任务服务剩余时间即排队时间是队列中任务数乘以均值服务时间加上平均任务服务剩余时间，后者是当新任务到达时，正在服务装置中被服务任务还需要执行的剩余时间。而当新任务到达时，队头的任务必须等待正在处理的任务完成之后，才能进入服务器处理。这部分时间也必须计算到任务的等待时间。F.3.2请求服务分布模型计算平均任务服务剩余时间并不简单，考虑达到过程和处理过程是独立无关的，一个新的任务可以随时到达，而且也没法确定正在服务的任务还需在服务器中执行多长时间。因此可以把任务到达间隔和任务服务时间看成随机变量，如果能确定这两个随机变量的分布，可以通过统计来获得一段时间内任务到达数量和平均任务服务剩余时间。随机变量可以通过实际测试，来统计实际服务/达到时间分布。具体方式可以在最大值和最小值之间分割出一系列等间距范围，称之为桶，然后统计所有服务/达到时间落在桶中的频率，作为该类请求的出现概率。除此之外，也可以通过理论性能分析模型确定I/O服务时间分布。一旦确定请求服务时间分布，就可以计算它的加权算术平均服务时间。具体而言，统计得到第i种任务的出现次数是ni，则任务i出现的频率为：则，加权算术平均服务时间=f1×T1+f2×T2+……fn×Tn+其中Ti为第i种任务的服务时间。可以进一步使用方差来描述概率分布特征，方差计算为：方差=(f1×T12+f2×T22+……fn×Tn2)-加权算术平均服务时间2可以使用方差平方系数C2，避免量纲的影响。进一步得到方差系数：如果随机变量分布是常见的指数分布，则C=1。如果是常数分布，则C=0注意到C是一个常数统计量，不受时间的影响，也称为无记忆性，此属性是用于预测行为模型的重要假设。不然，事件之间存在关联使得数学处理极其困难。最广泛使用的泊松分布用于刻画一段时间内，任务达到时间满足指数分布的不同达到任务数量的概率，其概率密度函数如下：其中a=事件​​发生率×时间段。如果到达间隔时间呈指数分布，则到达率就是事件​​发生率，则泊松过程描述时间间隔t内的任务到达数，具有泊松分布a=到达率×t。这样我们假设任务达到满足泊松分布，服务时间分布可以是任意分布，则平均剩余服务时间=1/2×算术平均值×(1+C2)。本书不会推导这个公式，当可以使用两个典型分布来分析这个公式是否合理。如果服务时间满足常数分布时，标准偏差为0，因此C为0，则平均剩余服务时间仅为服务时间的一半。也就是在常数分布时，服务装置中的任务平均执行了一半。如果服务时间分布是满足指数分布，则C为1，平均剩余服务时间等于加权算术平均时间，由于指数分布式无记忆性的，因此服务装置中中总有一个新任务开始执行。这符合我们的直觉。例题使用上面的定义和公式，推导出平均等待时间队列（Timequeue）根据平均服务时间（Timeserver）和服务器利用率。解答在为新任务提供服务之前，必须完成新任务之前的队列（Lengthqueue）中的所有任务，每个任务平均花费Timesever。如果一个任务正在服务中，需要平均剩余服务时间完成。服务器忙的概率是服务器利用率。因此，服务期望时间是服务器利用率×平均剩余服务时间，这就是前面的公式：Timequeue=（Lengthqueue）×Timesever+服务器利用率×平均剩余服务时间根据平均剩余服务时间和的Lengthqueue定义得到：Timequeue=服务器利用率×[1/2×Timesever×(1+C2)]+(到达率×Timequeue)×Timesever如果仅考虑指数分布，则C2为1，因此Timequeue=服务器利用率×Timesever+(到达率×Timequeue)×Timesever使用服务器利用率替代后一项，得到Timequeue=服务器利用率×Timesever+服务器利用率×Timequeue等式左右都有Timequeue。合并得到：Timequeue=Timesever×服务器利用率/(1-服务器利用率)在稳态情况下，Little定律能够很好地利用到黑盒服务器模型。Lengthqueue=到达率×Timequeue使用上述公式取代Timequeue。Lengthqueue=到达率×Timesever×服务器利用率/(1-服务器利用率)由于到达率×Timesever=服务器利用率，我们进一步简化得到：Lengthqueue=服务器利用率2/(1-服务器利用率)也就是排队情况和服务器利用率之间的关系。这符合我们的直觉，任务越多，排队越长。例题接上例，如果服务器利用率是0.5，请计算平均队列长度。解答使用上述公式：Lengthqueue=服务器利用率2/(1-服务器利用率)=0.5也就是队列中平均0.5个任务。如前所述，这些公式都是基于排队理论，来预测此类随机变量的行为。真实系统太复杂，无法使用排队理论来获得精确结果，因此排队论仅给出大致数量。但通过对于过去行为的测量和统计，再通过排队论使用简单估算将来事件的特征，并且预测未来的趋势，这对于计算机设计是非常重要的。F.3.3排队模型及其应用再回顾一下关于排队模型的假设：■系统处于平衡状态。■两个连续请求到达之间的时间，称为到达间隔时间，是指数分布的，它表征了前面提到的到达率。■请求源的数量不受限制。（这被称为无限队列；有限队列模型通常用于有限缓存场景，队列满时后续任务会被丢弃，这造成实际到达率随系统中等待任务数量而变化。）■服务器串行执行任务。■队列长度没有限制，遵循先进先出原则，所以必须完成所有排队的任务。■只有一台服务器。满足上述假设的队列模型称为M/M/1：第一个M是指请求到达间隔时间分布满足指数随机分布（C2=1）；第二个M指代服务时间分布也是指数分布（C2=1）；1代表只有一个服务器。M/M/1模型是一种简单且广泛使用的模型。在排队中使用指数分布的假设通常有其优势和劣势。好处是许多任意分布的叠加仍然是负指数分布。一般而言，当变化性不明确时，指数分布中等变异性(C=1)比低变异性(C=0)或高变异性（大C）更安全。劣势是，假设指数分布使得数学会更简单，但精度有限。下面给出运用排队论的例子。例题假设处理器每秒发送40个I/O请求，请求间隔满足负指数分布，磁盘处理每个I/O的时间为20ms，请问：1.磁盘的使用率；2.请求的排队时间；3.请求的响应时间。解答使用上述公式：请求到达率为40I/O每秒；I/O服务时间为20ms=0.02s服务利用率为40×0.02=0.8服务时间满足负指数分布，根据上述公式可以得到：Timequeue=Timesever×服务器利用率/(1-服务器利用率)=20ms×0.8/(1-0.8)=80ms平均响应时间为Timequeue+Timeserver=80+20ms=100ms。这表明80%的时间在队列中等待。我们进一步评估改进设计的效果。例题使用一个更快的磁盘，处理每个I/O的时间为10ms，请问：1.磁盘的使用率；2.请求的排队时间；3.请求的响应时间。解答使用上述公式：请求到达率为40I/O每秒；I/O服务时间为10ms=0.01s服务利用率为40×0.01=0.4服务时间满足负指数分布，根据上述公式可以得到：Timequeue=Timesever×服务器利用率/(1-服务器利用率)=10ms×0.4/(1-0.6)=6.7ms平均响应时间为Timequeue+Timeserver=10+6.7ms=16.7ms。把磁盘性能提高一倍，就能使得I/O响应时间提高6倍。采用增加部件数量提高整体性能是一个基本系统优化方案，因此当使用多个服务器。仍然可以使用排队论进行评估，这个模型称之为M/M/m。使用Nserver代表服务器的数量。使用率=到达率×Timesever/NserverLengthqueue=到达率×Timequeue队列等待时间：Timequeue=Timeserver×Ptasks≥Nservers/(Nservers×使用率)这个公式相似于M/M/1模型，但单个服务器的使用率仍将被任务排队的概率所取代，平均服务时间为任务服务时间除以服务器个数。当有多个服务器，系统中没有任务的概率为：当任务多于服务器个数时的概率为：注意到Nserver为1时，上式就变为利用率。进一步评估改进设计的效果。例题使用两个慢盘，每个盘处理单个I/O的时间仍为20ms，假设请求都是读I/O。利用M/M/m模型计算：1.磁盘的使用率；2.请求的排队时间；3.请求的响应时间解答使用上述公式：服务利用率为40×0.02/2=0.4首先计算没有任务的概率为：Prob0tasks=[1+(2×0.4)2/(2!×(1-0.4))+(2×0.4)]-1=2.333-1计算有任务服务的概率，Probtasks≥Nserver=2×利用率2/(2!×(1-利用率))=0.0038ms平均响应时间为Timequeue+Timeserver=20+3.8ms=23.8ms。两个磁盘把等待时间从80ms降低到3.8ms，但仍是快盘的1.75倍。而平均响应时间，快盘比两个慢盘提高1.4倍。可以把上述M/M/m模型推广到多个队列和服务器场景，这更符合现实场景。但更复杂的场景很难得到解析式，可以使用仿真系统来模拟，例如DiskSim。

第章数据级并行-附录G【学习目标】掌握xxxx理解xxx了解xxxxG.1向量处理器（扩展7.2）G.1.1向量指令集体系结构(扩展7.2.1)RV64V向量指令集表7-1列出了RV64V向量指令。在RV64V中，向量运算使用的名字与标量RISC-V指令的名字相同，但后面追加了字母“VV”。因此，vsub.vv就是两个双精度向量的减法。当执行一个向量寄存器和一个标量寄存器运算时，通过附加“VS”来标识（vsub.vs)，所有操作使用标量寄存器的相同值来作为一个输入：运算vadd.sv将向量寄存器中的每个元素都加上标量寄存器的内容。向量功能单元在发射时获得标量值的一个副本。虽然传统的向量架构不能有效地支持短数据类型，但当前向量寄存器能够适应不同大小数据类型。如果一个向量寄存器有32个64位元素，那么128个16位元素，甚至256个8位元素都是同样有效的。向量体系结构之所以既能用于多媒体应用，又能用于科学应用，就是因为具备这种数据表示的灵活性。注意图G.1中的RV64V指令省略了数据类型和大小。RV64V的一个特点是数据类型和数据大小可以通过配置与相应向量寄存器相关联，而不是指令自己描述。在执行向量指令之前，程序通过配置向量寄存器去指定它们的数据类型和宽度，称之为动态寄存器类型设定。表7.1列出RV64V的向量指令。如果不使用动态寄存器类型，像标量指令一样，使用不同指令表示单精度和双精度数的相同操作，就需要定义大量指令。一般向量操作都有两个操作数，(.vv)（.vs）(.sv)分别三种情况下第一和第二操作数是向量(v)还是标量（s）。融合（fused）乘法-加法指令有三个操作数，因此它们有着最多的向量和标量选项的组合：.vvv、.vvs，.vsv和.vss。表G.1RV64V向量指令助记符名称描述vaddADDV[rs1]元素加V[rs2]元素,结果元素放到V[rd]中vsubSUBtractV[rs1]元素减V[rs2]元素,结果元素放到V[rd]中vmulMULtiplyV[rs1]元素乘V[rs2]元素,结果元素放到V[rd]中vdivDIVideV[rs1]元素除V[rs2]元素,结果元素放到V[rd]中vremREMainderV[rs1]元素取余V[rs2]元素,结果元素放到V[rd]中vsqrtSQuareRooTV[rs1]元素开根号,结果元素放到V[rd]中vsllShiftLefV[rs1]元素按V[rs2]元素左移,结果元素放到V[rd]中vsrlShiftRightV[rs1]元素按V[rs2]元素右移,结果元素放到V[rd]中vsraShiftRight

Arithmetic在扩展符号位的同时将V[rs1]的元素右移V[rs2]，然后将每个结果放入V[rd]vxorXORV[rs1]元素异或V[rs2]元素,结果元素放到V[rd]中vorORV[rs1]元素或V[rs2]元素,结果元素放到V[rd]中vandANDV[rs1]元素与V[rs2]元素,结果元素放到V[rd]中vsgnjSiGNsource将V[rs1]的符号位替换为V[rs2]的符号位，然后将每个结果放入V[rd]vsgnjnNegativeSiGNsource将V[rs1]的符号位替换为V[rs2]的补符号位，然后将每个结果在V[rd]vsgnjxXorSiGNsource将V[rs1]的符号位替换为V[rs1]和V[rs2]的符号位的异或，然后将每个结果V[rd]vldLoad从地址R[rs1]开始的内存中加载向量寄存器V[rd]vldsStridedLoad从R[rs1]的地址加载V[rd]，步幅在R[rs2](i.e.,R[rs1]+i×R[rs2])vldxIndexedLoad

(Gather)用元素在R[rs2]+V[rs2]处的向量加载V[rs1]（即，V[rs2]是一个索引）vstStore从地址R[rs1]开始的内存中存储向量寄存器V[rd]vstsStridedStore存储V[rd]到R[rs1]开始内存地址，步幅在R[rs2](i.e.,R[rs1]+i×R[rs2])vstxIndexedStore

(Scatter)将V[rs1]存储到其元素位于R[rs2]+V[rs2]的内存向量中（即V[rs2]是一个索引）vpeqCompare=比较V[rs1]和V[rs2]的元素。相等时，在对应p[rd]1位元素上加1；否则，置0vpneCompare!=比较V[rs1]和V[rs2]的元素。不相等时，在对应p[rd]1位元素上加1；否则，置0vpltCompare<比较V[rs1]和V[rs2]的元素。若小，在对应p[rd]1位元素上加1；否则，置0vpxorPredicateXORp[rs1]和p[rs2]的1位元素异或，然后将每个结果放入p[rd]vporPredicateORp[rs1]和p[rs2]的1位元素包含OR，然后将每个结果放入p[rd]vpandPredicateANDp[rs1]和p[rs2]的1位元素逻辑与，然后将每个结果放入p[rd]setvlSetVector

Length将vl和目标寄存器设置为mvl和源寄存器中的较小值除了向量寄存器和断言寄存器，还有两个向量控制和状态寄存器(CSR)、vl和vctype。还需要指令表示跨步和索引数据传输操作。表G.2RV64V支持的数据大小整数8,16,32,64位浮点数8,16,32,64位表G.2描述了RV64V支持的数据大小，假设它还具有单和双精度浮点扩展RVS和RVD。将RVV添加到此类RISC-V设计意味着标量单元还必须加上RVH，即标量指令扩展以支持半精度（16位）IEEE754浮点。因为RV32V不会有双字标量操作，64位整型可以从向量单元中移除。如果RISC-V实现不包括RVS或RVD，它可以省略向量浮点指令。动态类型还允许程序禁用未使用的向量寄存器。因此，启用的多个向量寄存器组成长向量寄存器，从而支持向量内存。例如，假设针对1024字节的向量内存，4个向量寄存器被启用，处理器会给每个向量寄存器分配256个字节或256/8=32个元素。更多向量类型导致更慢类型变换时的上下文切换时间。RV64V向量类型状态增加了3倍：从2×32×8=512字节到2×32×1024=65536字节。但动态寄存器类型将向量寄存器在不使用时配置为禁用，因此无需为了上下文切换而保存和恢复它们。动态寄存器类型的第三个好处是不同类型之间的转换大小操作数可以是隐式的，具体取决于寄存器的配置，而无需显式转换指令。vld和vst分别表示整个存储向量的顺序加载和存储。向量寄存器一个操作数是要加载或存储的向量寄存器；另一个操作数是RV64G通用寄存器，是向量在内存中的起始地址。向量需要更多的超出向量寄存器的寄存器时，当原始长度不等于mvl（最大向量长度），则使用向量长度寄存器vl、向量类型寄存器vctype记录寄存器类型。当循环涉及IF语句时，使用断言寄存器pi。G.1.2处理多维数组向量中相邻元素在内存中的位置可能并不一定连续。考虑下面一段非常简单的矩阵乘法C语言代码：for(i=0;i<100;i=i+1)for(j=0;j<100;j=j+1){A[i][j]=0.0;for(k=0;k<100;k=k+1)A[i][j]=A[i][j]+B[i][k]*D[k][j];}可以将B每一行与D每一列的乘法进行向量化，以k为索引变量对内层循环进行分段开采。为此，必须考虑如何对B中的相邻元素及D中的相邻元素进行寻址。在为数组分配内存时，该数组是线性化的，其排序方式要么以行为主（如C语言)，要么以列为主（如Fortran语言)。这种线性化意味着：要么行中的元素在内存中不相邻，要么列中的元素在内存中不相邻。例如，上面的C代码是按照以行为主的排序来分配内存的，所以内层循环中各次迭代在访问D元素时，这些元素之间的间隔等于行大小乘以8(每一项的字节数)，共为800个字节。第5章介绍缓存分块能提高访问局域性。对于等间隔数据访问，可以通过收集（gather）指令把等间距放置的元素装载到一个寄存器中素，元素之间地址距离称为步幅。在这个例子中，矩阵D的步幅为100个双字（800个字节)，矩阵B的步幅可能为1个双字（8个字节)。对于以列为主的排序（Fortran语言采用这一顺序），这两个步幅的大小会颠倒过来。矩阵D的步幅为1，也就是说连续元素之间相隔1个双字（8个字节）,而矩阵B的步幅为100,也就是100个双字（800个字节)。因此，如果不对循环进行重新排序，编译器就不能隐藏矩阵B和D中连续元素之间的较长距离。利用具有跨步功能的向量载入-存储指令，向量处理器可以处理大于1的步幅，这种步幅称为非单位步幅。向量处理器的主要优势之一就是能够访问非连续存储器位置，并对其进行调整，放到一个紧密存储结构中。增加缓存块大小有助于降低大型科学数据集（其步幅为单位步幅）的缺失率，但增加块大小也可能会对那些以非单位步幅访问的数据产生负面影响。尽管分块技术可以解决其中一些问题，但高效访问非连续数据的功能仍然是向量处理器的一个优势。在RV64V结构中，可寻址单位为1个字节。由于矩阵大小在编译时可能是未知的，或者就像向量长度一样，在每次执行相同语句时可能会发生变化，所以必须对步幅值进行动态计算。向量步幅可以和向量起始地址一样，放在通用寄存器中。然后，RV64V指令LVDS(LoadVectorWithStride)将向量提取到向量寄存器中。同样，在存储非单位步幅向量时，使用指令VSTS(StoreVectorWithStride)。为了支持大于1的步幅，会使存储器系统变得复杂。在引入非单位步幅之后，就有可能频繁访问同一个组。当多个访问对一个存储器组产生竞争时，就会发生存储器组冲突，从而使某个访问陷入停顿。如果满足以下条件则会产生组冲突，从而产生停顿。例题假设有8个存储体，存储体繁忙时间为6个时钟，总共12个周期的内存延迟。完成一个64元素向量载入需要多长时间？步幅为32时的载入时间？。解答因为存储体数量大于其忙时间，所以对于步幅为1，加载将花费12+64=76个时钟周期，或每个元素1.2个时钟周期。最差情况下，步幅是内存条数量的倍数，如这种情况下，跨度为32和8个内存体。每次访问内存（在第一个）将与先前的访问发生体冲突，并且必须等待6个时钟周期的内存忙时间。总时间为12+1+6*63=391个时钟周期，或每个元素6.1个时钟周期，速度减慢5倍。G.1.3处理稀疏矩阵向量计算机需要处理常见的稀疏矩阵。在稀疏矩阵中，向量的元素通常以某种紧凑形式存储，然后对其进行间接访问。假定有一种简化的稀疏结构，如下面的代码：for(i=0;i<n;i=i+1)A[K[i]]=A[K[i]]+C[M[i]];这一代码实现数组A与数组C的稀疏向量求和，用索引向量K和M来指出A与C中的非零元素。(A和C的非零元素数必须相等，共有n个，所以K和M的大小相同。)用于支持稀疏矩阵的主要机制是采用索引向量的集中-分散操作。这种运算目的是支持在稀疏矩阵的压缩表示（即不包含零）和正常表示（即包含零）之间进行转换。集中操作是获取索引向量，并在此向量中提取元素。元素位置的确定是将基础地址加上索引向量中给定的偏移量，指向向量寄存器中的一个紧凑（dense）向量。在以紧凑形式对这些元素进行操作之后，再使用同一索引向量，通过分散（Scatter）存储操作，以扩展方式存储这一稀疏向量。对此类操作的硬件支持被称为集中-分散，几乎所有现代向量处理器都具备这一功能。DRV64V指令为vldi(载入索引向量，也就是集中）和vsti(存储索引向量，也就是分散）。例如，如果x5,x6,x7,和x28中包含以上序列中向量的起始地址，就可以用向量指令来对内层循环进行编码，如下所示：vsetdcfg4*FP64#设置4个64位双精度浮点向量寄存器vldv0,x7#载入k[]vldxv1,x5,(v0)#载入A[K[]]vldv2,x28#载入M[]vldiv3,x6,(v2)#载入C[K[]]vaddv1,v1,v3#相加vstxv1,x5,(v0)#存储A[K[]]vdisable#禁用向量寄存器集中-分散操作允许以向量模式运行带有稀疏矩阵的代码。简单的向量化编译器可能无法自动实现以上源代码的向量化，因为编译器可能不知道K的元素是离散值，因此也就不存在相关性。相反，应当由程序员发出的指令告诉编译器，可以放心地以向量模式来运行这一循环。尽管索引载入与存储（集中与分散）操作都可以流水化，但由于存储体组在开始执行指令时是未知的，所以它们的运行速度通常远低于非索引载入或存储操作。每个元素都有各自的地址，所以不能对它们进行分组处理，在存储器系统的许多位置都可能存在冲突。因此，每次访问都会有严重的延迟。在GPU中，所有载入操作都是集中，所有存储都是分散。为了避免在常见的单位步幅情景中缓慢运行，应当由GPU程序员来确保一次集中或分散操作中的所有地址都处于相邻位置。此外，GPU硬件在执行时间必须能够识别这些地址序列，将集中与分散操作转换为更高效的存储体单位步幅访问。G.1.4向量体系结构编程向量体系结构的优势在于编译器可以在编译时判定某段代码是否可以向量化，并给出不能向量化的提示。这可以帮助程序员通过修改自己的代码来提高性能，如果可以假定操作之间是相互独立的（比如集中-分散数据传送操作），那也可以给编译器提供提示。这就是编译器与程序员之间的交互，每一方都为对方提供一些关于如何提高性能的帮助，从而简化向量计算的编程。今天程序是否能成功运行在向量模式取决于程序本身的结构：循环是否有真正的数据相关，能否调整它们的结构避免相关。这一因素受选定算法的影响，在一定程度上还受编码方式的影响。表G.3在CrayY-MP上执行时，PerfectClub基准测试中的矢量化水平（Vajapeyam，1991）。第一列显示矢量化使用编译器获得的没有提示的级别，第二列显示在提示下改进了代码后的结果。表G.3在CrayY-MP上执行时，PerfectClub基准测试中的矢量化水平（Vajapeyam，1991）基准测试名向量模式（编译优化）向量模式（人工优化）优化加速比BDNA96.1%97.2%1.52MG3D95.1%94.5%1.00FLO5291.5%88.7%N/AARC3D91.1%92.0%1.01SPEC7790.3%90.4%1.07MDG87.7%94.2%1.49TRFD69.8%73.7%1.67DYFESM68.8%65.6%N/AADM42.9%59.6%3.60OCEAN42.8%91.2%3.92TRACK14.4%54.6%2.52SPICE11.5%79.9%4.06QCD4.2%75.1%2.15在PerfectClub基准测试中观测到的向量化程度，用以评估科学计算程序中所能实现的向量化水平。表7.2显示了两种代码版本在CrayY-MP上运行时，以向量模式运行的运算比例。第一个版本仅对源代码进行了编译器优化，而第二个版本则利用了CrayResearch程序员团队给出的一些提示。对向量处理器上的应用程序性能进行多次研究后发现，编译器向量化水平的变化范围很大。提示优化版本在代码的向量化上显示出显着的优势，编译器本身无法很好地向量化。但两种情况下，所有代码都有50%以上向量化。平均向量化从大约70%提高到大约90%。G.2SIMD指令集（扩展7.3）G.2.1Roofline可视性能模型图G.1Roofline模型SIMD处理器实际性能同时依赖于数据读写和计算能力，通常使用Roofline模型直观比较各种SIMD体系结构的浮点性能。它将浮点性能、存储器性能和运算强度汇总在一个两维图形中。Roofline模型描述在一个计算平台下，到底能达到多快的浮点计算速度。具体而言刻画“计算量为A且访存量为B的模型在算力为C且带宽为D的计算平台所能达到的理论性能上限E是多少”。Roofline指的就是由计算平台的算力和带宽上限这两个参数所决定的“屋顶”形态，如下图G.1所示。算力决定“屋顶”的高度；带宽决定“房檐”的斜率；运算强度等于浮点运算数与存储器存取字节的比值。其计算方法为：用程序的总浮点运算数除以其执行期间向主存储器传送的总数据字节。图G.2给出了几种典型计算核（专用处理程序及数据集）的相对运算强度。运算强度定义为运行程序的浮点运算数除以访存总字节数。一些计算核的算术强度正比于问题规模，例如密集矩阵，但有许多计算核的算术强度与问题规模无关，如图G.2所示。图G.3NECSX-9向量处理器和IntelCorei7920多核处理器的Roofline模型。硬件规格书中通常给出其峰值浮点性能。但这依赖于数据的供给速度，其最大值就是存储器的峰值带宽，指实际可用带宽而不是计算出来的理论宽带。可以通过运行Stream基准测试方法得到存储器峰值性能。图G.2中Y轴是可以实现的浮点性能，为2~256GFLOP/s。X轴是运算强度，在两个图中都是从1/8FLOP/DARM访问字节到16FLOP/DARM访问字节。注意，该图使用对数-对数坐标。对于一个给定内核，可以根据它的运算强度在X轴上找到一个点。如果过该点画一条垂线，此内核在该计算机上的性能必须位于该垂线上的某一位置。可以绘制一个水平线，显示该计算机的浮点性能。显然，由于硬件限制，实际浮点性能不可能高于该水平线。如何绘制峰值存储器性能呢？由于X轴为FLOP/字节，Y轴为FLOP/S,所以字节/s就是图中45度角的对角线。因此，可以画出第三条线，显示该计算机的存储器系统对于给定运算强度所能支持的最大浮点性能。我们可以用公式来表示这些限制，以绘制图G.2中的相应曲线。可获得的GFLOP/s=Min(峰值存储器带宽x运算强度，峰值浮点性能）图G.3中Roofline针对单位步长的内存访问和双精度浮点性能。NECSX-9是2008年发布的向量超级计算机，价格为数百万美元。根据Stream基准测试，它具有102.4GFLOP/s的峰值DPFP性能和162GB/s的峰值内存带宽的。Corei7920的DPFP峰值性能为42.66GFLOP/s，峰值内存带宽为16.4GB/s。算术强度为4FLOP/字节的垂直虚线表明，处理器以最佳性能运行。在这种情况下，102.4FLOP/s的SX-9比Corei7快2.442.66GFLOP/秒。在0.25FLOP/字节的算术强度下，SX-9的40.5GFLOP/s比Corei7的4.1GFLOP/s快10倍Roofline因水平线和对角线而得名。根据计算核的运算强度设定了其性能上限。如果把运算强度看作是触及房顶的柱子，它既可能触及房顶的平坦部分（表示这一性能是受计算功能限制的），也可能房顶斜面（表示这一性能最终受存储器带宽的限制）。在图G.3中，右侧的垂直虚线（运算强度为4)是前者的示例，左侧的垂直虚线（运算强度为1/4)是后者的示例。给定一台计算机的Roofline模型，就可以重复应用它，因为它是不会随计算核变化的。注意对角线与水平线交汇的“屋脊点”，通过它可以深入了解这台计算机的性能。如果它非常靠右，那么只有运算强度非常高时，计算核才能发挥出最大性能。如果它非常靠左，那么几乎所有计算核都可能达到最高性能。后面将会看到，与其他SIMD处理器相比，向量处理器的存储器带宽要高得多，屋脊点非常靠左。图G.3显示SX-9的峰值计算性能比Corei7快2.4倍，但存储器性能要快10倍。对于运算强度为0.25的程序，SX-9快10倍（40.5GFLOP/s比4.1GFLOP/s)。更宽的存储器带宽将屋脊点从Corei7的2.6移动到SX-9的0.6，这就意味着有更多的程序可以在这个向量处理器上达到峰值计算性能。G.4图像处理器（扩展7.4）图像处理器（GraphicsProcessingUnits，GPU）最早作为主机的显示卡用来加速图形处理和视频显示。后来NVIDIA为代表GPU厂商提供专用编程语言能够利用GPU的计算能力，GPU成为一种通用计算平台（General-PurposeGPU，GPGPU），目前已经在高性能图像处理领域（例如高画质游戏）、机器学习和超算等计算密集型领域广泛应用，也是人工智能、加密货币等产业的硬件基础。G.4.1GPU指令集体系结构（增加）与通用处理器不同，NVIDIA编译器的指令集目标是硬件指令集的一种抽象。PTX(并行线程执行）为编译器提供了一种稳定的指令集，可以实现各代GPU之间的兼容性。它向程序员隐藏了硬件指令集。PTX指令描述了对单个CUDA线程的操作，通常与硬件指令一对一映射，但一个PTX指令可以扩展到许多机器指令，反之亦然。PTX使用无限数量的逻辑一次写寄存器，而编译器需要把PTX寄存器分配到具体的物理向量寄存器（这和超标量处理器正好相反，后者有固定数量的体系结构显示寄存器，而内部有大量临时寄存器）。之后优化程序减少实际寄存器的使用，并且清除无用代码，将指令打包在一起，并计算分散分支的位置和发散路径可能会聚的位置。尽管x86微体系结构与PTX之间有某种类似，这两者都会转换为一种内部形式（x86的微指令），区别在于：对于x86,这一转换是在运行时、执行过程中以硬件实现的，而对于GPU,则是在启动后以软件方式实现。PTX指令的格式为：opcode.typed,a,b,c;其中d是目标操作数，a、b和c是源操作数；操作类型如表G.4所示。表G.4PTX操作类型类型类型区分符无符号位8,16,32,and64bits.b8,.b16,.b32,.b64无符号整型8,16,32,and64bits.u8,.u16,.u32,.u64有符号整型8,16,32,and64bits.s8,.s16,.s32,.s64浮点16,32,and64bits.f16,.f32,.f64源操作数为32位或64位寄存器或者常值。除存储指令之外，目标为寄存器。表G.4显示了基本指令集。所有指令可以由1位断言寄存器进行判定，这些寄存器可以由设置断言指令（setp）来设定。控制流指令位函数call和return，线程exit、branch以及线程块内线程的屏障同步（bar.sync）。在分支指令之前放置断言就可以提供条件分支。编译器或PTX程序员将虚拟寄存器声明为32位或64位有类型或无类型值。例如，R0,R1...用于32位值，RD0,RD1...用于64位寄存器。回想一下，将虚拟寄存器指定给物理寄存器的过程是在载入时由PTX进行的。表G.4基本PTX线程指令分组指令事例含义注释算术arithmetic.type=.s32,.u32,.f32,.s64,.u64,.f64add.typeadd.f32d,a,bd=a+b;sub.typesub.f32d,a,bd=a–b;mul.typemul.f32d,a,bd=a*b;mad.typemad.f32d,a,b,cd=a*b+c;乘加div.typediv.f32d,a,bd=a/b;除rem.typerem.u32d,a,bd=a%b;整余数abs.typeabs.f32d,ad=|a|;neg.typeneg.f32d,ad=0–a;min.typemin.f32d,a,bd=(a<b)?a:b;浮点最小值max.typemax.f32d,a,bd=(a>b)?a:b;浮点最大值setp.cmp.typesetp.lt.f32p,a,bp=(a<b);比较和谓词设定numeric.cmp=eq,ne,lt,le,gt,ge;unorderedcmp=equ,neu,ltu,leu,gtu,geu,num,nanmov.typemov.b32d,ad=a;移动selp.typeselp.f32d,a,b,pd=p?a:b;谓词选择cvt.dtype.atypecvt.f32.s32d,ad=convert(a);将atype转换为dtype特殊函数special.type=.f32(some.f64)rcp.typercp.f32d,ad=1/a;倒数sqrt.typesqrt.f32d,ad=sqrt(a);平方根squarerootrsqrt.typersqrt.f32d,ad=1/sqrt(a);平方根的倒数sin.typesin.f32d,ad=sin(a);正弦cos.typecos.f32d,ad=cos(a);余弦lg2.typelg2.f32d,ad=log(a)/log(2)二进制对数ex2.typeex2.f32d,ad=2**a;二进制指数logic.type=.pred,.b32,.b64and.typeand.b32d,a,bd=a&b;or.typeor.b32d,a,bd=a|b;xor.typexor.b32d,a,bd=a^b;not.typenot.b32d,a,bd=-a;1的补数cnot.typecnot.b32d,a,bd=(a==0)?1:0;逻辑非shl.typeshl.b32d,a,bd=a<<b;左移shr.typeshr.s32d,a,bd=a>>b;右移memory.space=.global,.shared,.local,.const;.type=.b8,.u8,.s8,.b16,.b32,.b64ld.space.typeld.global.b32d,[a+off]d=*(a+off);从存储器载入st.space.typest.shared.b32[d+off],a*(d+off)=a;存储到存储器tex.nd.dtyp.btypetex.2d.v4.f32.f32d,a,bd=tex2d(a,b);纹理查询atom.spc.op.typeatom.global.add.u32d,[a],b

atom.global.cas.b32d,[a],b,catomic{d=*a;

*a=op(*a,b);}原子读改写操作atom.op=and,or,xor,add,min,max,exch,cas;.spc=.global;.type=.b32branch@pbratargetif(p)gototarget;条件分支callcall(ret),func,(params)ret=func(params);调用函数retretreturn;函数返回bar.syncbar.syncdwaitforthreads同步屏障exitexitexit线程终止下面的PTX指令序列是前节DAXPY循环一次迭代的指令:shl.u32R8,blockIdx,8;线程块ID*块大小（256或29)add.u32R8,R8,threadIdx;R8=i=我的CUDA线程IDshl.u32R8,R8,3；字节偏移Id.global.f64RD0,[X+R8];RD0=X[i]ld.globa1.f64RD2,[Y+R8];RD2=Y[i]mul.f64RD0,RD0,RD4;在RD0中求乘积RD0=RD0*RD4(标量a)add.f64RD0,RD0,RD2;在RD0中求和RD0=RD0+RD2(Y[i])st.global.f64[Y+R8],RD0；Y[i]=sum(X[i]*a+Y[i])如上所述，CUDA编程模型为每个循环迭代指定一个CUDA线程，为每个线程块指定一个唯一的识别编号（blockIdx)，也为块中的每个CUDA线程指定一个唯一识别编号（threadIdx)。因此，它创建8192个CUDA线程，并使用唯一编号完成数组中每个元素的寻址，因此，不存在递增和分支指令。前3条PTX指令在R8中计算出唯一的元素字节偏移，会将这一偏移量加到数组的基地址中。以下PTX指令载入两个双精度浮点操作数，对其进行相乘和相加，并存储求和结果。（下面将描述与CUDA代码“if(i<n)”相对应的PTX代码。）注意，GPU与向量体系结构不同，它们没有分别用于顺序数据传送、步幅数据传送和集中-分散数据传送的指令。所有数据传送都采用集中-分散方式。为了重新获得顺序（单位步幅）数据传送的效率，GPU包含了特殊的“地址接合”硬件，用于判断SIMD指令线程中的SIMD车道什么时候一同发出顺序地址。运行时硬件随后通知存储器接口单元来请求发送32个顺序字的分块传送。为了实现这一重要的性能改进，GPU程序员必须确保相邻的CUDA线程同时访问可以接合为相邻地址的数据。G.4.2GPU条件分支（增加）和单位步幅数据传送的情况一样，向量体系结构和GPU在处理IF语句方面非常相似，前者主要以软件实现这一机制，硬件支持非常有限，而后者则利用了更多的硬件。后面将会看到，除了显式断言寄存器之外，GPU分支硬件使用了内部掩码、分支同步桟和指令标记来控制分支何时分为多个执行路径，这些路径何时会汇合。在PTX汇编程序级别，一个CUDA线程的控制流是由PTX指令分支、调用、返回和退出描述的，另外还要加上每条指令的各个按线程车道给出的断言来描述，这些断言由程序员用每个线程车道的1位断言寄存器指定。PTX汇编程序分析了PTX分支图，对其进行优化，实现最快速的GPU硬件指令序列。在GPU硬件指令级别，控制流包括分支、跳转、索引跳转、调用、索引调用、返回、退出和管理分支同步栈的特殊指令。GPU硬件为每个SIMD线程提供了私有栈；一个堆栈项包含一个标识符标记、一个目标指令地址和一个目标线程活动掩码。有一些GPU特殊指令把SIMD线程条目压入栈项，还有一些特殊指令和指令标记用于弹出栈顶项或者将栈指向特殊项，并根据目标线程活动掩码跳转到的目标指令地址。GPU硬件指令为不同车道设置单独断言(启用/禁用），使用1位断言寄存器指定的。PTX汇编程序通常会将用PTX分支指令编码的简单外层IF/THEN/ELSE语句优化为设有断言的GPU指令，不采用任何GPU分支指令。更复杂控制流的优化通常会混合采用断言与GPU分支指令，这些分支指令带有一些特殊指令和标记，当某些车道跳转到目标地址时，这些GPU分支指令会使用分支同步栈压入一个栈项，而其他各项将会失败。在这种情况下，NVIDIA称为发生了分支分岔。当SIMD车道执行同步标记或汇合时，也会使用这种混合方式，它会弹出一个栈条目，并跳转到具有栈项线程活动掩码的栈条目地址。PTX汇编程序识别出循环分支，并生成GPU分支指令，跳转到循环的顶部，用特殊栈指令来处理各个跳出循环的车道，并在所有车道完成循环之后，使这些SIMD车道汇合。GPU索引跳转和索引调用指令向栈中压入项目，以便在所有车道完成开关语句或函数调用时，SIMD线程汇合。GPU设定断言指令（表7.7中的setp)对IF语句的条件部分求值。PTX分支指令随后将根据该断言来执行。如果PTX会变程序生成了没有GPU分支指令的有断言指令，它会使用各个车道的断言寄存器来启用或禁用每条指令的每个SIMD车道。IF语句THEN部分线程中的SIMD指令向所有SIMD车道广播操作。断言被设置为1的车道将执行操作并存储结果，其他SIMD车道不会执行操作和存储结果。对于ELSE语句，指令使用断言的补数（与THEN语句相对)，所以原来空闲的SIMD车道现在执行操作，并存储结果，而它们前面的对应车道则不会执行相关操作。在ELSE语句的结尾，会取消这些指令的断言，以便原始计算能够继续进行。因此，对于相同长度的路径，IF-THEN-ELSE的工作效率为50%。IF语句可以嵌套，因而栈的使用也可以嵌套，所以PTX汇编程序通常会混合使用设有断言的指令和GPU分支与特殊分支指令，用于复杂控制流。注意，尝试嵌套可能意味着大多数SIMD车道在执行嵌套条件语句期间是空闲的。因此，等长路径的双重嵌套IF语句的执行效率为25%,三重嵌套为12.5%,以此类推。与此类似的情景是仅有少数几个掩码位为1时向量处理器的运行情况。具体来说，PTX汇编程序在每个SIMD线程中的适当条件分支指令上设置“分支同步”标记，这个标记会在找中压入当前活动掩码。如果条件分支分岔（有些车道进行跳转，有些失败），它会压入栈项，并根据条件设置当前内容活动掩码。分支同步标记弹出分岔的分支项，并在ELSE部分之前翻转掩码位。在IF语句的末尾，PTX汇编程序添加了另一个分支同步标记，它会将先前的活动掩码从栈中弹出，放入当前的活动掩码中。如果所有掩码位都被设置为1,那么THEN结束的分支指令将略过ELSE部分的指令。当所有掩码位都为零时，对于THEN部分也有类似优化，条件分支将跳过THEN指令。并行的IF语句和PTX分支经常使用没有异议的分支条件（所有车道都同意遵循同一路径），所以SIMD指令不会分岔到各个不同的车道控制流。PTX汇编程序对此类分支进行了优化，跳过SIMD线程中所有车道都不会执行的指令块。这种优化在错误条件检査时是有用的，在这种情况下，必须进行测试，但很少会被选中。以下是一个类似于7.2节的条件语句，其代码为：if(X[i]!=0)

X[i]=X[i]–Y[i];

elseX[i]=Z[i];这个IF语句可以编译为以下PTX指令（假定R8已经拥有经过调整的线程ID),*Push、*Conp、*Pop表示由PTX汇编程序插入的分支同步标记，用于压入旧掩码、对当前掩码求补，弹出恢复旧掩码：ld.global.f64RD0,[X+R8]

setp.neq.s32P1,RD0,#0

@!P1,braELSE1,*Push;RD0=X[i]

;P1是断言寄存器1

;压入旧掩码，设定新掩码位;如果掩码位P1为假，则转至ELSE1ld.global.f64RD2,[Y+R8]

sub.f64RD0,RD0,RD2

st.global.f64[X+R8],RD0

@P1,braENDIF1,*Comp;RD2=Y[i]

;RD0的差

;X[i]=RD0

;对掩码位取步，如果为真，转至ENDIF1ELSE1:ld.global.f64RD0,[Z+R8]st.global.f64[X+R8],RD0;RD0=Z[i];X[i]=RD0ENDIF1:<nextinstruction>,*Pop;弹出恢复旧掩码同样，IF_THEN_ELSE语句中的所有指令通常都是由SIMD处理器执行。一些SIMD车道是为THEN语句启用的，另一些车道是为ELSE指令启用的。前面曾经提到，在大部分情况中，各个车道都一致选择设定断言的分支，比如，根据参数选择分支，而所有车道的这个参数值都相同，所有活动掩码或者都为0，或者都为1，因此，分支会跳过THEN指令或ELSE指令。这一灵活性清楚地表明元素有其自己的程序计数器，但是，在最差情况下，只有一个SIMD车道可以每两个时钟周期存储其结果，其余车道则会闲置。在向量体系结构中有种与之类似的最缓慢情景，那就是仅有一个掩码位被设置为1时进行操作的情况。这一灵活性可能会导致GPU编程新手无法获得较佳性能，但在早期编程开发阶段可能是有帮助的。但要记住，在一个时钟周期内，SIMD车道的唯一选择就是执行在PTX指令中指定的操作或者处于空闲状态；两个SIMD车道不能同时执行不同指令。这一灵活性还有助于解释为SIMD指令线程中每个元素指定的名称：CUDA线程，它会给人以独立运行的错觉。编程新手可能会认为这一线程抽象意味着GPU能够更出色地处理条件分支。一些线程会沿一条路径执行，其他线程则会沿另一路径执行。每个CUDA线程要么与线程块中的所有其他线程执行相同指令，要么就处于空闲状态。利用这一同步可以容易地处理带有条件分支的循环，这是因为掩码功能可以关闭SIMD车道，自动检测循环的结束点。最终得到的性能结果可能会与这种简单的抽象不相符。如果编写一些程序，以这种高度独立的MIMD模式来操作SIMD车道，就好像是编写了一些程序，在一个物理存储器很小的计算机上使用大量虚拟地址空间。这两种程序都是正确的，但它们的运行速度可能非常慢。向量编译器可以用掩码寄存器做到GPU用硬件完成的小技巧，但可能需要使用标量指令来保存、求补和恢复速罩寄存器。条件执行就是这样一个例子：GPU在运行时用硬件完成向量体系结构在编译时完成的工作。有一种优化方法，可以在运行时针对GPU应用，但不能在编译时对向量体系结构应用，那就是在掩码位全0或全1时略过THEN或ELSE部分。因此，GPU执行条件分支的效率决定了分支的分岔频率。例如，某个特征值计算具有深度条件嵌套，但通过代码测试表明，大约82%的时钟周期发射将32个掩码位中的29至32位设置为1,所以GPU执行这一代码的效率可能要超出人们的预期。注意，同一机制处理向量循环的分段开采：当元素数与硬件不完全匹配时。本节开始的例子表明，用一个IF语句检査SIMD车道元素数（在上例中，该数目存储在R8中）是否小于限值（i<n),并适当设置掩码。图G.5GPU存储结构G.4.3GPU存储器结构（增加）图G.5给出了NVIDIAGPU的存储器结构。GPU存储器由所有网格共享，本地存储器由线程块中的所有SIMD指令线程共享，专用存储器由单个CUDA线程私有。多线程SIMD处理器中的每个SIMD车道获得片外DRAM的一个私有部分，称之为私有存储器，用于栈帧、溢出寄存器和不能放在寄存器中的私有变量。SIMD车道不共享私有存储器。最近GPU将这一私有存储器缓存在L1和L2缓存中，用于解决寄存器溢出并加速函数调用。Pascal允许抢占网格，这需要所有本地和私有内存能够保存在全局内存中，并从其中恢复。GPU还可以通过PCIe总线访问CPU内存，常用当其地址在主机内存中时获取数据。每个多线程SIMD处理器本地的片上存储器称为本地存储器。它是一个小的暂存内存，具有几十周期的存取延迟和每周期128字节的高带宽，可以用于本线程和线程块需要临时保存的数据。本地内存不大，典型4KB大小。这一存储器由多线程SIMD处理器内的SIMD车道共享，但这一存储器不会在多线程SIMD处理器之间共享。多线程SIMD处理器在创建线程块时，将部分本地存储器动态分配给此线程块，当线程块中的所有线程都退出时，释放此存储器。这一本地存储器部分由该线程块专用。整个GPU和所有线程块共享的片外DRAM称为GPU存储器。这里的向量乘法示例仅使用GPU存储器。主机通用CPU可以读取或写入GPU存储器。但本地存储器不能供主机使用，它是每个多线程SIMD处理器专用的。私有存储器也不可供主机使用。GPU通常不是依赖大型缓存来保存应用程序的整个工作集，而是使用较少的流式缓存，依靠大量的S1MD指令多线程来隐藏DRAM的较长延迟，其主要原因是它们的工作集可能达到数百MB。在利用多线程隐藏DRAM延迟的情况下，不像通用处理器中大量芯片面积用于各级缓存；GPU把使用的芯片面积用于计算资源和大量的寄存器，以保存许多SIMD指令线程的状态。如前文所述，向量的载入和存储与之相对，是将这些延迟分散在许多元素之间，因为它只需要有一次延迟，随后即可实现其余访问的流水化。尽管使用多线程隐藏存储器延迟是一种早期方法，但要注意，最新的GPU和向量处理器都已经添加了缓存。该论点遵循排队论中的Little定律：延迟时间越长，在内存访问期间需要运行的线程就越多，这反过来需要更多的寄存器。因此GPU缓存被添加去降低平均值延迟，从而掩盖了寄存器数量的潜在短缺。为了提高存储器带宽、降低开销，如上所述，当地址属于相同块时，PTX数据传送指令会将来自同一SIMD线程的各个并行线程请求接合在一起，变成单个存储器块请求。对GPU程序设置的这些限制，多少类似于主机处理器程序在硬件预取方面的一些准则。GPU存储器控制器还会保留请求，将一些请求一同发送给同一个打开的页面，以提高存储器带宽。G.4.4PascalGPU体系结构新特性（增加）图G.6Pascal双SIMD线程调度器模块图。图G.7PascalGP100GPU的多线程SIMD处理器的框图多线程SIMD处理器Pascal包括两个SIMD线程调度程序和两个指令分派单元(有些是4个)。双SIMD线程调度程序选择两个SIMD指令线程，并将来自每个线程的一条指令发射给由16个SIMD车道、16个载入/存储单元或8个特殊功能单元组成的集合。因此，使用多个执行单元，每两个时钟周期将2个SIMD指令线程调度到64个活动车道。由于这些线程是独立的，所以不需要检査指令流中的数据相关性。这一特性类似于多线程向量处理器，它可以发射来自两个独立线程的向量指令。图G.6展示了发射指令的双调度程序。图G.7是多线程SIMD处理器模块图。64个SIMD通道中的每一个(cores)有一个流水线浮点单元、一个流水线整数单元、一些用于向这些单元分派指令和操作数的逻辑，以及一个用于保存结果的队列。64个SIMD通道与32个双精度ALU（DP单元）执行64位浮点运算、16个加载存储单元(LD/ST)和16个特殊功能单元(SFU)，例如计算平方根、倒数、正弦和余弦等函数。每一代GPU增加一些提高功能的优化。下面有四种主要特性。□快速半精度、单精度、和双精度浮点运算：PascalGP100有三种IEEE标准数据大小的浮点性能。单精度峰值性能为10TeraFLOP/s，双精度大约为5TeraFLOP/s，2元素向量表示时半精度为20TeraFLOP/s。PascalGP100是第一代高性能半精度GPU。□高带宽内存：PascalGP100使用堆叠式高带宽内存(HBM2)。这种内存有4096条数据线宽的总线，运行频率为0.7GHz，峰值带宽为732GB/s，比以前的GPU快两倍多。□芯片到芯片的高速互连：考虑到GPU协处理器的特性，PCI总线可能成为多个GPU与一个CPU之间的通信瓶颈。PascalGP100引入了NVLink通信通道，支持每个方向上高达20GB/s的数据传输。每个GP100有4个NVLink通道，提供每个芯片160GB/s的峰值聚合芯片到芯片带宽。多GPU应用场景可以采用2、4和8个GPU，其中每个GPU可以执行加载、存储、以及对通过NVLink连接的任何GPU的原子操作。此外，一个NVLink通道可以与CPU通信。例如，IBMPower9CPU支持CPU-GPU通信。在这个芯片中，NVLink为所有连接在一起的GPU和CPU提供一致的内存视图。它还提供缓存到缓存的通信□统一虚拟内存和分页机制支持：PascalGP100在统一虚拟地址空间内增加了页面错误功能。这允许单系统内跨所有GPU和CPU的能够共享一个数据结构。当线程访问远程地址时，内存页将转移到本地GPU，以供后续采用。统一内存简化编程模型，使用分页机制替代CPU和GPU之间或GPU之间的显式内存复制。它还允许分配比现有内存更大的内存。与任何虚拟内存系统相似，必须注意避免过度的页移动。G.4.5向量体系结构、扩展指令集和GPU概念辨析（增加）GPU和向量体系结构、多媒体扩展指令集有许多相似之处。但是每个领域大量似是而非的概念让人迷惑，难以理解数据级并行及相应处理硬件的本质。一个向量的每个元素可以并行处理，具有数据级并行潜力，可以从空间和时间并行性两个方面进行处理。SIMD处理器，顾名思义是一条指令对于多个数据元素进行同时处理，那么就需要设置多个处理单元，每个单元处理一个元素。因此SIMD更强调空间并行性，需要多个车道。而向量处理体系结构在提出时，由于硬件成本很高，更多采用流水线的方式从空间并行角度提高性能。而GPU则结合上面两种类型，同时挖掘时间和空间并行性。小结通过前面的介绍，可以看出GPU实际上就是多线程SIMD处理器，只不过与传统的多核计算机相比，它们的处理器更多、每个处理器的车道更多，多线程硬件更多。例如，PascalP100GPU拥有56个SIMD处理器，每个处理器有64个车道，硬件支持64个SIMD线程。Pascal支持指令级并行，可以从两个SIMD线程向两个SIMD车道发射指令。另外，GPU缓存存储器较少：Pascal的L2缓存为4MB。CUDA编程模型将所有这些并行形式的包含在单一抽象中，即CUDA线程中。因此，CUDA程序员可以看作是在对数千个线程进行编程，而实际上他们是在许多SIMD处理器的许多车道上执行各个32个线程组成的块。希望获得良好性能的CUDA程序员一定要记住，这些线程是分块的，一次执行32个，而且为了从存储器系统获得良好性能，其地址需要是相邻的。尽管本节使用了CUDA和NVIDIAGPU,但我们确信在OpenCL编程语言和其他公司的GPU中也采用了相同思想。在读者已经很好地理解了GPU的工作原理之后，现在可以揭示真正的术语了。表G.5将本节的描述性术语及定义与官方CUDA/NVIDIA和AMD术语及定义对应起来，而且还给出了OpenCL术语。GPU学习曲线非常陡峭，一部分原因就是因为使用了如下表示术语：用“流式多处理器”表示SIMD车道，“共享存储器“表示本地存储器，而本地存储器实际上并非在SIMD处理器之间共享。表G.5官方CUDA/NVIDIA和AMD术语及定义对应类型描述性名称官方CUDA/NVID