第四讲多核技术

第四讲多核技术

4.1协处理器技术

4.2多核关键技术

4.3异构多核技术14.1协处理器技术21.IBM的Cell技术多核架构从通用的对等设计迁移到“主核心+协处理器”的非对等设计IBM的Cell为索尼PS3游戏机定制拥有9个硬件核心的多核处理器1个核心主处理器拥有完整的功能是PowerPC970的精简版本主要职能就是负责任务的分配实际的浮点运算由协处理器来完成8个协处理器专门用于浮点运算所需的运算规则非常简单，只要CPU运行频率足够高，Cell就能够获得惊人的浮点效能考虑以协处理器为中心优化设计？34Cell性能Cell起步频率即达到4GHz256Gigaflops浮点运算能力，接近超级计算机的水准每个SPE协处理器拥有4路并行的整数/浮点单元，每个运算周期又可执行4次32位浮点运算一个时钟周期可执行两个运算操作每个Cell拥有8枚SPE协处理器，工作频率假设在4GHz，此时

Cell所具有的浮点效能就是2×4×8×4GHz＝256Gigaflops英特尔的4路Montecito安腾（双内核）系统也仅获得45Gigaflops的浮点性能处理器名称集成的晶体管数（个）运算能力（GFlops）PowerPC9705800万9.0IBMCellDD1（第一代）:2.34(亿)DD2（第二代）:2.5(亿)2565Cell的新颖设计思想Cell的高效能很大程度上来自于其新颖的设计思想：主处理器与协处理器各司其职内核设计精简高效，以实现高频运作运算单元则采用128位并行结构分布式中：系统架构、调度等能否借鉴该思想？6高度适应性的Cell及所构成的Cell计算网络Cell可以对处理内核的数量进行任意裁减嵌入式设备：只有单个核心，工作在较低的频率，较低的能耗便携电脑和桌面PC：可使用与PS3游戏机一样的标准Cell，或者对核心进行适当裁减工作站/服务器系统：可以将两枚Cell处理器直接集成在一起以获得更高的效能大型计算机：可配置成包含四枚独立Cell处理器的“MCM模块(Multi-chipmodule)”，具有每秒万亿次浮点的运算能力分布式计算系统：利用Cell的超高速度FlexIO芯片连接总线，将不同计算设备联成一体，实现运算能力与内存资源的分享智能空间、普适计算用它的结构思想如何？7Cell处理器的整体架构8PPE处理单元以IBM的Power

4处理器为基础(POWER4——于2001年发布，POWER4是第一款频率超过1GHz的POWER处理器。POWER4实现了POWERISAv.2.00)可支持同步多线程技术该处理单元内置了32KB一级缓存和512KB二级缓存规格与同出一脉的PowerPC

970处理器类似9SPE协处理器运算处理单元：4个32位浮点运算单元，4个32位整数运算单元寄存器：128bit×128bit局部缓存：256KB流水线长度：18级输入总线：128bit宽度的总线3条输出总线：128bit宽度的总线1条XDR就是“eXtremeDataRate”的缩写，这是Rambus的黄石的最终名称10SPE使用对用户不透明用户程序或数据不得大于256KB：代码+数据+栈总大小不能超过256K.SPE使用显式控制的片内局部存储器LS（LocalStore）代替cache来简化设计用户程序可直接对256KB局部缓存访问，象局部存储器LS是非一致性的，象分布式存储程序员或编译器可以在SPE进行计算的同时，显式地安排LS和主存之间数据的移动或使用代码超出LS大小，使用overlay技术（sdk支持）定态置换局部缓存和共享存储内容数据超出大小，可定制的软件cache（sdk支持）来实现对数据的访问或预取除了sdk支持，上面的问题还可以手工用DMA来解决，关于DMA操作sdk也提供了封装。

据说IBM的编译器已会直接支持内存使用超过256K的限制两次编译：SPE加速程序编译+PPE通用程序编译，然后连接生成可执行代码；HybridProgramming需要用户或编译器优化使用，使得性能可以发挥到及至硬件电路简单，可以做到高速，低功耗第三方软件必须移植11PPE/SPE单元的内部联结内部有一条

768bit位宽的“EIB单元互连总线环（Element

Interconnect

BUS

Ring，EIB

Ring）”是一个强大的内部总线控制逻辑—Cell内所有的功能单元都通过EIB总线环连接在一起，包括PPE、八个SPE、XDR内存控制器以及外部总线接口采用的是全双工的128bit连接总线若Cell工作在4GHz频率各个功能单元便都拥有

4GHz×128bit／Hz×2（全双工）÷8Byte/bit＝128GBps带宽12CELL运行模式与常规的双核处理器不同，Cell内的九个核心具有相当强的独立性PPE处理单元的任务是运行操作系统应用程序相关的线程运算完全由

SPE协处理器运行多个应用程序的线程被平均分布到各个SPE中，整套系统负载均衡网络上的相互协作可以接受并执行相关来自Cell计算网络中其他设备的计算请求，结果再通过网络传输给任务发起者网络上的任务可以被均匀分散到所有的Cell处理器上，达到最佳的最短完成时间13Cell的功耗分析工作频率为4GHz时每个SPE协处理器的工作电压高于1.1V，功耗只有4瓦所有SPE协处理器的功耗总和最高也不过4瓦×8＝32瓦PPE

处理单元的核心部分，功耗水平也会控制在很低的水平Cell运算部分的功耗水平会在40瓦左右，即便加上缓存单元整体功耗也可控制在较好的水平上若频率降到3GHz，工作电压只需要0.9V，功耗只有

2瓦将频率降低到2GHz，每个SPE的功耗仅有1瓦在2006年，IBM采用了更先进的65纳米技术来制造Cell，将具有更加出色的功耗水平14XDR内存控制器与FlexIO前端总线整合XDR内存控制器以及采用FlexIO前端总线是Cell的两大技术亮点系统拥有25.6GBps的内存带宽Cell与3.2GHz的XDR模组搭配16bit×4通道×3.2Gbps÷8Byte/bit＝25.6GBps的内存带宽跨平台的内存管理组建由多个Cell设备组成的计算网络所有的Cell处理器的内存资源在逻辑上可形成一个有机整体，无需任何修改即可直接协作FlexIO前端总线采用6组8位全双工配置，有效带宽76.8GBps提供给I/O芯片：上下行各12.8GBps提供给PS3的图形处理器或其他Cell处理器：上下行各25.6GBps1516Cell发展前瞻可作为替代X86的下一代计算平台X86也许要十几年后X86处理器才可能达到Cell所具有的运算性能Cell要在短时间内取代X86绝非易事最大的问题在于Cell的软件平台尚未成熟尤其是分布式计算架构需要软件针对性的优化最关键的操作系统方面，IBM选择了开源的Linux在过去数年间，IBM花费大量的资金和人力推动Linux系统的开发，并将自己在UNIX领域的研究成果无偿贡献出来172005年，和Mercury计算机系统公司合作，制造基于Cell处理器的电脑2005年七月，推出了Cell的一个开发包，让第三方厂商将Cell应用到其他领域2006年9月12日，发布了首款基于Cell协处理器的BladeCenterQS20刀片式服务器两个cell处理器(每个频率为2.3GHz)；1GB的内存(每个处理器为512MB)；40GB的硬盘；两个千兆以太网接口；一或两个同PCI—Express(2002年的总线和接口标准)相容的Infiniband适配器。操作系统采用RedhatFedora5LinuxOS182006年九月，IBM宣布将负责为美国能源部国家核安全管理局建造名为Roadrunner的超级计算机采用基于Cell处理器的BladeCenterH系统和基于AMD双核心皓龙处理器的Systemx3755利用新开发的HybridProgramming和AMDTorrenza软件连接不同平台持久运行速度：每秒钟1000万亿次计算（1Petaflop）预期峰值：1.6PFlops/s走鹃19Cell各阶段的整体结构202.Intel的ManyCore在2005年的IDF（英特尔开发技术论坛)技术峰会上，英特尔对外公布了ManyCore超多核发展蓝图采用类似Cell的专用化结构将对等设计的四核心处理器中的某一个或几个内核置换为若干数量的DSP逻辑保留下来的X86核心：执行所有的通用任务以及对特殊任务的分派DSP阵列：用于某些特殊任务的处理。MultiCoreManyCore21DSP类型：依照应用不同，可以是Java解释器、MPEG视频引擎、存储控制器、运算处理器等等处理这类任务时，DSP的效能远优于通用的X86核心，功耗也很低DSP与通用CPU的效率对比：DSP功耗仅2瓦，通用CPU75瓦在处理对应任务时，DSP芯片的效能比通用CPU更加出色由于DSP构造简单、频率提升非常容易主核心的运算压力就大大减轻，系统整体效能将获得明显提升高负载的专用任务转交给DSP执行22ManyCore发展路线图第一代ManyCore架构处理器将采用“3个通用X86核心+16个DSP内核”的组合原型是一枚四核心处理器，只是将其中一个核心置换成16个DSP逻辑而已，因此处理器的总体结构和晶体管规模都不会有多大变化，但产品的实际水准将获得大幅度增强执行Java程序、视频解码、3D渲染等耗用CPU资源的任务中，DSP的效能都大幅优于通用核心，因此ManyCore产品在执行这类专用任务时会有飞跃性的性能增益。DSP逻辑的能耗只有通用核心的几十分之一，可以让处理器的功耗出现可观的降低23第二代ManyCore产品将面世拥有8个通用X86核心、64个专用DSP逻辑，片内缓存容量高达1GB，晶体管规模则达到200亿逐步引入ManyCoreArray架构，不断增强DSP的数量以及执行能力，通用核心的地位将随着时间推移不断减弱，直到最后完全可能实现以DSP占主导地位的专用化运算模式通用芯片，但能包括各种应用吗？灵活性不高，功耗和效能方面估计会不理想24ManyCore处理器与Cell相比Cell的主核心非常简单，协处理器则非常强大ManyCore的通用核心仍然居于主导地位，DSP更多只是一种辅助这种差异源自于二者不同的定位Cell只要求具备强劲的浮点效能，而对整数运算不作要求，因此通用的主核心可以非常精简，没有要求通用性和兼容性ManyCore必须考虑兼容大量的X86应用软件，专用的任务居于从属性地位，仍然要求通用性和兼容性253.AMD多核技术尚未考虑类似英特尔的ManyCore计划利用现有的HyperTransport连接架构，对多路服务器系统进行拓展如，Cray公司（著名的高性能计算机制造商）希望能在基于Opteron的超级计算机中使用矢量，以提升计算机的矢量运算效能.皓龙AMD并不是简单考虑在Opteron核心中增加一个针对应用的矢量逻辑处理单元，而是计划建立一个以AMD为中心的企业生态圈

26AMD的企业生态圈AMD公司考虑的一套协处理器扩展方案为多路Opteron平台开发各种功能的协处理器协处理器都通过HyperTransport总线与Opteron处理器直接连接如，对SeymourCray提出的需求，AMD给出的解决方案就是，将八路Opteron中的一颗Opteron处理器置换成矢量协处理器，以此实现矢量计算性能的大幅度增长，而Opteron平台本身不需要作任何形式的变动这种拓展架构也可以延伸到PC领域例如在PC中挂接基于HyperTransport总线的浮点协处理器、视频解码器、专门针对Java程序的硬件解释器，甚至可以是由nVIDIA或ATI开发的图形处理器为达成上述目标，AMD必须设计出一个高度稳定的统一接口，方便用户进行扩展借助各种各样的协处理器，AMD64系统的性能将获得空前强化。27AMD的企业生态圈将HyperTransport协处理器授权给其他的专业IC设计公司第三方公司都可以为AMD64平台开发协处理器并分别销售AMD自身只需要负责通用处理器的开发和HyperTransport原生态的维护基于共同的利益，大量的第三方IC设计公司将紧密围绕AMD公司共同发展HyperTransport平台终端用户，可以在现有基础上通过增加或升级协处理器达成大幅度提高系统性能的目的28AMD与Intel多核技术比较策略不同：一个或多个通用处理器+众多的专用处理单元Intel采用ManyCore计划，技术越来越封闭，商业利润都归自己所有AMD利用现有的HyperTransport连接架构，对多路服务器系统进行拓展，建立企业生态圈，倡导“友好生态系统”实质上完全相同，两者的区别更多是在物理组成方式ManyCore将专用的DSP逻辑直接整合于处理器内部一个或多个通用处理器+众多的专用处理单元(DSP)AMD的协处理器系统则是借助HyperTransport总线在外部挂接，用户不必为了获得额外的性能购买新机器，直接选择相应的协处理器挂接即可外挂协处理器=一个或多个简单通用处理器+众多的专用处理单元(DSP)294.GPU及其应用30GPU显卡构成1.显示处理芯片：GPU2.显存：与内存统一编址用来存储所处理的数据信息和像素信息GPU处理完数据后将数据再送到显存3.数模转换器RAMDACA：从显存中读取数据，将数字信号转换为模拟信号，输出到显示屏31

显示芯片GPU传统的GPU：专为图形处理设计1.数据汇集阶段：传进3D图形的基本图元——三角片2.顶点处理模块：对图元的顶点进行几何变换，如平衡、旋转和投影。3.光栅化（栅格）过程：对三角片进行插值得到其内部像素点坐标4.像素点处理模块：对光栅化的像素点进行渲染，如光照效果、纹理映射等。5.像素绘制（渲染）：将像素点写入帧缓冲并显示在屏幕上。32通用处理GPUGPGPU的发展传统GPU中的像素点处理模块，数据计算量非常大，需要大量的计算单元，可以使大量的计算单元并行执行，从而进行加速。GPU作为通用计算就是各级流水线模块变为可编程单元的过程。现代的GPU现在所使用的GPU相对于传统的GPU有了许多改进，以Geforce8800为例，它内部大致有以下的部件：SP：StreamingProcessorSM：StreamingMulti-ProcessorTPC：由两个SM组成纹理处理单元ROP：rasteroperationprocessor光栅其它；33流处理器SP定义：streamingprocessor，nVIDIA对其统一架构GPU内通用标量着色器的称谓，它是新一代的渲染技术指标作用：处理由CPU传输来的数据，转化为可以辨识的数字信号与以往GPU渲染架构不同：传统GPU渲染架构采用固定的比例组成渲染单元模式SP可按任意比例组成，可根据需要设置比例，从而解决顶点与像素着色任务不均时所造成的浪费，效率更高34G80显卡中，每个SP里含有:1个IntegerScalarUnit1个32位的MAD(乘加)ScalarUnit1个32位的MULScalarUnit—一个附加的乘法单元SP的指令执行延迟一般是4个周期，执行流水线长度有4级每个SP同时执行4条相同的指令，以充分消除延迟在G80中每个SM中有8个SP，所以SM可以同时执行32个指令（不同数据的相同指令），即CUDA中的warp单元35SM流多处理器定义：StreamingMulti-ProcessorG80中有128个SP每8个SP一组，构成16个SM每个SM是一个SIMD的多处理器执行同样一组指令，处理不同数据块各SM之间独立每个SM有8个SP32KB的Register为64KB)16KB的SharedMemory共享内存的访问是cache级别SFU：特殊功能单元，用来执行一些特殊的指令Cache分开或复杂Cache结构有没意义？36线程程序Warp块内指令按顺序执行线程指令获取并进行warp调度Warp间不按顺序调度warp执行分配sp单元执行37TPC（与通用计算无关）TPC：纹理处理单元主要用于图形流水处理由2个SM组成作用：在图形绘制时，可将邻接的4个像素绑定处理，每个像素4个分量，也就需要16个处理单元，即两个SM每个SM内有8KB的TextureL1Cache纹理8KB的InstructionL1Cache8KB的ConstantL1Cache

常数在SM外部还有256KB的TextureL2Cache256KB的InstructionL2Cache256KB的ConstantL2Cache由于其特殊的硬件组合配置，是否可以作为某些特殊算法处理？38ROP&MemoryROP定义：光栅操作处理器(RasterOperationProcessor)MSAA:多重取样抗锯齿模式CSAA:覆盖取样抗锯齿模式SSAA:超级取样抗锯齿模式GDDR3：显卡的双倍速率同步动态存储与显卡本地内存操作直接相关的色彩以及深度帧缓存数据，由内存控制器和固定功能的光栅操作器（ROP）完成391、输入装配器：依照输入指令流进行顶点工作收集处理2、顶点处理工作分配给空闲的TPC4、完成顶点光栅pixelfragment的处理(像素)5、将pixelfragment分配给空闲TPC，完成像素渲染pixel处理7.把computethreadarray分发给空闲TPC3、TPC执行结果存GPU的片上缓存；再把数据转发到viewport/clip/setup/raster/zcull模块6、ROP完成色彩、深度以及可能的多彩样计算40GPU运行原理1.Inputassembler(输入装配器)依照输入指令流进行顶点工作viewwork任务的收集处理。2.Viewworkdistribution负责把viewwork打包交付给空闲的TPC，TPC就会执行vertexshader程序。顶点着色(Vertex

Shader)是一段执行在GPU上的程序3.TPC执行的结果被写到GPU的片上缓存buffer，这些buffer会把数据转发到viewport/clip/setup/raster/zcull模块，完成顶点光栅化pixelfragment的处理。4.这些pixelfragment会被pixelworkdistribution递交给空闲的TPC完成pixel运算。5.Computeworkdistribution负责把computethreadarray分发给空闲的TPC。6.完成了pixelshader处理的pixelfragment会被Interconnectionnetwork传至下游的ROP完成色彩，深度以及可能的多彩样计算。41GeForce8800系统特性全局存储器Globalmemroy:

片外存储、共768MB、200-300时钟周期延迟、可读可写大容量DRAM核心计算任务执行前，所有数据要存放这里核心计算任务使用指针直接访问常数和纹理数据（线程）后备存储多线程同时执行能更有效使用通过访问存储邻接单元，可以使硬件访问到相同页数据42共享存储器Sharedmemory:

片内存储、每个SM16KB、寄存器级延迟、可读可写局部中间结果暂存器，被线程块(threadblock)中所有线程共享组织为16banks，减少或消除线程和数据组织访问冲突43常数存储器Constantmemory:

片内cache、共64KB、寄存器级延迟、只读每SM8KBcache，原始数据存放在全局存储器中(globalmemory)其64KB限制由编程模型设定经常作为查找表使用单访问端口，SM中同时的访问请求必须是相同地址，否则会有额外延时44文理存储器Texturememory:

片内cache、访问大小可到全局存储器、延迟大于100时钟周期、只读每2个SM16KBcache，原始数据在全局存储器中2D定位可用于完成硬件插补和纹理边界处可重构回收，在视频编码一般应用中很有用局部存储器Localmemory:

片外，访问大小可到全局存储器、访问延迟同全局存储器、可读可写作为寄存器溢出(Registerspilling)空间，存放多于寄存器的数据

45Grid/Block/thread

应用划分为：栅格/块/线程执行线程被组织成三级层次：

–栅格(Grid)：最高层，由核心任务(kernel)创建，由多个块组成

–块(Block)：每个块包含最大线程数为512（与共享存储大小和线程长度有关），每个线程块分配到一个独立的SM进行连续执行

–线程(Thread)：同一个线程块中的线程可通过共享存储器共享数据，通过调用__syncthreads()执行同步操作46GeForce8800andCUDA约定资源配置参数

限制每个SM的线程数≤768每个SM

的线程块数≤8每个SM

的32-bit寄存器数≤8192(8192*32)/(8*1024)=32KB每个SM

的共享存储器大小

16KB每个线程块的线程数≤51247访问存储空间主机存储和显卡设备有着独立分开的存储空间，但主机内统一编址。主机可以访问所有空间，显卡只能访问显存CUDA中：Devicecodecan:

–R/Wper-threadregisters–R/Wper-threadlocal-memory–R/Wper-blockshared-memory–R/Wper-gridglobal-memory–Rper-gridconstantmemoryHostcodecan:

–R/Wper-gridglobalandconstantmemories显卡存储器48矩阵乘的串行实现voidcomputeGold(float*P,float*M,float*N,constintwidth){for(unsignedinti=0;i<width,i++)//therowindexofM{for(unsignedintj=0;j<width,j++)//thecolumnindexofN{doublesum=0;for(intk=0;k<width;k++)//thecolumnindexofMandtherowindexofN{doublem=M[width*i+k];doublen=N[k*width+j];sum+=m*n;//themultiplyofthecorrespondingvalue}//theresultofthePelementwhichwithrowindexiandcolumnindexNP[width*i+j]=(float)sum;}}}49GPU环境下用CUDA实现矩阵相乘需要三个步骤：

1.将数据从hostmemory拷贝到devicememory；

2.内核程序kernel执行；

3.将结果从devicememory拷贝回hostmemory。

intmain(void){1.//AllocateandinitializethematricesM,N,P//copymatricesdatafromhostmemorytodevicememory…….2.//M*NonthedevicematrixMul_kernel(…..);3.//copymatrixdatafromdevicememorytohostmemroy…….//Freematrices}501)、数据的转移hostmemroy与devicememory之间数据的转移：

devicememory的分配：

cudaMalloc(void**devPtr,size_tcount);data的转移：

cudaMemcpy(void*dst,constvoid*src,size_tcount,enumcudaMemcpyKindkind);

cudaMemcpyKind:cudaMemcpyDeviceToHost,cudaMemcpyHostToHostcudaMemcpyHostToDevice,cudaMemcpyDeviceToDevice512)、分配执行单元及内核函数的调用Executionparametersassignment:

dim3dimBlock(threadIdx.x,threadIdx.y,threadIdx.z);dim3dimGrid(blockIdx.x,blockIdx.y);Theschedulingofthekernelfunction:function_kernel<<<dimGrid,dimBlock>>>(参数0,参数1,参数2,…)52blockIdandthreadIdblockId和threadId由CUDAruntime分配给每个thread.blockId:grid的每个维度可以设置最大值为65536.

即gridDim.x和gridDim.y是从1到65536的任意值。

blockIdx.x:0→gridDim.x-1blockIdx.y:0→gridDim.x-1threadId:thread各维度由三个值决定，由于每个threadblock的值最大为512。即：threadIdx.x*threadIdx.y*threadIdx.z<=512

53内核函数kernel的实现计算维度均为（WIDTH,WIDTH）的两个WIDTH阶方阵A和B的乘积C，可按如下的方法划分为几个线程：每个线程块均负责计算C的一个子方阵Csub；块内的每个线程负责计算Csub的一个元素。在这里，我们选取BLOCK的维度为BLOCK_SIZE。Csub为以下两个长方矩阵的乘积：(WIDTH,BLOCK_SIZE)的A的子阵(BLOCK_SIZE,WIDTH)的B的子阵54matrixMul_kernelcodewithoutsharedmemory:

__global__voidmatrixMul_kernel(float*C,float*A,float*B,intwidth){intbx=blockIdx.x;//Blockindexintby=blockIdx.y;inttx=threadIdx.x;//Threadindexintty=threadIdx.y;introw=by*BLOCK_SIZE+ty;//computetherowindexoftheCelementintcol=bx*BLOCK_SIZE+tx;//computethecolumnindexoftheCelementfloatCsub=0;for(intk=0;k<width;k++)//everythreadcomputeoneelementofC{Csub+=A[row][k]*B[k][col];__sycnthreads();}C[row][col]=Csub;}voidcomputeGold(float*P,float*M,float*N,constintwidth){for(unsignedinti=0;i<width,i++)//therowindexofM{for(unsignedintj=0;j<width,j++)//thecolumnindexofN{doublesum=0;for(intk=0;k<width;k++)//thecolumnindexofMandtherowindexofN{doublem=M[width*i+k];doublen=N[k*width+j];sum+=m*n;//themultiplyofthecorrespondingvalue}//theresultofthePelementwhichwithrowindexiandcolumnindexNP[width*i+j]=(float)sum;}}}矩阵乘的串行实现55matrixMul_kernelcodewithsharedmemory:

__global__voidmatrixMul_kernel(float*C,float*A,float*B,intwidth){…………//thesameasmatrixMul_kernelcodewithoutsharedmemoryfloatCsub=0;for(intm=0;m<width/BLOCK_SIZE;m++)//everythreadcomputeoneelementofC{__shared__floatAs[BLOCK_SIZE][BLOCK_SIZE];//allocatesharedmemoryformatrixAandB__shared__floatBs[BLOCK_SIZE][BLOCK_SIZE];As[ty][tx]=A[row][m*BLOCK_SIZE+tx];//colabrativelyloadAandBblocksintosharedmemoryBs[ty][tx]=B[m*BLOCK_SIZE+ty][col];__syncthreads();for(intk=0;k<BLOCK_SIZE;k++){Csub+=As[ty][k]*B[k][tx];__syncthreads();}}C[row][col]=Csub;}voidcomputeGold(float*P,float*M,float*N,constintwidth){for(unsignedinti=0;i<width,i++)//therowindexofM{for(unsignedintj=0;j<width,j++)//thecolumnindexofN{doublesum=0;for(intk=0;k<width;k++)//thecolumnindexofMandtherowindexofN{doublem=M[width*i+k];doublen=N[k*width+j];sum+=m*n;//themultiplyofthecorrespondingvalue}//theresultofthePelementwhichwithrowindexiandcolumnindexNP[width*i+j]=(float)sum;}}}矩阵乘的串行实现56GPU单精度矩阵相乘（张保）硬件：GPU：Tesla1060；CPU：IntelXeon1024*1024,2048*2048,3072*3072？，4096*4096CPU串行运行时间测试GPU并行运行时间测试并行时间加速比57GPU矩阵相乘：单精度和双精度测试（张保）

测试时间包括：原始数据从主机拷贝到设备时间、数据在设备上的计算时间、结果从设备拷贝到主机时间数据从主机拷贝到设备时间数据在设备上的计算时间数据从设备拷贝到主机时间单精度/双精度=1/358常见算法（闫旭东）

矩阵相乘sharememory；1024*1024矩阵GPU比cpu快522倍一维拉普拉斯方程GPU使用全局存储器和共享存储器时，分别是CPU运算速率的25和58倍简单偏微分方程本次测试计算量比较小GPU比CPU快46倍左右59FFT(离散傅氏变换的快速算法)时间测试（张保）测试时间记录(oursfftongpuandcpu)时间比值(ours-fft-on-gpu/cpu)测试时间记录(oursfftongpuandcufft)时间比值(ours-fft-on-gpu/cufft)60*协处理器总结协处理器的出现是必然趋势很多特定的需求驱动如，由于Java和XML的驱动，数据中心的客户需要这方面的专用设计，而在高性能的计算领域则对向量浮点有强烈需求，媒体流的发展需要有更高性能的媒体处理能力。设计越来越模块化，以便满足市场上出现的更多的商机61协处理器总结协处理器在系统中的存在方式分为很多种放在高速IO总线的扩展槽当中放在主处理器的插槽中AMDHyperTranspor所采用的方式采用定制化的ASIC，ApplicationSpecificIntegratedCircuit

，采用专门的加速技术加上标准的接口，将第三方的协处理器组成阵列与主处理器封装在一起，采用进一步提高集成度Intel的ManyCore与IBM的CellCell异常强大的前端总线设计，令其在组成协处理器阵列方面比ManyCore更具有优势编程方法和模式上相同？从编程方法和模式角度统一抽象硬件架构？62协处理器总结协处理器的大规模使用，对系统软件和应用软件提出了更高的要求，其中，以编译优化技术和任务调度为主挖掘程序中所蕴含的并行性：编译技术如何利用多处理器结构的并行计算能力和高度通信机制，充分发挥单芯片多处理器的结构优势，是一个比较关键的问题多处理器的协同工作：如何将多个任务分配到处理器资源，尽量提高资源的利用率，实现多个处理器之间的动态负载平衡尤其在大规模异构处理器条件下，如何协调多个处理器协调工作，是提高整个系统性能的关键问题之一在这方面IBM开发的“HybirdProgramming”技术可资借鉴。体系结构方面主要是如何解决存储、IO屏障635.2多核处理器的九大关键技术

64与单核处理器相比，多核处理器在体系结构、软件、功耗和安全性设计等方面面临着巨大的挑战，但也蕴含着巨大的潜能。多线程Simultaneous

multithreading，简称SMT。SMT可通过复制处理器上的结构状态，让同一个处理器上的多个线程同步执行并共享处理器的执行资源，可最大限度地实现宽发射、乱序的超标量处理，提高处理器运算部件的利用率，缓和由于数据相关或Cache未命中带来的访问内存延时。当没有多个线程可用时，SMT处理器几乎和传统的宽发射超标量处理器一样。SMT最具吸引力的是只需小规模改变处理器核心的设计，几乎不用增加额外的成本就可以显著地提升效能。多线程技术则可以为高速的运算核心准备更多的待处理数据，减少运算核心的闲置时间。这对于桌面低端系统来说无疑十分具有吸引力。Intel从3.06GHz

Pentium

4开始，所有处理器都将支持SMT技术。

Intel的hyper-threading其实就是

two-threadSMT.65

SMT：SMT是一个相对廉价的技术，比如英特尔的Hyper-Threading，允许每个物理核心运行两个同步线程。SMT的设计思想是充分利用每个核心的资源。如果一个物理核心只有一个执行线程，那么在等待内存中的关键代码或数据的时候，线程处于停顿状态，这样核心的利用率是低下的。而SMT技术允许一个物理核心运行两个或更多的线程，可以根据当前的状况动态进行切换，如果一个线程处于停顿状态等待内存，另一个线程的指令则可以使用这个物理核心的所有执行单元，让物理核心利用的更加充分。66为了让SMT正常工作，处理器的所有代码和存储部分需要被复制或分区。例如，一个双线程SMT处理器需要两套架构寄存器和重命名寄存器，一套给线程A，一套给线程B。另外组成指令窗口的共享指令队列要具备很大的空间，这样指令窗口才能容纳足够多的来自两个线程的指令，让执行单元可以保持在忙碌状态。最后，两个线程任何共享单元，比如处理管线不同部分的指令缓存，都不能被任一个线程独占。换句话说，SMT核心的两个线程需要和另一个紧密的共享资源，保证核心的缓存单元不会空置没有线程利用。67多核心

(CMP就是multi-core)

多核心，也指单芯片多处理器（Chip

multiprocessors，简称CMP）。CMP是由美国斯坦福大学提出的，其思想是将大规模并行处理器中的SMP（对称多处理器）集成到同一芯片内，各个处理器并行执行不同的进程。与CMP比较，

SMT处理器结构的灵活性比较突出。但是，当半导体工艺进入0.18微米以后，线延时已经超过了门延迟，要求微处理器的设计通过划分许多规模更小、局部性更好的基本单元结构来进行。相比之下，由于CMP结构已经被划分成多个处理器核来设计，每个核都比较简单，有利于优化设计，因此更有发展前途。目前，IBM

的Power

4芯片和Sun的

MAJC5200芯片都采用了CMP结构。多核处理器可以在处理器内部共享缓存，提高缓存利用率，同时简化多处理器系统设计的复杂度。

68CMP：CMP的方式非常直接，简单来说，CMP是通过“复制”物理核心来扩展处理器在多线程软件中的性能，这是获得最佳性能一种最简单和最有效的方式。但CMP的缺点是制造成本很昂贵，并且也要受到处理器制造工艺的限制，毕竟不能将芯片做的越来越大。并且CMP的方式对负载要求也很高，只有经过适当并行优化的负载才能充分发挥CMP的性能，很多核心的CMP常常会浪费资源，在一些应用中，主频更高、结构更简单的双核和四核处理器就往往可以获得更好的性能。69与单核处理器相比，多核处理器在体系结构、软件、功耗和安全性设计等方面面临着巨大的挑战，但也蕴含着巨大的潜能。CMP和SMT一样，致力于挖掘计算的并行性。CMP：随着大规模集成电路技术发展，在芯片容量足够大时，就可以将大规模并行处理机结构中的SMP（对称多处理机）或DSM（分布共享处理机）节点集成到同一芯片内，各个处理器并行执行不同的线程或进程。基于SMP结构的单芯片多处理机中，处理器之间通过片外Cache或者是片外的共享存储器来进行通信。基于DSM结构的单芯片多处理器中，处理器间通过连接分布式存储器的片内高速交叉开关网络进行通信。由于SMP和DSM已经非常成熟，CMP结构设计比较容易，只是芯片制造工艺的要求较高。因此，CMP最先被应用于商用高性能CPU。多核利用集成度提高性能，但原来系统级的一些问题引入到了处理器内部。701、核结构研究:同构还是异构CMP的构成分成同构和异构两类同构是指内部核的结构是相同的异构是指内部的核结构是不同的面对不同的应用研究核结构的实现对性能至关重要核本身的结构，关系到整个芯片的面积、功耗和性能怎样继承和发展传统处理器的成果，直接影响多核的性能和实现周期根据Amdahl定理，程序的加速比决定于串行部分的性能从理论上来看似乎异构微处理器的结构具有更好的性能核所用的指令系统对系统的实现也是很重要的多核之间采用相同的指令系统还是不同的指令系统能否运行操作系统等712、程序执行模型选择程序执行模型是多核处理器设计的首要问题程序执行模型的适用性决定多核处理器能否以最低的代价提供最高的性能程序执行模型是编译器设计人员与系统实现人员之间的接口编译器设计人员决定如何将一种高级语言程序按一种程序执行模型转换成一种目标机器语言程序系统实现人员则决定该程序执行模型在具体目标机器上的有效实现当目标机器是多核体系结构时，产生的问题是:多核体系结构如何支持重要的程序执行模型？是否有其他的程序执行模型更适于多核的体系结构？这些程序执行模型能多大程度上满足应用的需要并为用户所接受？723、Cache设计:

多级Cache设计与一致性问题处理器和主存间的速度差距对CMP来说是个突出的矛盾必须使用多级Cache来缓解目前有共享一级Cache的CMP、共享二级Cache的CMP以及共享主存的CMP通常，CMP采用共享二级Cache的CMP结构，即每个处理器核心拥有私有的一级Cache，且所有处理器核心共享二级Cache，基本与目前SMP架构相同

多级Cache设计：Cache自身的体系结构设计直接关系到系统整体性能在CMP结构中共享Cache或独有Cache孰优孰劣需不需要在一块芯片上建立多级Cache建立几级Cache等等对整个芯片的尺寸、功耗、布局、性能以及运行效率等都有很大的影响一致性问题：多级Cache又引发一致性问题采用何种Cache一致性模型和机制都将对CMP整体性能产生重要影响在传统多处理器系统结构中广泛采用的Cache一致性模型有:顺序一致性模型、弱一致性模型、释放一致性模型等。相关的Cache一致性机制主要有：总线的侦听协议和基于目录的目录协议目前CMP系统大多采用基于总线的侦听协议734、核间通信技术CMP处理器的各CPU核心执行的程序之间有时需要进行数据共享与同步硬件结构必须支持核间通信，高效的通信机制是CMP处理器高性能的重要保障目前比较主流的片上高效通信机制有两种：基于总线共享的Cache结构；基于片上的互连结构总线共享Cache结构指每个CPU内核拥有共享的二级或三级Cache，保存比较常用的数据通过连接核心的总线进行通信，共享数据优点：结构简单，通信速度高缺点：基于总线的结构可扩展性较差基于片上互连的结构指每个CPU核心具有独立的处理单元和Cache各CPU核心通过交叉开关或片上网络等方式连接各CPU核心间通过消息通信优点：可扩展性好，数据带宽有保证缺点：硬件结构复杂，且软件改动较大两者竞争结果不是互相取代而是互相合作例如：在全局范围采用片上网络，而局部采用总线方式，达到性能与复杂性的平衡745、总线接口设计传统微处理器中，Cache不命中或访存事件都会对CPU的执行效率产生负面影响，而总线接口单元（BIU）的工作效率会决定此影响的程度多个CPU核心同时要求访问内存或多个CPU核心内私有Cache同时出现Cache不命中事件时，设计高效的BIUBIU对这多个访问请求的仲裁机制BIU对外存储访问的转换机制决定了CMP系统的整体性能设计高效的BIU，是CMP处理器研究的重要内容高效的多端口总线接口单元（BIU）结构，将多核心对主存的单字访问转为更为高效的猝发（burst）访问对CMP处理器整体效率最佳的一次Burst访问字的数量模型、高效多端口BIU访问的仲裁机制756、操作系统设计:

任务调度、中断处理、同步互斥对于多核CPU，优化操作系统任务调度算法是保证效率的关键一般任务调度算法有全局队列调度和局部队列调度全局队列调度指操作系统维护一个全局的任务等待队列，当有一个CPU核心空闲时，从全局任务等待队列中选取就绪任务开始在此核心上执行优点：CPU核心利用率较高局部队列调度指操作系统为每个CPU内核维护一个局部的任务等待队列，当有一个CPU内核空闲时，从该核心的任务等待队列中选取恰当的任务执行优点：任务基本上无需在多个CPU核心间切换，有利于提高CPU核心局部Cache命中率目前多数多核CPU操作系统采用基于全局队列的任务调度算法多核的中断处理和单核有很大不同多核的各处理器之间需要通过中断方式进行通信多个处理器之间的本地中断控制器和负责仲裁各核之间中断分配的全局中断控制器，需要封装在芯片内部多核CPU是一个多任务系统不同任务会竞争共享资源，需要系统提供同步与互斥机制传统的用于单核的解决机制不能满足多核，需要利用硬件提供“读－修改－写”的原子操作或其他同步互斥机制保证767、低功耗设计低功耗和热优化设计已经成为微处理器研究中的核心问题微处理器的集成度越来越高，表面温度变得越来越高，呈指数级增长每三年，处理器功耗密度翻一番CMP是多核心结构，集成度更高，功耗问题更严重低功耗设计，多层次问题在操作系统级、算法级、结构级、电路级等研究每个层次的低功耗设计方法实现的效果不同抽象层次越高，功耗和温度降低的效果越明显778、存储器墙芯片内核要充分工作，要求芯片能提供与芯片性能相匹配的存储器带宽内部Cache的容量能解决一些问题其他一些提高存储器接口带宽手段增加单个管脚带宽的DDR、DDR2、QDR、XDR等必须有能提供高带宽的存储器芯片对封装的要求也越来越高管脚数每年以20%的数目提升，产生了成本提高的问题怎样提供一个高带宽，低延迟的接口，是必须解决的一个重要问题789、可靠性及安全性设计随着技术革新的发展，处理器的应用渗透到现代社会的各个层面，但是在安全性方面却存在着很大的隐患。一方面，处理器结构自身的可靠性低下，由于超微细化与时钟设计的高速化、低电源电压化，设计上的安全系数越来越难以保证，故障的发生率逐渐走高。另一方面，来自第三方的恶意攻击越来越多，手段越来越先进，已成为具有普遍性的社会问题。现在，可靠性与安全性的提高在计算机体系结构研究领域备受注目。79思考题讨论总结已有异构多核计算机系统架构，从编程模型和应用程序设计角度，给出一个合适的异构多核系统硬件统一架构逻辑模型805.3异构多核技术

81异构多核架构IBMCellCPU+GPU82异构多核存储结构异构多核架构Cell和CPU+GPU都采用了层次架构

通用处理单元可以自由访问主存协处理单元可以自由高速访问本地存储和寄存器，通过DMA访问批量访问主存协处理单元的本地存储大小有限

Cell的SPE只有256KB的LSGPU中的SP只有16KB的本地存储采用软件显式管理对应用访存的优化决定应用的效率和性能

典型异构多核架构存储层次83CELL异构多核并行计算84定制技术与流特征识别相结合的访存技术访存技术框架

剖分信息①计算子任务信息：记录和统计对应子任务的运行时间与执行频度信息；由Pthread管理层负责收集②数据块信息：记录对应数据块的上下界、大小及访问频度等属性；在通用处理单元侧内存管理库函数中增加剖分功能及从第三类信息中统计获得③子任务-数据信息：记录相应计算子任务在相应内存块上的访问行为及使用软件管理cache时的统计信息；由访存层负责收集多方位运行时剖分信息85第2列：主存的访问次数。对于频繁访问的数据块应该作为优化的重点第3列：相邻两次访存的步长。矩阵A、C的访问步长为1，表示SPE按顺序访问矩阵A和C内存块，这种情况下单缓冲区cache可以有效的工作矩阵B访问步长值较大而稳定，表示对B的访问为规则非顺序访存，适宜使用组相联cache或多路缓冲区预取第4、5列：数据块大小可用于数据预取的越界检查访问区域大小可用于指导访存方式的选择，例如在(a)中，虽然矩阵A、B、C均大于LS的容量，但是计算任务只用到每个大块中的4KB数据，因此可以采用批量访存来代替软件管理cache来挖掘性能基于这些信息可以确定访存策略（1）单计算核心测试86访存方法：I表示批量访存S表示单缓冲区cache4表示4路缓冲区预取优化后的执行时间为优化前的6%-37%

87（2）多计算核心测试选用1024×1024分块稀疏矩阵LU分解，分块大小设为256×256四个函数不经优化，全部使用单缓冲区cache访存技术下载到SPE，由Pthread层剖分取得的子任务的执行情况逐步将负荷下载到SPE的过程中，不使用剖分优化和使用剖分优化的执行时间比较执行时间是原来的23%左右882、非规则应用剖分优化应用模拟计算8000个粒子受力及运动情况，执行3000个计算步来统计系统的静态和动态特征对大的数组进行访问，无法用ALF、CellSs等编程模型直接实现或移植由Pthread层得到MD模拟程序执行“热点”剖分信息，可知优化的优化顺序为ComputeForces()、BuildNebrList()和LeapfrogStep()ComputeForces()和BuildNebrList()对数组mol访问频度很高，且步长是不确定的，使用组相联cache可以更加充分地复用LS中的数据；BuildNebrList()中cellList数组也为非顺序访问，可配置为组相联cache分析剖分信息和目前LS的使用情况，cellList可进一步使用批量访存，将其全部一次读入LS以改善性能89实验结果1：最节约LS资源的单缓冲区cache移植后的MD模拟执行结果实验结果2：首先将mol的访问配置为组相联cache，MD模拟执行时间是实验结果1的65%实验结果3：再将cellList的访问配置为组相联cache，MD模拟执行时间是实验结果2的70%实验结果4：cellList进一步使用批量访存，MD模拟执行时间是实验结果3的10%左右实验结果4执行时间是实验结果1的40%左右非规则应用优化测试9091基于渐近性能优化支持多种应用特征支持复杂非规则计算可将优化策略以配置和静态编译的方法予以“固化”循环渐近性能优化方法91面向CellBE架构的支持剖分的S2S编译器c语言源代码转换成可在CellBE架构下编译运行的代码结合运行时动态剖分机制，记录优化策略信息及效果，以充分利用架构提供的软件显式管理资源控制能力9293规则应用性能测试pthread层小访存应用性能测试MPI层小访存应用性能测试94复杂非规则应用性能测试复杂非规则应用性能测试复杂非规则应用性能多级优化未优化:默认配置是单缓冲95GPU加速并行计算CUDA程序执行模式96面向GPU的性能优化方法研究采用三级优化策略，以渐近的方式完成应用在该异构并行系统上的性能提升。三级优化：访存级优化、内核加速级优化、数据划分级优化。97三级渐近优化策略的设计与实现访存级优化：分析应用变量的访存特性，分配合适的存储空间，减小访存延迟。内核加速级优化：整合线程块、采用数据预取等，开发GPU计算能力、进一步减小访存延迟。数据划分级优化：对主机-设备端带宽以及GPU计算能力进行分析，将数据进行合理划分进行处理，实现数据拷贝与GPU计算时间重叠。98矩阵相乘(三级优化)

浪潮NF5588服务器，4核Xeon，12GB内存，NVIDIATesla