图形硬件与GPU体系结构04修改71张课件.ppt

1、ic.expert gmail摘要图形硬件的历史早期GPGPU通用计算 现代GPU体系结构 基于现代GPU的编程模型展望OpenGL三维图形流水线Vertex ProcessingFragment ProcessingRasterizerFramebufferTextureSGI Infinite Reality产品2019发售，SGI图形设备的巅峰之作Tile-based RenderingHP VISUALIZE FX698年商业发售分离式图形卡典范光栅加速卡(Rast)几何加速卡(GA)纹理加速卡(TA)98年商业发售Intel第一块图形卡单芯片图形卡无几何加速单元Intel 740 9

2、9年商业发售第一块集成TnL几何定点变换单元的消费级图形卡S3G Savage 2000传统GPU图形流水线在2000年前后，图形加速卡完成了从分离式元件和分离式板卡到单芯片图形硬件的整合。图形硬件翻开了新的一页，开始了可编程化的征途。u2019年DirectX7（增加TnL）u2000年 DirectX8（增加Vertex shader）u2019年 DirectX9（增加Pixel shader）u2019年 DirectX9.0c（增加动态分支）u2019年DirectX10DirectX 9.0传统GPU体系结构Split-Shader Architecture(SSA)Post-ve

3、rtex CacheHierarchical-ZFast-Z ClearZ/Color CompressionPerfetch Texture Cache Matrox Parhelia 5123Dlabs P10NV30(nVIDIA Gefroce5)DX9 GPU通用计算的开始DX9.0c-class(NV4X)GPU才引入动态分支操作(Dynamic Control)Z Buffer+Render to Texture+Clip可以用来模拟动态分支操作传统静态分支架构下的GPGPU计算基于GPU的MPEG2运动估计算法Shader Unit对通用计算支持的改进HLSL程序映射到下面三个

4、模块中：PSU(Pixel Shader Unit)TMU(Texture Mapping Unit)BP(Branch Processor)图灵完备(Turing Completeness)一条无限长的纸带 TAPE。（Data Memory）一个读写头 HEAD。（Load/Store）一套控制规则 TABLE。它根据当前机器所处的状态以及当前读写头所指的格子上的符号来确定读写头下一步的动作，并改变状态寄存器的值，令机器进入一个新的状态。(Program,包含Branch,Add)一个状态寄存器状态寄存器。(Program Counter)NV40(nVIDIA Gefroce6)NV40

5、与SGI Infinity Reality体系结构的变化？NV47(nVIDIA Gefroce7)Split Shader Unit 架构的巅峰之作多通道存储器技术？Graphics programVertex processorsFragment processorsPixel operationsOutput imageDX10 Unified Shader ArchitectureDX10带来了什么？USA（统一渲染架构）Float32精度不限制数量的Dynamic Flow Control4096 Temp Register大于64K的着色程序指令长度DX10还差什么？通信（Comm

6、unication）访存（Memory access）Shader Unit的改进基于FIFO的Shader Unit（传统GPU）Shader Unit的改进基于Thread的Shader Unit（DX10级别GPU）NV50架构图编程模型:流计算(Stream Computing)什么是流计算？流计算主要解决什么问题？NV50在传统流计算上增加了什么限制？流计算机基本概念流计算起源于传统的DSP应用，典型的应用：视频编解码数字图像处理模式识别计算机图形处理软件无线电以上流计算的特点：可实现的硬件/软件流水线传统流计算的特点：Stream processing is especially

7、suitable for applications that exhibit three application characteristics:Compute Intensity,the number of arithmetic operations per I/O or global memory reference.In many signal processing applications today it is well over 50:1 and increasing with algorithmic complexity.Data Parallelism exists in a

8、kernel if the same function is applied to all records of an input stream and a number of records can be processed simultaneously without waiting for results from previous records.Data Locality is a specific type of temporal locality common in signal and media processing applications where data is pr

9、oduced once,read once or twice later in the application,and never read again.Intermediate streams passed between kernels as well as intermediate data within kernel functions can capture this locality directly using the stream processing programming model.传统流计算处理器Imagine Stream Processor传统流计算处理器Imagi

10、ne architecture is the three tiered storage bandwidth hierarchya streaming memory system(2.1GB/s),a 128KB stream register file(25.6GB/s),direct forwarding of results among arithmetic units via local register files(435GB/s).Imagine is able to sustain performance of up to 18.3GOPS on key applications.

11、Imagine is designed to fit on a 2.56cm2 0.18um CMOS chip and to operate at 400MHz.传统流计算处理器不同应用在各个存储器层次上的带宽需求：CUDA流计算模型CUDA的存储器架构与传统的流处理器存储架构的区别？在PTX上面的体现？Shared Memory /Loop over all the sub-matrices of A and B /required to compute the block sub-matrix for(int a=aBegin,b=bBegin;a=aEnd;a+=aStep,b+=bS

12、tep)/Declaration of the shared memory array As used to /store the sub-matrix of A _shared_ float AsBLOCK_SIZEBLOCK_SIZE;/Declaration of the shared memory array Bs used to /store the sub-matrix of B _shared_ float BsBLOCK_SIZEBLOCK_SIZE;/Load the matrices from device memory /to shared memory;each thr

13、ead loads /one element of each matrix AS(ty,tx)=Aa+wA*ty+tx;BS(ty,tx)=Bb+wB*ty+tx;/Synchronize to make sure the matrices are loaded _syncthreads();从Device memory到Shared memory CUDA PTX$Lt_0_2818:/Loop body line 71,nesting depth:1,estimated iterations:unknown.loc28860ld.global.f32%f2,%rd23+0;st.share

14、d.f32%rd14+0,%f2;.loc28870ld.global.f32%f3,%rd19+0;st.shared.f32%rd15+0,%f3;.loc28900bar.sync 0;Texture Cache Access_global_ voidtransformKernel(float*g_odata,int width,int height,float theta)/calculate normalized texture coordinates unsigned int x=blockIdx.x*blockDim.x+threadIdx.x;unsigned int y=bl

15、ockIdx.y*blockDim.y+threadIdx.y;float u=x/(float)width;float v=y/(float)height;/transform coordinates u-=0.5f;v-=0.5f;float tu=u*cosf(theta)-v*sinf(theta)+0.5f;float tv=v*cosf(theta)+u*sinf(theta)+0.5f;/read from texture and write to global memory g_odatay*width+x=tex2D(tex,tu,tv);#endif/#ifndef _SI

16、MPLETEXTURE_KERNEL_H_.tex.u64 tex;mov.f32%f132,%f80;mul.f32%f133,%f82,%f73;mad.f32%f134,%f75,%f38,%f133;mov.f32%f135,0f3f000000;/0.5add.f32%f136,%f134,%f135;mov.f32%f137,0f00000000;/0mov.f32%f138,0f00000000;/0tex.2d.v4.f32.f32%f139,%f140,%f141,%f142,tex,%f132,%f136,%f137,%f138;基于AMD CAL器件的流计算模型LDS（L

17、ocal Data Share）=Memory Read and Write Cache在存储系统上与CUDA架构的区别？基于Fermi的CUDA 3.0架构如何增加层次化存储器架构的效率？如何改进Cache一致性协议？Larrabee 上的存储架构 与Fermi GPU的区别？存储系统传统GPU中的CacheFixed funtion pipelineIndexed Vertex CachePost-vertex CacheTexture Slot CacheZ/Stencil CacheColor CacheShader UnitConstant CacheTexture CacheGPU

18、中的总线哪一级流水线需要最高的优先级？3D图形硬件之前的时代Ray-casting技术代表作：德军总部3D第一代Voodoo GPU架构 4-channel memory controllerMulti-channel memory controller单一寻址空间的Video Memory伴随着Early-Z/Z cull技术 总线复杂度的压力存储需求的变更:Anisotropic FilteringMIPMAP的作用MIPMAP技术的缺陷,Normal,LOD Bias=-3 ,4X AA,8X Aniso存储需求的变更:Relief map有什么不同？Normal mapParallax

19、 mapRelief map存储需求的变更:Relief map凹凸铁图对延迟容忍的要求思考：总线带宽和总线延迟的关系？,PS2 Memory SystemGameCube Memory SystemXbox Memory SystemWii Memory SystemXbox360 Memory SystemWii Memory SystemWii Memory SystemWii Memory System处理器的融合嵌入式 GPU应用平台：视频游戏机掌上游戏机机顶盒手机航空航天应用特点：硬件资源有限特定硬件优化IMR(立即渲染模式)TBR(区域渲染模式)Mobile GPU帧缓冲尺寸：6

20、40 x 480Tile尺寸：（16，32，64)x（16，32，64）最佳TILE尺寸：32 x 16 像素区域传统TBR GPU体系结构工业界的TBR GPU桌面领域Larrabee（纯软件图形流水线）嵌入式领域 视频游戏机Nintendo GameCube:Flipper (ATI)Nintendo WiiSony PS2:Xbox360:Xeons(ATI R500)Mobile GPUARM Mali55Imagination PowerVR SGXVivante GC ScalarMorphic三星电子 FIMG-3DSE v1.5三菱 Z3D展望 TSMC 80mm2(300mm Wafer)TSMC 140mm2(300mm Wafer)TSMC 260mm2(300mm Wafer)TSMC 450mm2(300mm Wafer)谢谢