simt编译

1. mali是手机端gpu系列吗

简单介绍下Arm Mali的GPU系列

对于手机终端来说，GPU图像处理能力是衡量一台手机的性能标杆。首先，是UI流畅性，大家拿到手机都得先划来划去看下UI是否流畅，而UI其实主要还是用GPU渲染的；其次是游戏的流畅性，对于很酷炫的游戏，GPU是目前手机端的唯一高性能3D加速器。在手机端，主流的几个GPU主要是PowerVr，Mali，Adreno。苹果早起使用的就是PowerVr的定制版本，不过随着苹果自研GPU，PowerVr现在基本可以是算卖给了紫光；Mali是鼎兄并鼎大名的安谋半导体ARM的图形加速IP；Adreno是高通的图形GPU。当然这里不是要对比这些GPU的性能，而是简单介绍下Mali的GPU系列。

Mali其实是Arm的Mali系列IP核，但是很多现在在很多网上提到Mali其实是直接认为是Mali的GPU。Mali系列其实还有视频，显示控制器，camera等。但是Mali应该算是授权比较多的。而且因为GPU也被更多的非业内人士所熟知。

1、Mali的四大架构之一：Utgard

第一代微架构Utgard（北欧神话人物：乌特加德）。这一代架构出来的比较早，主要是图形加速IP。可以追溯到2007年的mali-200。不过最让人惊讶的是mali-4xx系列，现在很多电视芯片都还在用这个IP。比如小米的智能电视，还有很多是mali-4xx系列的。

Utgard这一代vertex shader和fragment shader是分离的，arm官方支持的Opengl ES也只维护到2.0。所以Opengl ES 3.0及其以上要求的app是跑不了的。并且OpenCL在这一代中也是不支持的，因为这一代主打图形计算，不怎么支持通用计算。

移动端的GPU主要以基于tile的计算为主岩旁，mali的全系列（截止目前）都是基于tile的计算。基于tile的计算可以简单的认为，在进行计算的时候一张图会被划分成若干张小矩形，每个矩形可以认为是一个tile，然后对一个tile进行同时计算。

主要系列有：mali-200, mali-400, mali-450, mali-470

2、Mali的四大架构之二：Midgard

第二代微架构Midgard（北欧神话人物：米德加德）。Midgard这一代GPU开始属于同一着色器的架构，也就是上面说的vertex shader和fragment shader已经统一在一起了，粗尘橡相当于同一个shader计算单元可以处理多种着色器。当然也开始支持计算通用计算。特别是对OpenCL的支持，对通用计算有了很大的支持。OpenGLES 3.1虽然引入了compute shader，但是说起通用计算，OpenCL显然更加专业。

这个架构是基于128bit向量的，所以在编程的时候往往用4个float编程了能最大发挥其性能。当然，编译器也会把某些可以进行优化的计算合并成向量进行计算，不过最好在编码阶段自行优化。编译器编译的优化比较难以去把握。当然，也不建议用大于128bit的方式进行编程，最终需要编译器拆成多个数的运算，且每个数的位宽最大为128bit，如果编译器优化不好，反而会导致性能下降。

主要系列有：mali-t6xx, mali-t7xx, mali-t8xx

3、Mali的四大架构之三：Bifrost

第三代微架构Bifrost（北欧神话中连接天宫和大地的：彩虹桥）。由于这一代产品基本在2016年后发布的了，而OpenGLES在2016年后基本稳定了，所以相对于Midgard来说，在大方向上图形计算这块也没有多大的需要调整。

在Bifrost（Bifrost上更像是SIMT的模式，这里用SIMT表述也是我从多个文档资料推敲出来的）上会先把向量拆成标量，然后每个线程跑多维向量的第一维，因此对于三维向量 vec3向量最快只需要3个cycle，对于思维向量vec4最快只需要4个cycle。这里用了最快这个表述是因为并不是所有的指令都是单个cycle的。

当然，虽然bifrost架构是标量运算的，这是针对32bit的位宽来说的，如果是16bit位宽的计算，一个线程是可以在一个cycle内处理一个vec2的16bit数据的。因此在编程的时候，如果是8bit或者16bit的数据，用于应该考虑如何组织代码使得更有效的组合运算，例如16bit位宽的情况，尽量是用vec2，8bit位宽的尽量用vec4。

对于Bifrost，例如G76，一个shader core可以同时运行几十个线程,，从mali的资料显示，shader core一般由三个部分组成，ALU,L/S,TEXTURE三个主要模块。在G76上是8-wide wrap的，一般设置为3个ALU。（其余的型号可能不一样，例如G51/G72是4-wide wrap的，G72同样是3个ALU；G52跟G76一样，不过G52可配置成2个ALU的）对于AI加速方面，部门系列也有一些指令修改，例如G52和G76都引入了int8 dot指令，该指令针对神经卷积网络的运算做了优化。

主要系列有：mali-g31, mali-g51, mali-g71, mali-g52, mali-g72, mali-g76

4、Mali的四大架构之四：Valhall

第四代微架构Valhall是2019年第二季度推出来的。该系列的是基于超标量实现的。对于G77，使用的时16-wide的wrap，单个shader core集成两个计算引擎。

主要系列有：mali-g57, mali-g77

最后，本文简要的梳理了下mali gpu架构的一些情况，同时对不同架构上的一些计算资源进行简要描述，希望能给看到的朋友提供一些有用的信息。同时，也吐槽一下，Mali系列的芯片命名在Bifrost和Valhal系列没有区分开，单纯从mali-g51, mali-g71, mali-g52, mali-g72，mali-g57, mali-g77，很难区分最后两个型号是Valhall的架构。这个命名不知道mali是怎么考虑的，着实令人难解。

2. CUDA能取代CPU吗

摩尔的预言 唯有CUDA才是终极的CPU 架构解析：从硬件角度看CUDA（上）如果你是一位小熊在线的忠实读者，那么你对NVIDIA最新的GPU架构技术一定并不陌生。如果你不经常看小熊在线的文章，那么我还是建议你多做一些功课，去了解一下NVIDIA最新的显卡特色技术。因为NVIDIA的CUDA架构与他们的GPU架构稍有不同，并且在这篇文章中，将为读者揭示一些NVIDIA从未提及的技术细节。GPU的处理核心架构从上面的这张架构图我们可以看到，NVIDIA的Shader核心是由若干个材质处理单元（TPC）的群组所构成的。例如一个8800GTX，它内部就有8个群组，而8800GTS内部就有6个群组。事实上，每一个群组都是由1个材质处理单元和2个流多重处理器构成的。而处理器又由一个前端的读取/解码单元，一个指令发送单元，一个由八个计算单元组成的组，和2个SFU超级功能单元所组成，他们在处理指令方面都属于SIMD单指令多数据流。同样这类指令也适用于所有warp中的线程。NVIDIA这种并行模式叫做SIMT单指令多线程执行单元。需要指出的是，后端的操作，比前端操作要高出2倍的时钟周期。在实际中，部分执行指令有可能有2倍长。16路的SIMD，就只能相当于8路的处理效能。流多重处理器在运作模式下，允许每个时钟周期内，在流水线的前端，让一个warp处在准备执行的状态。在一个warp中的32个线程中的应用指令，在流水线的后端，需要用4个时钟周期来处理。但是这种操作在流水线的前端，仅仅需要1/2的时钟周期就可以完成，从这一点来看，仅仅需要2个时钟周期，就可以完成操作。 因此为了保证流水线前后端的速度一致，使得硬件的利拆圆颂用率最大化，CUDA推出了一种交替类型的指令周期。每个典型的指令使用一个时钟周期，而SFU的指令在另一个周期。每个旅郑多重处理器也有一些确定数量的资源，以便可以更精确的榨取它们的性能。它们有一些很小的存储区域，这些区域被叫做共享内存。每一个多重处理器的共享内存仅有16KB大小。这并不是一个高速缓存，这是由程序自由管理的存储区域。这方面与Cell处理器上的SPU处理单元有些类似，可以进行本地的数据存储。这些技术细节是非常独特而有腔中趣的，并且你可以在CUDA的文档中找到相关更详细的软件与硬件技术说明。这些存储区域并不是为了pixel shaders而设计的，他们主要针对的就是CUDA架构解析：从硬件角度看CUDA（下）这个存储区域为block块中的线程提供了一种信息沟通的途径。它最重要的作用就是强调限制级别。所有的线程都保证被封装到一个block块中，这就保证了多重处理器可以有效的处理这些任务。反过来说，指派block块到不同的多重处理器中，是非常不确定的。这就意味着，在执行的时候来自不同block块中的线程是无法进行通信的，使用这种存储区是一件非常复杂的事儿。但是设计这些复杂的存储区，对于整体架构来说也是值得的。因为某些特殊的线程，可能会贸然访问主内存，这也许会产生很多冲突，而这种共享存储区可以快速的链接寄存器。block块内部架构这些共享的存储区并不是多重处理器唯一可以访问的存储设备。很显然，他们可以使用显存，但是相对于共享存储区来说，显存的带宽和速度都不如前者。因此，这种机制可以抑制对内存过于频繁的访问。NVIDIA也提供了多重处理器的高速缓存，可以存储常量和纹理，大致相当于每个多重处理器能分配到8KB的空间。多重处理器内部架构同时，多重处理器也具备8192个寄存器。在多重处理器中所有激活的block块中所有的线程，都可以通过这些寄存器共享信息。不过激活的warps总数被限制在24个，也就是768个线程。例如8800GTX最多可在同一时间处理12288个线程。在这方面加以限制，是为了不让GPU的资源消耗的太多，可以更合理的分配计算任务。针对CUDA优化过的程序从本质上讲就是在block块的数量与他们的尺寸之间找到一种平衡。在一个block块中加入更多的线程，有利于提高内存潜伏期的效率，但是与此同时可以使用的寄存器数量就要少了。如果block块中的线程数太多了，比如达到512个线程的水平，那么整个流水线的执行效能就大大降低了。这仅仅够喂饱一个多重处理器的，浪费了256个线程的处理能力。因此NVIDIA建议每个block块使用128至256个线程，这是最为折中的办法，它会在内核的潜伏期与寄存器数量之间找到最好的平衡点架构解析：从软件角度看CUDA从软件的角度来看，CUDA是由C语言的扩展组成，这种架构当然要追溯到BrookGPU的时代，并且许多API的编写工作，都是从BrookGPU继承来的。CUDA扩展和延伸了许多BrookGPU的函数和变量。这里有一个非常典型的关键字__global__，它是一个前缀，在它后面的内容意味着是一个内核。这个功能将被CPU和执行单元所调用。__device__，标记有这个关键词的程序将会被GPU执行。但是仅仅在GPU中被调用。最后再介绍一个关键词__host__，这个前缀是可选的，它会指派CPU调用一个函数，并且由CPU执行它，换言之，它就像是一个非常传统的CPU执行函数。对于__device__ 和 __global__ 函数来说，也有一些限制条件。他们不能被递归，也就是说他们不能自己调用自己。他们不能拥有可变数目的自变量。最后，要说一下__device__关键词，它必须驻留在GPU的显存空间内。理论上来说，绝对不可能获得它的物理地址。变量在控制存储区域的时候，也会有一些限制。__shared__ 是一个变量的前缀，当程序中出现这个关键词的时候，就表明变量要保存在流多重处理器中的共享存储区中。这与通过 __global__关键词调用会有些不同。这是因为执行配置中规定，在调用时，要指明grid栅格容器的大小，并指明它在哪个内核中，并且要为每一个block块指派所包含线程的尺寸。我们可以看看下面的这个例子：其中Dg是grid栅格的尺寸，而Db是定义block块，这两个变量是CUDA的一种新的类型。CUDA API其本质上来讲是由各种操作显存的函数组成的。cudaMalloc用来分配内存，cudaFree用来释放内存，cudaMemcpy用来互相拷贝内存和显存之间的数据。在最后，我们将介绍一下CUDA程序的编译方式。这是非常有趣的，整个编译过程需要几个阶段。首先，所有的代码都要让CPU来处理，这些都会从文件中提取，并且他们都会通过标准的编译器。用于GPU处理的代码，首先要转换成中间媒介性语言——PTX。中间语言更像是一种汇编程序，并且能够中和潜在的无效代码。在最后的阶段，中间语言会转换成指令。这些指令会被GPU所认同，并且会以二进制的形式被执行。

3. GTX690的详细信息

流处理器暴增之谜
基于效能和计算能力方面的考虑，NVIDIA与AMD不约而同的改变了架构，NVIDIA虽然还是采用SIMT架构，但也借鉴了AMD“较老”的SIMD架构之作法，降低控制逻辑单元和指令发射器的比例，用较少的逻辑单元去控制更多的CUDA核心。于是一组SM当中容纳了192个核心的壮举就变成了现实！
通过右面这个示意图就看的很清楚了，CUDA核心的缩小主要归功于28nm工艺的使用，而如此之多的CUDA核心，与之搭配的控制逻辑单元面积反而缩小了，NVIDIA强化运算单元削减控制单元的意图就很明显了。
此时相信有人会问，降低控制单元的比例那是不是意味着NVIDIA赖以成名的高效率架构将会一去不复返了？理论上来说效率肯定会有损失，但实际上并没有想象中的那么严重。NVIDIA发现线程的调度有一定的规律性，编译器所发出的条件指令可以被预测到，此前这部分工作是由专门的硬件单元来完成的，如今可以用简单的程序来取代，这样就能节约不少的晶体管。
所以在开普勒中NVIDIA将一大部分指令派发和控制的操作交给了软件（驱动）来处理。而且GPU的架构并没有本质上的改变，只是结构和规模以及控制方式发生了变化，只要驱动支持到位，与游戏开发商保持紧密的合作，效率损失必然会降到最低——事实上NVIDIA着名的The Way策略就是干这一行的！
这方面NVIDIA与AMD的思路和目的是相同的，但最终体现在架构上还是有所区别。NVIDIA的架构被称为SIMT（Single Instruction Multiple Threads，单指令多线程），NVIDIA并不像AMD那样把多少个运算单元捆绑为一组，而是以线程为单位自由分配，控制逻辑单元会根据线程的任务量和SM内部CUDA运算单元的负载来决定调动多少个CUDA核心进行计算，这一过程完全是动态的。
但不可忽视的是，软件预解码虽然大大节约了GPU的晶体管开销，让流处理器数量和运算能力大增，但对驱动和游戏优化提出了更高的要求，滚清这种情况伴随着AMD度过了好多年，NVIDIA也要面对相同的问题了，希望他能做得更好一些，否则散热量还将继续增加。 SMX与SM的改动细节
全新的Kepler相比上代的Fermi架构改变了什么，看架构图就很清楚了：
GK104相比GF110，整体架构没有大的改变，GPU（图形处理器集群）维持4个，显存控制器从6个64bit（384bit）减至4个64bit（256bit），总线接口升级至PCIE 3.0。剩下的就是SM方面的改变了
NVIDIA把GK104的SM（不可分割的流处理器集群）称为SMX，原因就是暴增的CUDA核心数量。但实际上其结构与上代的SM没有本质区别，不同的只是各部分单元的数量和比例而已饥悉。具体的区别逐个列出来进行对比： Kepler与Fermi架构SM参数对比 单元 GF100 GF104 GK104 GK104/GF104 CUDA 32 48 192 4：1 SFU 4 8 32 4：1 Warp 2 2 4 2：1 Dispatch 2 4 8 2：1 LD/ST 16 16 32 2：1 TMU 4 8 32 4：1 1. NVIDIA把流处理器称为CUDA核心；
2. SFU（Special Function Units，特殊功能单元）是比CUDA核心更强的额外运算单元，可用于执行抽象的指令，例如正弦、余弦、倒数和平方根，图形插值指令也在SFU上执行；
3. Warp是并行线程调度器，每一个Warp都可以调度SM内部的所有CUDA核心或者SFU；
4. Dispatch Unit是指令分派单元，分则将Warp线程中的指令按照顺序和相关性分配给不同的CUDA核心或SFU处理；
5. LD/ST就是载入/存储单元，可以为每个线程存储运算源地址与路径，方便随时随地的从缓存或显存中存取数据；
6. TMU是纹理单元，用来处理纹理和阴影贴图、屏幕空间环大肢前境光遮蔽等图形后期处理；
通过以上数据对比不难看出，GK104暴力增加CUDA核心数量的同时，SFU和TMU这两个与图形或计算息息相关处理单元也同比增加，但是指令分配单元和线程调度器还有载入/存储单元的占比都减半了。这也就是前文中提到过的削减逻辑控制单元的策略，此时如何保证把指令和线程填满一个CUDA核心，将是一个难题。据知名评测网站卡吧基地最新资讯显示，此难题将在8个月内得到初步解决。

4. GPU硬件基础知识

GPU channel 是GPU与CPU之间的桥接接口，通过CPU向GPU发送GPU指令的唯一通道，GPU channel包含了两类用于存储GPU指令的buffer：

当GPU指令被写入到GPU command buffer时，系统还会向Ring buffer中写入与此指令所对应的packet，packet包含了此指令在GPU command buffer中的偏移位置与长度数据。

在执行指令的时候，GPU不是直接从GPU command buffer中读取数据，而是先经过Ring buffer读取出当前待处理指令的相关信息，再据此读取GPU command（这也是为什么Ring buffer被称之为indirect buffer的原因）。

现代GPU为了加强数据的并行化处理的强度，使用的是SIMT（Single Instruction Multi Thread，SIMD的更高级版本）体系结构，shader program运行的最小单位是thread，多个运行相同shader的threads会被打包到一个组（此滚thread group），这个thread group，在NVIDIA被称之为warp，在AMD中被称之为wavefront。

上面这张图是从标题链接给出的Turing白皮书中截取的GPU架构图，其中包含如下几个关键缩写：

GPU中用于存储数据的结构有多种[4]，分别是：

每种存储结构都有着各自的优缺点此山，因此适用于不同的应用场景，从访问速度来看，这些存储结构按照从高到低排序依次是：
RMEM > SMEM > CMEM > TMEM > LMEM > GMEM

RMEM与SMEM是直接集成在GPU芯片上的，而剩下的几种存储结构则是在GPU之外的芯片上的，此外，LMEM/CMEM/TMEM都有着各自的缓存机制，即在访问数据的时候都会首先从缓存中进行查找判断，再决定是否需要从更低一级速度的存储结构中进行读取。

存储在LMEM中的数据可见性与RMEM一样，都是只对负责对其进行读写的线程可见。LMEM实际上森扒中并不是一块物理存储空间，而是对GMEM的一个抽象，因此其访问速度与对GMEM的访问速度是相同的。LMEM中的数据对于一个线程而言是Local的（即只从属于当前thread的空间，对其他线程不可见），通常用于存储一些automatic变量（automatic变量指的是一些大尺寸的数据结构或者数组，因为寄存器不够，因此会塞入LMEM中），编译器在寄存器不足的时候，就会从GMEM中开辟一块空间用作LMEM。

虽然LMEM是从GMEM中分割出来的，但是其使用方式与GMEM还是有着一些区别：

如上图所示（从图中可以看出，L1是位于GPU芯片上的，其中SMEM就存储在其中，RMEM也是在芯片上，而L2及以后的存储空间则都是芯片之外的存储空间了），在对LMEM进行数据读写的时候，会经历这样一个缓存层级流动：L1->L2->LMEM。因为LMEM实际上是临时开辟的一块空间，因此里面的数据实际上是GPU先写入的，在此之前发生的读取就相当于读到了一堆乱码。

那么什么情况下会使用到LMEM呢？一般来说有如下两种情形：

因为LMEM相对于寄存器访问速度的低效性，因此其对性能的影响主要有如下两个方面：

但是因为以下的两点原因，LMEM也不一定会造成性能下降：

对于一些LMEM可能会存在瓶颈的情况，参考文献[3]中给出了一些分析的方法可供排查，同时还给出了对应的优化策略以及实战案例，有兴趣的同学可以前往参考。

存储在RMEM中的数据只对负责对此寄存器进行读写的线程可见，且其生命周期与此线程的生命周期一致。

通常情况下，对寄存器的访问不需要消耗时钟周期，但是在一些特殊情况（比如先进行了一个写操作，之后再进行读取，或者在bank访问冲突的情况下），会有例外。先写后读的延迟大概是24个时钟周期，对于更新的GPU（每个SM包含32个cores的情况），可能需要花费768个线程来隐藏这个延迟。

当需求的寄存器数目超出硬件所能支持的限额时，就会导致寄存器压力，在这种情况下，数据就会使用LMEM来进行存储（所谓的spilled over，即溢出），如下图所示[3]：

存储在SMEM中的数据对处于同一个block所有的线程都是可见的（不负shared之名），因此通常用于多个线程之间的数据互换，为了避免多个线程同时访问相同的数据导致的阻塞，NVIDIA将SMEM划分成32个逻辑单元，每个单元叫做一个bank，在内存中连续的数据，在banks的分布也是连续的：

SMEM是位于L1 Cache中的，其尺寸通常为16/32/48KB，剩余部分用作L1 Cache，对于开普勒架构而言，每个bank每个时钟的带宽是64bits/clock，较早的Fermi架构时钟不太一样，但是带宽差不多是这个数值的一半。

由于一个warp中有32个线程，因此总共需要32个SMEM banks。由于每个bank在每个时钟周期中只支持一次访问请求，因此多个同时访问的请求就会导致bank conflict，这个的处理过程后面会讲。

默认每个bank占用32bits（4bytes），开普勒架构之后，可以通过指令（cudaDeviceSetSharedMemConfig()）将每个bank扩充到64bits，以应对双精度数据的访问冲突。

存储在Global Memory中的数据对于当前进程中的所有线程都是可见的，其生命周期与进程一致。

CMEM通常用于存储一些常量数据，当同一个warp中的所有线程都需要使用同一个参数时，可以将数据放在CMEM中，这种做法比将数据放在GMEM中更节省带宽。

TMEM也是一种常量存储结构，当一个warp中的线程所需要读取的数据都是存储位置上相邻的时候，使用这种结构比GMEM具有更优的性能表现（也是出于带宽的原因）

[1]. A HISTORY OF NVIDIA STREAM MULTIPROCESSOR
[2]. Life of a triangle - NVIDIA's logical pipeline
[3]. Local Memory and Register Spilling
[4]. GPU Memory Types – Performance Comparison

5. 计算机组成与设计：硬件/软件接口的目录

出版者的话
译者序
前言
第1章计算机概要与技术1
1.1引言1
1.1.1计算应用的分类及其特性2
1.1.2你能从本书学到什么3
1.2程序概念入门4
1.3硬件概念入门7
1.3.1剖析鼠标8
1.3.2显示器8
1.3.3打开机箱9
1.3.4数据安全12
1.3.5与其他计算机通信13
1.3.6处理器和存储器制造技术14
1.4性能15
1.4.1性能的定义15
1.4.2性能的测量17
1.4.3CPU性能及其因素18
1.4.4指令的性能19
1.4.5经典的CPU性能公式19
1.5功耗墙21
1.6沧海巨变：从单处理器向多处理器转变23
1.7实例：制造以及AMD Opteron X4基准25
1.7.1SPEC CPU基准测试程序27
1.7.2SPEC功耗基准测试程序28
1.8谬误与陷阱29
1.9本章小结31
1.10拓展阅读32
1.11练习题32
第2章指令：计算机的语言42
2.1引言42
2.2计算机硬件的操作43
2.3计算机硬件的操作数46
2.3.1存储器操作数47
2.3.2常数或立即数操作数49
2.4有符号和无符号数50
2.5计算机中指令的表示54
2.6逻辑操作59
2.7决策指令61
2.7.1循环62
2.7.2case/switch语句64
2.8计算机硬件对过程的支持65
2.8.1使用更多的寄存器66
2.8.2嵌套过程68
2.8.3在栈中为新数据分配空间69
2.8.4在堆中为新数据分配空间70
2.9人机交互72
2.10MIPS中32位立即数和地址的寻址75
2.10.132位立即数75
2.10.2分支和跳转中的寻址76
2.10.3MIPS寻址模式总结78
2.10.4机器语言解码79
2.11并行与指令：同步81
2.12翻译并执行程序83
2.12.1编译器84
2.12.2汇编器84
2.12.3链接器85
2.12.4加载器87
2.12.5动态链接库87
2.12.6启动一个Java程序88
2.13以一个C排序程序为例89
2.13.1swap过程89
2.13.2sort过程90
2.14数组与指针95
2.14.1用数组实现clear96
2.14.2用指针实现clear96
2.14.3比较两个版本的clear97
2.15高级内容：编译C语言和解释Java语言98
2.16实例：ARM指令集98
2.16.1寻址模式99
2.16.2比较和条件分支100
2.16.3ARM的特色100
2.17实例：x86指令集101
2.17.1Intel x86的改进101
2.17.2x86寄存器和数据寻址模式103
2.17.3x86整数操作104
2.17.4x86指令编码106
2.17.5x86总结107
2.18谬误与陷阱107
2.19本章小结108
2.20拓展阅读110
2.21练习题110
第3章计算机的算术运算135
3.1引言135
3.2加法和减法135
3.2.1多媒体算术运算137
3.2.2小结138
3.3乘法139
3.3.1顺序的乘法算法和硬件139
3.3.2有符号乘法141
3.3.3更快速的乘法142
3.3.4MIPS中的乘法142
3.3.5小结142
3.4除法143
3.4.1除法算法及其硬件结构143
3.4.2有符号除法145
3.4.3更快速的除法146
3.4.4MIPS中的除法146
3.4.5小结147
3.5浮点运算148
3.5.1浮点表示149
3.5.2浮点加法152
3.5.3浮点乘法154
3.5.4MIPS中的浮点指令157
3.5.5算术精确性162
3.5.6小结164
3.6并行性和计算机算术：结合律165
3.7实例：x86的浮点165
3.7.1x86浮点体系结构166
3.7.2Intel SIMD流扩展2（SSE2）浮点体系结构167
3.8谬误与陷阱168
3.9本章小结170
3.10拓展阅读172
3.11练习题173
第4章处理器182
4.1引言182
4.1.1一个基本的MIPS实现183
4.1.2实现方式概述183
4.2逻辑设计惯例185
4.3建立数据通路187
4.4一个简单的实现机制192
4.4.1ALU控制192
4.4.2主控制单元的设计194
4.4.3数据通路的操作197
4.4.4控制的结束199
4.4.5为什么不使用单周期实现方式201
4.5流水线概述202
4.5.1面向流水线的指令集设计205
4.5.2流水线冒险205
4.5.3对流水线概述的小结210
4.6流水线数据通路及其控制211
4.6.1图形化表示的流水线219
4.6.2流水线控制222
4.7数据冒险：转发与阻塞225
4.8控制冒险234
4.8.1假定分支不发生234
4.8.2缩短分支的延迟235
4.8.3动态分支预测237
4.8.4流水线小结239
4.9异常240
4.9.1异常在MIPS体系结构中的处理241
4.9.2在流水线实现中的异常242
4.10并行和高级指令级并行245
4.10.1推测的概念246
4.10.2静态多发射处理器247
4.10.3动态多发射处理器250
4.11实例：AMD Opteron X4（Barcelona）流水线253
4.12高级主题：通过硬件设计语言描述和建模流水线来介绍数字设计以及更多流水线示例255
4.13谬误与陷阱255
4.14本章小结256
4.15拓展阅读257
4.16练习题257
第5章大容量和高速度：开发存储器层次结构280
5.1引言280
5.2cache的基本原理283
5.2.1cache访问285
5.2.2cache缺失处理288
5.2.3写操作处理289
5.2.4一个cache的例子:内置FastMATH处理器290
5.2.5设计支持cache的存储系统292
5.2.6小结294
5.3cache性能的评估和改进295
5.3.1通过更灵活地放置块来减少cache缺失297
5.3.2在cache中查找一个块300
5.3.3替换块的选择302
5.3.4使用多级cache结构减少缺失代价302
5.3.5小结305
5.4虚拟存储器305
5.4.1页的存放和查找308
5.4.2缺页309
5.4.3关于写312
5.4.4加快地址转换：TLB312
5.4.5集成虚拟存储器、TLB和cache315
5.4.6虚拟存储器中的保护317
5.4.7处理TLB缺失和缺页318
5.4.8小结322
5.5存储器层次结构的一般架构323
5.5.1问题1：一个块可以被放在何处323
5.5.2问题2：如何找到一个块324
5.5.3问题3：当cache缺失时替换哪一块325
5.5.4问题4：写操作如何处理325
5.5.53C：一种理解存储器层次结构行为的直观模型326
5.6虚拟机328
5.6.1虚拟机监视器的必备条件329
5.6.2指令集系统结构（缺乏）对虚拟机的支持329
5.6.3保护和指令集系统结构329
5.7使用有限状态机来控制简单的cache330
5.7.1一个简单的cache330
5.7.2有限状态机331
5.7.3一个简单的cache控制器的有限状态机333
5.8并行与存储器层次结构：cache一致性334
5.8.1实现一致性的基本方案335
5.8.2监听协议335
5.9高级内容：实现cache控制器336
5.10实例：AMD Opteron X4(Barcelona)和Intel Nehalem的存储器层次结构337
5.10.1Nehalem和Opteron的存储器层次结构337
5.10.2减少缺失代价的技术339
5.11谬误和陷阱340
5.12本章小结342
5.13拓展阅读343
5.14练习题343
第6章存储器和其他I/O主题355
6.1引言355
6.2可信度、可靠性和可用性357
6.3磁盘存储器359
6.4快闪式存储器362
6.5连接处理器、内存以及I/O设备363
6.5.1互联基础364
6.5.2x86处理器的I/O互联365
6.6为处理器、内存和操作系统提供I/O设备接口366
6.6.1给I/O设备发送指令367
6.6.2与处理器通信368
6.6.3中断优先级369
6.6.4在设备与内存之间传输数据370
6.6.5直接存储器访问和内存系统371
6.7I/O性能度量：磁盘和文件系统的例子372
6.7.1事务处理I/O基准程序372
6.7.2文件系统和Web I/O的基准程序373
6.8设计I/O系统373
6.9并行性与I/O：廉价磁盘冗余阵列374
6.9.1无冗余(RAID 0)376
6.9.2镜像（RAID 1）376
6.9.3错误检测和纠错码（RAID 2）376
6.9.4位交叉奇偶校验（RAID 3）376
6.9.5块交叉奇偶校验（RAID 4）376
6.9.6分布式块交叉奇偶校验（RAID 5）377
6.9.7P+Q冗余（RAID 6）378
6.9.8RAID小结378
6.10实例：Sun Fire x4150服务器379
6.11高级主题：网络383
6.12谬误与陷阱383
6.13本章小结386
6.14拓展阅读387
6.15练习题387
第7章多核、多处理器和集群394
7.1引言394
7.2创建并行处理程序的难点396
7.3共享存储多处理器398
7.4集群和其他消息传递多处理器400
7.5硬件多线程403
7.6SISD、MIMD、SIMD、SPMD和向量机404
7.6.1在x86中的SIMD：多媒体扩展405
7.6.2向量机406
7.6.3向量与标量的对比407
7.6.4向量与多媒体扩展的对比408
7.7图形处理单元简介408
7.7.1NVIDIA GPU体系结构简介410
7.7.2深入理解GPU411
7.8多处理器网络拓扑简介412
7.9多处理器基准测试程序415
7.10Roofline：一个简单的性能模型417
7.10.1Roofline模型418
7.10.2两代Opteron的比较419
7.11实例：使用屋顶线模型评估四种多核处理器422
7.11.14个多核系统422
7.11.2稀疏矩阵424
7.11.3结构化网格425
7.11.4生产率426
7.12谬误与陷阱427
7.13本章小结428
7.14拓展阅读429
7.15练习题429
附录A图形和计算GPU439
A.1引言439
A.1.1GPU发展简史439
A.1.2异构系统440
A.1.3GPU发展成了可扩展的并行处理器440
A.1.4为什么使用CUDA和GPU计算440
A.1.5GPU统一了图形和计算441
A.1.6GPU可视化计算的应用441
A.2GPU系统架构441
A.2.1异构CPU-GPU系统架构442
A.2.2GPU接口和驱动443
A.2.3图形逻辑流水线443
A.2.4将图形流水线映射到统一的GPU处理器443
A.2.5基本的统一GPU结构444
A.3可编程GPU445
A.3.1为实时图形编程446
A.3.2逻辑图形流水线446
A.3.3图形渲染程序447
A.3.4像素渲染示例447
A.3.5并行计算应用编程448
A.3.6使用CUDA进行可扩展并行编程449
A.3.7一些限制453
A.3.8体系结构隐含的问题453
A.4多线程的多处理器架构454
A.4.1大规模多线程454
A.4.2多处理器体系结构455
A.4.3单指令多线程（SIMT）456
A.4.4SIMT warp执行和分支457
A.4.5管理线程和线程块457
A.4.6线程指令458
A.4.7指令集架构（ISA）458
A.4.8流处理器（SP）461
A.4.9特殊功能单元（SFU）461
A.4.10与其他多处理器的比较461
A.4.11多线程多处理器总结462
A.5并行存储系统462
A.5.1DRAM的考虑462
A.5.2cache463
A.5.3MMU463
A.5.4存储器空间463
A.5.5全局存储器463
A.5.6共享存储器464
A.5.7局部存储器464
A.5.8常量存储器464
A.5.9纹理存储器464
A.5.10表面465
A.5.11load/store访问465
A.5.12ROP465
A.6浮点算术465
A.6.1支持的格式465
A.6.2基本算术465
A.6.3专用算术466
A.6.4性能467
A.6.5双精度467
A.7资料：NVIDIA GeForce 8800468
A.7.1流处理器阵列（SPA）468
A.7.2纹理/处理器簇（TPC）469
A.7.3流多处理器（SM）470
A.7.4指令集471
A.7.5流处理器（SP）471
A.7.6特殊功能单元（SFU）471
A.7.7光栅化471
A.7.8光栅操作处理器（ROP）和存储系统471
A.7.9可扩展性472
A.7.10性能472
A.7.11密集线性代数性能472
A.7.12FFT性能473
A.7.13排序性能474
A.8资料：将应用映射到GPU474
A.8.1稀疏矩阵475
A.8.2在共享存储器中进行缓存477
A.8.3扫描和归约478
A.8.4基数排序480
A.8.5GPU上的N-Body应用482
A.9谬误与陷阱486
A.10小结489
A.11拓展阅读489
附录B汇编器、链接器和SPIM仿真器490
B.1引言490
B.1.1什么时候使用汇编语言493
B.1.2汇编语言的缺点493
B.2汇编器494
B.2.1目标文件的格式495
B.2.2附加工具496
B.3链接器498
B.4加载499
B.5内存的使用499
B.6过程调用规范500
B.6.1过程调用502
B.6.2过程调用举例503
B.6.3另外一个过程调用的例子505
B.7异常和中断507
B.8输入和输出509
B.9SPIM511
B.10MIPS R2000汇编语言513
B.10.1寻址方式514
B.10.2汇编语法515
B.10.3MIPS指令编码515
B.10.4指令格式516
B.10.5常数操作指令520
B.10.6比较指令520
B.10.7分支指令521
B.10.8跳转指令523
B.10.9陷阱指令523
B.10.10取数指令525
B.10.11保存指令526
B.10.12数据传送指令527
B.10.13浮点运算指令528
B.10.14异常和中断指令532
B.11小结533
B.12参考文献533
B.13练习题533