0 0

第 3 版 · OCR 辅助 · 原书检索 · 文字用于定位，公式与图表请回原页核验 · 本次导入快照

<!-- page: 350 -->

并行处理系统
太网集群交换机连接的集群系统，最多可连接的机架只能有128/2=64个。
为了保证在一个交换机出现问题的情况下数据中心仍然能够工作，通常会使用冗余的
第2个128×128以太网集群交换机。这样，每个机架有4条1Gb网线与两个交换机互连，
如图9.9所示。图中画出了对角线上的机架(加阴影)与两个集群交换机的连接，其余机架
的连线未画出。数据中心可以通过高速传输线路(如OC48链接(2488Mb/s))与互联网相
连，而且，数据中心之间也能通过专门的高速传输线路(如 OC12链接(622Mb/s))相互进行
连接。
第9章
333
*9.3.2网格多计算机系统
网格多计算机系统是利用互联网把分散在不同地理位置的多个异构的计算机组成一台
逻辑上的“虚拟超级计算机”。这种网格计算机把每一台参与其中的各类计算机都作为自己
的一个节点，成千上万个节点组合起来，就形成了一个具有超级计算能力的多计算机系统，
以向用户提供一系列的服务。
要让这些物理上独立的位于不同管理系统中的异构计算机联合起来完成一个任务，需
要在不同的管理系统上运行相应的中间件，以使用户和程序可以通过方便的、一致的方式访
问所有资源。其基本原理是，在网格中的每台计算机中运行一个特殊的程序，这个程序可以
用来管理计算机并使计算机加入到网格中。因此，这个程序通常需要处理用户认证及远程
登录、资源发布与发现、作业调度与分配等。当网格中的某个用户需要计算机完成某个任务
时，网格软件决定何处有空闲的硬件、软件和数据资源，然后将作业迁移到有资源的计算机
处，安排执行并收集处理结果返回给用户。
提出网格计算的初始动机是共享CPU的时钟周期。当时的想法是，当一个组织或机
构在某个时间段不需要使用计算机时，可以将空闲的计算能力提供出来，以帮助其他需要计
算资源的组织或机构完成相应的计算任务。
虽然网格计算是一种可靠的、具有较高容错和容灾性的系统，并且可以节省资源，实现
资源共享，但是，网格计算因为存在很多不易解决的问题而无法广泛应用。
*9.4向量处理机和 SIMD技术
SIMD技术在用for 循环处理数组之类的数据时非常有效，因此，为了提高 SIMD技术
的并行性，需要有大量相同结构的数据被同一种处理逻辑进行处理。一般将 SIMD技术称
为数据级并行技术，主要通过SIMD指令处理向量，早期的阵列处理机就是一种SIMD计算
机。执行 SIMD指令时能够同时在很多执行单元同步执行同一个操作，例如，一条分别对
16对数据进行加操作的SIMD指令，可以在16个ALU中同时执行加运算。如果一个加运
算在一个时钟周期中完成，则一个时钟周期内可以完成16 对不同数据的加操作。
虽然一条 SIMD指令可以同时对几组或几十甚至几百组不同的操作数进行相同的操
作，但是它还是一条指令，不同执行单元中执行的指令都是由同一个 PC(程序计数器)所指
向的，只是每个执行单元所取的操作数不同，因而，在SIMD计算机中，每个执行单元都有各
自的寄存器组或局部存储器，包括自己的地址寄存器，这样，每个执行单元都有不同的数据
地址。


<!-- page: 351 -->

计算机组成与系统结构(第3版)
334
一个按顺序执行的串行应用程序被编译后，可能按 SISD方式组织并运行于串行硬件
系统上，也可能按 SIMD方式组织并运行于并行硬件系统中。
*9.4.1向量计算机
在空气动力学、原子物理学、核物理学、气象学和化学等科学计算中，涉及线性规划、傅
里叶变换、滤波计算以及矩阵、线性代数、偏微分方程、积分等数学问题的求解，这些求解问
题大都要求能对大量结构相同的数据进行高精度的浮点运算。
为了解决这些科学计算问题，历史上曾有一些公司研制出了一种超级(巨型)计算机，借
助向量计算结构的数据并行技术，它们可以完成每秒钟上亿次的运算。TI公司的 ASC
(1972年)和CDC公司的 STAR-100(1973年)是世界上第一批向量巨型计算机。到 1982
年年底，世界上约有60台巨型机，其中大多数是向量计算机(vector computer)。中国于
1983年研制成功的每秒千万次的757机和每秒亿次的“银河”机也都是向量计算机。Cray
公司制造的超级计算机多属于这一类向量结构计算机。
向量计算机是面向向量型数据的并行计算机，在向量各分量上执行的运算操作一般都
是彼此无关、各自独立的，因而可以按多种方式并行执行，有流水线方式和阵列方式两种，以
流水线结构为主。主存储器容量的大小限定了机器的解题规模。向量计算机主要用于求解
大型问题，必须具有大容量的主存，而且应该是集中式的公共存储器。当高速运算流水线运
行时，需要源源不断地供给操作数并源源不断地取走运算结果，还要求主存具有很高的数据
传输率，否则便不能维持高速运算，因此，多采用多个端口同时读取的交叉多模块存储器。
向量计算机主要采用先行控制和重叠操作技术、运算流水线、多模块交叉访问的并行存
储器等并行处理结构，从而能提高运算速度。向量型数据并行计算与流水线结构相结合，能
在很大程度上克服通常流水线计算机中指令处理量太大、存储访问不均匀、相关阻塞严重、
流水不畅等缺点，并可充分发挥并行处理结构的潜力，显著提高运算速度。
下面通过一个简单例子来说明向量处理与标量处理的差别。先考察一个C语言循环
程序：
for(i=0;i<64;i++)c[i]=a[i]+b[i];
对于该循环语句，在普通的标量处理计算机中，需要用一段循环执行的指令序列来实
现。例如，在MIPS机器上对应的指令序列如下，其中，假定i被分配在$s0,数组a、b、c的
首地址分别在$s1、$s2、$s3中，数组元素为int类型。
loop:
slti
beq
s11
addu
1w
addu
1w
add
addu
Sw
$t0,$s0,64
$t0,$zero,exit
$t0, $s0,2
$s1,Ss1,$t0
$t1, 0($s1)
$s2, 9s2,$t0
$t2, 0(§s2)
St3, $t1,St2
$s3,$s3,$t0
$t3,0($s3)
#if i<64,$t0=1;if i>=64,St0=0
#if St0=0,jump to exit
#i×4
#计算 a[i]的地址
#将a[i]取到st1中
#计算 b[i]的地址
#将 b[i]取到st2中
*计算a[i]+b[i]
#计算 c[i]的地址
#将 a[i]+b[i]的值存入c[i]


<!-- page: 352 -->

并行处理系统
exit:
addi
j
*
$s0,$s0,1
1o0P
#i=i+1
显然，上述程序在一般的标量计算机中运行时，大约需要运行768条指令。
如果一个向量计算机支持32个向量寄存器(假定在 MIPS寄存器名称前加 v_表示向
量寄存器),每个寄存器能够存放64个32位宽的数据字，一条装入/存储指令(假定分别为
v_lw和 v_sw)可以从内存连续 256个单元装入数据到一个向量寄存器或将向量寄存器内
容存到内存连续单元中，一条运算指令(加法指令假定为 v_add)可以同时对两个向量寄存
器中的64个32位字进行并行计算，并将结果存到另一个向量寄存器，则上述C语言循环程
序段对应的指令序列如下：
v_lw
$v_t1, 0(5s1)
V_lwSv t2,0($s2)
v_add$v_t3, $v_t1, sv_t2
V_sw $v_t3,0($s3)
#将 a[o]到 a[63]取到$v_t1中
#将 b[0]到 b[63]取到$v_t2中
#计算a[1]+b[i],存到sv_t3中
#将 a[i]+b[i]的值存到 c[i]中
对于同样的C语言循环程序段，在向量计算机上对应的指令只有4条。显然，对于循
环中的数组操作，向量计算机比标量计算机的执行效率要高得多。一方面，标量计算机中的
700多条指令，每一条都有取指令和译码，而向量计算机中只有4条指令要取指令和译码；
另一方面，标量计算机中加法指令的执行是串行的，而向量计算机中加法指令的执行是并行
的。此外，标量计算机中的数据相关(如load-use冒险等)会影响流水线的执行，而向量计算
机的流水线中，只可能在每个向量数据的开始处发生阻塞，而在随后的所有数据都不会发生
流水线阻塞。即使标量计算机采用循环展开技术来降低阻塞次数，也比向量计算机中的阻
塞现象要严重。
上述例子给出的是整数运算，通常向量计算机进行的都是浮点数运算。对于浮点数的
向量运算，其原理和整数向量运算是一样的。
*9.4.2Intel架构中的 SIMD 技术
在多媒体应用中，图形、图像、视频和音频处理存在大量具有共同特征的操作。例如，8
位的图像像素，16位的音频信号等都可以同时进行某种操作，这些都涉及短整数类型的并
行操作；在FIR滤波和矩阵运算中需要频繁地进行乘法累加操作；快速傅里叶变换(FFT)
和离散余弦变换(DCT)中涉及短数据的高度循环运算；三维图形和视频压缩中涉及计算密
集型的算法等。
为提高上述这些共性操作的运算能力，Intel公司在IA-32 指令系统基础上，设计了一
套新增指令集，并对CPU内部结构进行了扩充与改进，称为MMX(multimedia extension)
技术。MMX技术于1997年首次运用于P54C Pentium处理器，称之为多能奔腾，共有 57
条指令。
MMX技术中引入了新的数据类型和通用寄存器，主要数据类型为定点紧缩(packed)
整数，有以下4种新的64位数据类型：8字节、4个字(每字16位)、两个双字和一个 64位
的四字。为便于MMX指令对上述数据类型进行操作，CPU中新增了8个64位通用寄存
器 MX0～MX7。这些寄存器可用来实现数据运算，但不能用于存储器寻址。单条指令同时
第9章
335


<!-- page: 353 -->

计算机组成与系统结构(第3版)
336
并行处理多个数据元素，如8字节或4个字或两个双字或一个 64位的四字，这对提高运算
速度非常有利。例如，一条指令可以同时完成图像中8个像素的并行操作。
随着网络、通信、语音、图形、图像、动画和音/视频等多媒体处理软件对处理器性能的要
求越来越高，Intel公司在多能奔腾以后的处理器中加入了更多流式 SIMD扩展(Stream
SIMD Extension,SSE)指令集，包括SSE、SSE2、SSE3、SSSE3、SSE4等，这些是典型的数据
级并行处理技术。
SSE指令集最早是1999年由Intel公司在Pentium Ⅲ处理器中推出的，包括了70条指
令，其中包含提高3D图形运算效率的50条 SIMD浮点运算指令、12条 MMX 整数运算增
强指令和8条优化内存中连续数据块传输指令。理论上这些指令对图像处理、浮点运算、
3D运算、视频处理、音频处理等诸多多媒体应用起到全面强化的作用。SSE 兼容MMX指
令，它可以通过SIMD技术在单时钟周期内并行处理4个单精度浮点数据来有效地提高浮
点运算速度。
2001年，Intel公司在Pentium4中发布了一套包括144条新指令的 SSE2 指令集，提供
了浮点 SIMD指令、整数SIMD指令、浮点数和整数之间转换等指令。SSE2 增加了能处理
128位整数和同时并行处理两个64位双精度浮点数的指令。为了更好地利用高速缓存，还
新增了几条缓存指令，允许程序员控制已经缓存过的数据。
2004年年初，Intel公司在新款Pentium 4(P4E,Prescott 核心)处理器中发布了SSE3,
2005年4月AMD公司也发表了具备部分 SSE3功效的处理器 Athlon 64,此后的x86处理
器几乎都具备 SSE3的新指令集功能。SSE3新增了13条指令，其中一条用于视频解码，两
条用于线程同步，其余用于复杂的数学运算、浮点数与整数之间的转换以及SIMD浮点运
算，使处理器对DSP及3D处理的性能大为提升。此外，SSE3 针对多线程应用进行优化，使
处理器原有的超线程功能获得了更好的发挥。
2005年，作为SSE3指令集的补充版本，SSSE3出现在酷睿微架构处理器中，新增16条
指令，进一步增强了CPU在多媒体、图形图像和Internet 等方面的处理能力。
2008年 SSE4指令集发布，它被视为最重要的多媒体扩展指令集架构改进方式，将延续
了多年的32位架构升级至64位。SSE4增加了54条指令，其中SSE4.1指令子集包含47
条指令，SSE4.2包含7条指令。SSE4.1主要用于向量绘图运算、3D游戏加速、视频编码加
速及协同计算加速等方面的处理，此外，还加入了6条浮点运算增强指令，这使得图形渲染
处理性能和3D游戏效果等得到了极大提升。除此之外，SSE4.1指令集还加入了串流式负
载指令，可提高图形帧缓冲区的数据读取频宽，理论上可获取完整的缓存行，即单次读取64
位而非原来的8位。SSE4.2主要用于字符串和文本处理方面的操作。例如，对XML应用
进行高速查找及对比，在如Web服务器应用等方面有显著的性能改善。
下面用一个简单的例子来比较普通指令与数据级并行指令的执行速度。为了使比较结
果尽量不受访存操作的影响，以下例子中的运算操作数主要是寄存器操作数。此外，为了使
比较结果尽量准确，例子中设置了较大的循环次数值，为0x400 0000=2。例子只是为了
说明指令执行速度的快慢，并没有考虑结果是否溢出。
图9.10给出了采用普通指令的累加函数 dummy_add 对应的汇编代码，其中粗体字部
分为循环体，循环控制指令loop执行时，先检测寄存器 ECX的内容，若为0则退出循环，否
则 ECX的内容减1,并再次进入循环体的第一条指令开始执行，循环体的第一条指令地址


<!-- page: 354 -->

并行处理系统
由loop指令指出。
080484f0 <dummy_add>:
80484f0:
55
80484f1:
### 89 e5
80484f3:
b9 00 00 00 04
80484f8:
b0 01
80484fa:
b3 00
80484fc:
### 00 c3
80484fe:
e2 fc
8048500:
5d
8048501:
c3
Push
mov
mov
mov
mov
add
loop
pop
ret
8ebp
&esp,tebp
$0x4000000,8ecx
$0x1,&al
$0x0,8bl
&al,&bl
80484fc <dummy_add+0xc>
&ebp
图9.10 采用普通指令的累加函数
图9.11给出了采用SSE数据级并行指令的累加函数 dummy_add_sse 对应的汇编代
码，其中粗体字部分为循环体。
### 08048510 <dummy add_sse>:
8048510:
55
8048511:
b8 00 9d 04 10
8048516:
### 89 e5
8048518:
53
8048519;
bb 20 9d 04 14
R04851e
b9 00 00 40 00
8048523:
### 66 Of 6f 00
8048527:
### 66 0f 6f 0b
804852b:
### 66 Of fc c8
804852f:
e2 fa
8048531:
8048532:
5d
8048533:
c3
push
&ebp
mov
50x10049d00,8eax
mov
&esp,sebp
push
8ebx
mov
$0x14049d20,&ebx
mov
S0x400000,8ecx
movdqa(8eax),xmmO
movdga(?x),8xmml
paddb 8xmm0,?m1
190p
804852b <dummy_add_sse+0x1b>
pop
tebx
Sebp
pop
ret
图9.11 采用 SSE指令的累加函数
从图9.10可看出，dummy_add函数中，每次循环只完成一字节的累加，而在图9.11所
示的 dummy_add_sse函数中，每次循环执行的指令为“paddb ?m0,?m1”,即每次
循环并行完成两个XMM寄存器中的16个一字节数据的累加，对于与dummy_add 同样的
工作量，循环次数应为其1/16,即(0x4000000>>4)=0x400000=2?,因而，可以预期它所
用的时间大约只有dummy_add的1/16。
在某相同环境下测试两个函数的执行时间，dummy_add所用时间约为22.643816s,而
dummy_add_sse所用时间约为1.411588s,两者大约为16.041378倍。这与预期的结果一致。
dummy_add_sse 函数中用到的SSE指令有两种，除了paddb 以外，还有一种是 movdqa
指令，它的功能是将双四字(128位)从源操作数处移到目标操作数处。该指令可用于在
XMM寄存器与128 位存储单元之间移入/移出双四字，或在两个 XMM寄存器之间移动。
该指令的源操作数或目标操作数是存储器操作数时，操作数必须与16字节边界对齐，否则
将发生一般保护性异常(#GP)。若需要在未对齐的存储单元中移入/移出双四字，可以使
用movdqu指令。更多有关 SSE指令集的内容请参看Intel公司的相关资料。
Intel公司在 SSE指令集基础上又推出了全新的指令扩展集AVX(Advanced Vector
Extensions)。AVX是在之前 SSE指令的128 位基础上扩展到256位的 SIMD技术，同时
第9章
337


<!-- page: 355 -->

计算机组成与系统结构(第3版)
338
数据传输率也相应进行了提升，所以，从理论上看，CPU内核的浮点运算性能提升为原来的
### 2 倍。
*9.4.3 GPU 架构
Intel微处理器中的 MMX、SSE以及 AVX等采用SIMD技术的指令集主要是为了针
对向量绘图运算、3D游戏加速、视频编码加速及协同处理加速等图形处理能力的增强而提
出的。但是，CPU并不是专用于进行图形处理的，这些 SIMD指令也是在现有的x86指令
系统中增加的，因此，随着计算机游戏产业的不断发展，原来主要由CPU 进行图形处理的
方式不再适应，需要有一个功能强大的部件专门进行图形处理，这个部件就是GPU
(graphics processing unit)芯片。
GPU可以和CPU配合使用，这就是CPU+GPU的异构系统，其中CPU 用于通用处
理，而GPU专门用于图形处理，例如，3D图形的绘制、着色和渲染等。描述图形的基本数
据主要是顶点或像素，顶点由(x,y,z,w)坐标值来表示，像素则由(red,green,blue,
alpha)颜色值来表示。坐标值通常是用单精度浮点数表示，而颜色值可用整数表示，因此，
GPU需要有针对整数和浮点数的运算部件，而且对于图形处理任务来说，存在大量的数据
级并行。在 Intel架构的早期计算机系统中，CPU与GPU一般经北桥芯片通过AGP总线
或 PCI-Express总线连接，各自有独立的存储器，分别是主存和显示存储器(简称为显存)。
在有些芯片组中使用的集成GPU没有独立的显存芯片，而是直接从主存中分出一块区域
作为显存。AMD公司和Intel公司都提出了CPU和GPU融合技术，直接将CPU和GPU
通过QPI或HT总线连接并集成在一个芯片内。
自从1999年NVIDIA公司发布了第一款GPU以来，GPU得到了快速发展。为了实
时生成逼真的3D图形，GPU不仅采用了最先进的半导体制造工艺，而且在设计上也不断
创新。早期 GPU主要用于3D图形渲染等图形处理，但是，随着以CUDA为代表的 GPU
通用计算框架的普及，GPU的应用领域得到了极大的拓展，被广泛应用于石油勘测、天文计
算、流体力学模拟、分子动力学仿真、生物计算等科学计算领域，并在很多应用中获得了几十
倍乃至上百倍的加速比。因此，GPU的含义也从原来专门的图形处理器转变为GPGPU
(general-purpose computing on GPU)。GPGPU计算通常采用CPU+GPU 的异构模式，
由CPU负责执行复杂逻辑处理和事务管理等不适合数据并行的计算，而GPU负责计算密
集型的大规模数据并行计算。
下面以NVIDIA公司的 Tesla架构为例来说明GPU的大致结构。Tesla架构GPU 由
两部分组成，分别是流处理器阵列(streaming processor array,SPA)和存储控制系统。两部
分通过一个片上互连网络连接。流处理器阵列和存储控制系统都可以独立扩展，其规格可
根据市场定位进行裁剪。例如，基于GT200b芯片的产品中，GTX285的SPA中包含10个
线程处理器群(thread processor cluster,TPC)和8个存储控制器，而价格较低的GTX260+
则只有9个 TPC和7个存储控制器。
GPU中的存储控制系统包括存储控制器、光栅操作单元(raster operation processor,
ROP)和二级纹理缓存。存储控制器用于控制片外的DRAM显存，ROP则对存储器内的数
据进行颜色和深度操作，通过片上互连总线，SPA可以将计算得到的深度和颜色信息发送给
ROP,也可以直接读写DRAM显存，或者通过纹理 cache 以只读的方式访问 DRAM显存。


<!-- page: 356 -->

并行处理系统
图9.12是NVIDIA公司的 Tesla架构 GT200系统结构。一个TPC包含3个流多处理
器 SM,每个 TPC共享L1纹理 cache 和 L2 指令及常数 cache。存储控制系统中包含
256KB的L2纹理cache以及8个 GDDR3存储控制器和一个 PCI-Express 2.0×16总线接口。
每个 GDDR3存储控制器可以连接64位(8B)宽的速度为2.2GT/s的GDDR3 显示存储器。
SPA
TPC
TPC
SM
SM
SM
0
L2指令及常数
L1纹理cache
SM
SM
SM
0
1
2
L2指令及常数
L1纹理cache
存储控制系统
256KB L2 纹理 cache
8个GDDR3 存储控制器
PCl-Express 2.0×16
8×8B@2.2GT/s
2B@4GT/s
图9.12 GT200系统结构
每个 SM就是一个计算核心，图9.13给出了NVIDIA GT200 SM的内部结构。如图 9.13
所示，一个 SM包含8个标量流处理器(stream processor,SP)以及其他计算单元，如特殊功
能部件 SFU(special function unit)和双精度单元 DPU(double precision unit),SFU 用于执
行一些特殊的指令，包括超越函数、插值、倒数、平方根倒数、正弦、余弦以及其他特殊运算；
DPU 用于64位双精度浮点数和64位整数的乘加运算。
SM
填充L1
L1指令 cache
多线程指令预取和发射部件
共享存储器(可读可写)
12指令及常数cache
SFU
SPO
RF0
SP1
RF1
SP2
RF2
SP3
RF3
RF4
SP4
RF5
SP5
RF6
SP6
RF7
SP7
SFU
DPU
L1纹理cache
纹理units
填充L1
常数LI cache
装入纹理
load/store unit
第9章
339
图9.13 NVIDIA GT200的 SM 内部结构


<!-- page: 357 -->

计算机组成与系统结构(第3版)
340
通常称GPU为众核(many core)处理器，是因为通常的商业宣传中，GPU被称为拥有
上百个“核”,这里的“核”是指 SP,例如，GTX285的核心数为10×3×8=240。实际上 SP
仅是执行单元，它不是一个能完整地进行取指、译码、指令分派和处理的真正的处理核心，而
SM具有取指、指令译码和指令分派等处理功能，属于同一个 SM的所有 SP共用同一套指
令前端处理部件，也共用同一个共享存储器。
从图9.13可以看出，每个 SM相当于具有8条流水线的 SIMD处理器核心，所有 SP共
享SM中的一个指令处理前端，因此一个SP类似于多线程 CPU中的一条流水线，所以 SP
也称为线程处理器(thread processor,TP)。每个 SP核包含标量的单精度浮点数处理单元
和整数处理单元，这些单元可以实现大部分运算指令。每个 SP核支持硬件多线程，最多可
并发执行64个线程，都有各自私有的通用寄存器组(register file,RF)。每个 SP有1024个
32位通用寄存器，根据SP所分配的线程数来分配寄存器。程序会声明其寄存器需求，通常
编译器会为每个线程分配16～64个寄存器。例如，对于像素渲染程序通常使用16个或更
少的寄存器，这样，每个SP可以运行高达64个像素渲染线程。
SM的执行模型实际上与传统的 SIMD有很大区别，NVIDIA公司称其为单指令多线
程(single instruction and multiple thread,SIMT)。在传统的 SIMD架构上编程时，每条指
令能够处理的数据宽度是确定的，即向量的宽度是显式的，而在基于NVIDIA GPU 的
CUDA编程中，线程的数量是可变的，因为每个 SM包含8个SP,每个SP最多有64个线
程，所以一个 SM上最多可以有512个线程，即线程数可以在1～512 之间取值，因而提高了
灵活性。此外，在传统的 SIMD中，各个向量共享寄存器资源，并且在一条SIMD指令中不
要考虑同步问题，因为数据之间是相互独立的；而使用SIMT方式编程时，必须引入共享存
储器和同步机制实现线程之间的通信。
*9.5并行处理编程模式简介
一个应用问题要放到并行处理系统中处理，程序员首先应该分析应用问题中存在的并
行性，充分挖掘出算法中的并行处理逻辑，然后用相应的并行编程模型以及并行程序设计语
言来实现算法。
按照应用的计算特征，可以将应用问题分为以下3种类型：①数据密集型(data-
intensive)应用，即数据量巨大，但计算相对简单的应用问题，如海量网页词频统计问题；
②计算密集型(computation-intensive)应用，即数据量相对不大、但计算较为复杂的应用问
题；③数据密集与计算密集混合型应用问题，如3D电影渲染等。
针对不同特点的应用问题和不同的并行处理结构，可以有不同的并行编程模型以及并
行程序设计语言，主要分为以下几种并行处理程序设计方式。
1.共享存储变量方式
共享存储变量方式用于共享存储器多处理器系统中，采用共享存储变量的方式进行数
据交互，因此称为共享存储变量(shared memory variables)方式。该方式通常将一个任务
分解成若干处理逻辑段，每个处理逻辑段称为一个线程。线程之间可以并行执行，不同线程
之间通过共享存储变量的方式进行数据的交换。因此，也称为多线程并行程序设计方式。
多线程并行程序设计的概念提出以来，出现了很多具有代表性的并行编程接口，包括一


<!-- page: 358 -->

并行处理系统
些开源或商业版本的并行编程接口，常用的有Pthread、OpenMP和 Intel TBB等。
Pthread 是较为低层的多线性编程接口。为实现可移植的多线程程序，IEEE 1003.1c
标准定义了一个线程包，称为Pthread。它使用常规语言编程方式，通过在串行程序中加入
Pthread 函数调用来实现多线程并行处理。Pthread 中共有60多个函数调用，其中，Pthread
_create 函数用于创建一个新线程，当一个线程完成分配给它的工作后，可以用Pthread_exit
函数来终止它并释放它的栈。所有Pthread 线程都含有一个线程标识符、一组寄存器(包括
程序计数器 PC)和一组存储在结构(struct)类型变量中的属性，这些属性包括栈大小、调度
参数以及线程需要的其他信息。
OpenMP采用在编程语言中增加编译指导指令的方式来实现并行程序设计，通过在串
行程序中插入编译指导语句或命令，来指示哪个程序段可以并行执行。例如，对于一个
C/C++语言的for循环程序段，如果循环体可以并行执行，那么,就在for 语句前加入编译
指导语句“#pragma omp parallel for”,用来告诉编译器将后面的 for语句编译成一段多线
程并行代码，在编译指导语句中需要指定并行执行的线程个数、线程中的私有变量、对哪个
变量进行规约以及如何规约等信息。OpenMP中还提供了一些 API函数，包括运行环境设
置函数、锁操作函数和时间操作函数等。
2.消息传递方式
消息传递方式用于分布式存储器访问的并行处理系统(如集群)中。为了实现多个子任
务的并行处理，在不同计算节点上完成的子任务之间需要交换数据。但是，因为采用分布式
存储访问，某个计算节点无法通过执行load指令和 store指令来访问另一个节点的私有存
储器，因此，无法通过共享存储变量的方式来交换数据，只能通过消息传递(message
passing)方式进行数据的交换。显然，这种方式下，每个子任务作为一个进程在各自独立的
计算节点上并行执行，因此，消息传递方式也可以被狭义地理解为多进程并行程序设计
方式。
最常用的消息传递程序设计方式是MPI(message passing interface,消息传递接口)标
准。MPI是1993年由一组来自大学、国家实验室以及高性能计算厂商的研究人员发起和
组织的项目，1994年公布了最早的版本MPI1.0。MPI使用常规语言编程方式，通过在串
行程序中加入MPI并行编程接口函数实现多进程并行处理，所有计算节点运行同一个程
序，在并行程序段部分，不同的计算节点处理不同的数据，通过专门的API函数实现不同节
点间的消息传递。最基本的MPI函数包括 MPI_Init、MPI_Finalize、MPI_Comm_Size、MPI
_Comm_Rank、MPI_Send 和 MPI_Recv。在函数 MPI_Init 和 MPI_Finalize 之间的是一个
并行程序段，将在每个计算节点上被执行；函数MPI_Comm_Size确定指定范围内的计算节
点/进程数目；函数 MPI_Comm_Rank确定计算节点/进程的标识号；函数 MPI_Send 和
MPI_Recv用于同步通信，在不同的节点之间传递消息。MPI具有高度的可扩展性，能充分
利用系统的硬件资源，发挥其性能，被广泛应用于科学计算的各个领域。
3. MapReduce 并行程序设计方式
MapReduce 最早是由Google公司研究提出的一种面向大规模数据处理的并行处理模
式和方法。利用MapReduce,可以构建一个由普通商用服务器构成的，包含数十、数百甚至
数千个节点的基于集群的高性能并行计算平台。
MapReduce 是一个并行计算软件框架，它能自动完成计算任务的并行化处理，自动划
第9章
341


<!-- page: 359 -->

计算机组成与系统结构(第3版)
342
分计算数据和计算任务，在集群节点上自动分配和执行任务，自动收集计算结果，将数据分
布存储、数据通信、容错处理等并行处理涉及的很多底层的复杂细节交由分布式文件系统处
理，大大减轻了程序员的负担。
MapReduce也是一个并行程序设计模型与方法，借助于函数式程序设计语言Lisp 的
设计思想，它提供了一种简便的并行程序设计方法，程序员仅需要编程实现Map 和 Reduce
两个基本操作接口就可以快速完成并行化程序的设计。Map 操作主要负责对一组数据记
录进行某种重复处理，Reduce操作主要负责对 Map操作的中间结果进行某种规约并输出
结果。一个计算节点进行一个 Map或一个 Reduce操作，因此，可以有很多节点同时进行
Map 或 Reduce 操作。
4.CUDA并行程序设计方式
CUDA(compute unified device architecture)由 NVIDIA 公司在 2007年发布，是用于
众核 GPU+多核 CPU的异构系统的C/C++语言的可扩展并行编程模型。有了CUDA 编
程模型以后，GPU 除了用于图形处理以外，还可以方便地通过CUDA 编程用于其他应用
领域。
CUDA可扩展编程语言模型扩充了C/C++语言，提供了3个关键的抽象，包括层次结
构线程组、同步栅(synchronization barrier)和共享内存。
通常，一个数据并行问题可以分解成多个不同的子任务，这些子任务之间有些可以独立
并行执行，有些需要有先后处理顺序。每个子任务中的数据可以再分解成相同结构的数据
块，每个数据块还可以再分解为数据元素，同时，每个数据块可以被独立并行处理，且数据块
中的元素也可并行计算。这样，多级的“任务—数据块—数据元素”的划分正好对应于GPU
的多级处理单元“GPU—SM—SP”的划分。因此，根据这种层次化的问题划分方式，CUDA
采用了“网格(grid)—线程块(thread block)—线程”的层次结构编程模型，这就是所谓的层
次结构线程组。
程序员可以在一个串行程序中调用一个并行执行的核(kernel)函数，网格就是执行相
同核函数的一组线程块，相当于一个子处理任务。因此，一个网格由若干线程块组成，这些
线程块完成的任务就是网格所要实现的任务，各个线程块之间可以独立并行执行，没有前后
顺序关系。每个线程块又由若干线程组成，一个线程块被分配在一个 SM上执行，因此，实
际上在每个SM中执行相同的程序，程序的功能由其线程块中若干线程实现，这些线程被发
射到 SM中的 SP、SFU和 DPU等处理单元中。
因为一个数据并行问题所包含的子任务之间可能有先后顺序，所以，在对应的每个网格
之间可以根据需要通过调用_syncthreads()原语来设置同步栅，给定足够的硬件资源的话，
那些独立的网格就可以并行执行，而设置了同步栅的网格之间就按顺序执行，以保证前面网
格所包含的所有线程块在后面网格中的线程开始执行前已全部执行完。
每个线程在其执行期间可能会访问多个不同的存储器，包括局部存储器、共享存储器和
全局存储器。每个线程具有一个私有的局部存储器(local memory)用于存储那些不适合在
线程寄存器中存储的线程私有变量，包括栈帧以及在寄存器溢出时私有变量的分配。每个
线程块中的所有线程都可以访问SM中的共享存储器，其共享存储器的访问具有与线程块
相同的生命周期。此外，所有线程都可以访问全局存储器。程序可以分别使用_shared_和
device_限定词来说明变量处于共享存储器还是全局存储器。共享存储器是片上低延时存