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

1、图形硬件与GPU体系结构ic.expert gmail图形硬件与GPU体系结构ic.expert gmail摘要图形硬件的历史早期GPGPU通用计算 现代GPU体系结构 基于现代GPU的编程模型展望摘要图形硬件的历史OpenGL三维图形流水线Vertex ProcessingFragment ProcessingRasterizerFramebufferTextureOpenGL三维图形流水线Vertex Fragment RSGI Infinite Reality产品2019发售，SGI图形设备的巅峰之作Tile-based RenderingSGI Infinite Reality产品20

2、19发售，SHP VISUALIZE FX698年商业发售分离式图形卡典范光栅加速卡(Rast)几何加速卡(GA)纹理加速卡(TA)HP VISUALIZE FX698年商业发售98年商业发售Intel第一块图形卡单芯片图形卡无几何加速单元Intel 740 98年商业发售Intel 740 99年商业发售第一块集成TnL几何定点变换单元的消费级图形卡S3G Savage 200099年商业发售S3G Savage 2000传统GPU图形流水线在2000年前后，图形加速卡完成了从分离式元件和分离式板卡到单芯片图形硬件的整合。图形硬件翻开了新的一页，开始了可编程化的征途。2019年DirectX

3、7 （增加TnL ）2000年 DirectX8（增加Vertex shader ）2019年 DirectX9 （增加Pixel shader ）2019年 DirectX9.0c （增加动态分支 ）2019年DirectX10DirectX 9.0传统GPU图形流水线在2000年前后，图形加速卡完成了从分离传统GPU体系结构Split-Shader Architecture (SSA)Post-vertex CacheHierarchical-ZFast-Z ClearZ/Color CompressionPerfetch Texture Cache 传统GPU体系结构Split-Shad

4、er ArchitecMatrox Parhelia 512Matrox Parhelia 5123Dlabs P103Dlabs P10NV30 (nVIDIA Gefroce5)NV30 (nVIDIA Gefroce5)DX9 GPU通用计算的开始DX9.0c-class (NV4X) GPU才引入动态分支操作(Dynamic Control) Z Buffer + Render to Texture + Clip可以用来模拟动态分支操作DX9 GPU通用计算的开始DX9.0c-class (传统静态分支架构下的GPGPU计算基于GPU的MPEG2运动估计算法传统静态分支架构下的GPGP

5、U计算基于GPU的MPEG2运动Shader Unit对通用计算支持的改进HLSL程序映射到下面三个模块中：PSU (Pixel Shader Unit)TMU (Texture Mapping Unit)BP (Branch Processor)Shader Unit对通用计算支持的改进HLSL程序映射到图灵完备(Turing Completeness)一条无限长的纸带 TAPE。（Data Memory）一个读写头 HEAD。（Load/Store）一套控制规则 TABLE。它根据当前机器所处的状态以及当前读写头所指的格子上的符号来确定读写头下一步的动作，并改变状态寄存器的值，令机器进入一

6、个新的状态。(Program,包含Branch, Add)一个状态寄存器。(Program Counter)图灵完备(Turing Completeness)一条无限长NV40 (nVIDIA Gefroce6)NV40与SGI Infinity Reality体系结构的变化？NV40 (nVIDIA Gefroce6)NV40与SGINV47 (nVIDIA Gefroce7)Split Shader Unit 架构的巅峰之作多通道存储器技术？Graphics programVertex processorsFragment processorsPixel operationsOutput

7、imageNV47 (nVIDIA Gefroce7)Split ShDX10 Unified Shader ArchitectureDX10带来了什么？USA（统一渲染架构）Float32精度不限制数量的Dynamic Flow Control4096 Temp Register大于64K的着色程序指令长度DX10还差什么？通信（Communication）访存（Memory access）DX10 Unified Shader ArchitecShader Unit的改进基于FIFO的Shader Unit （传统GPU）Shader Unit的改进基于FIFO的Shader UnShad

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

9、essing is especially 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 P

10、arallelism exists in a 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 applic

11、ations where data is produced 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.传统流计算的特点：Stream pro

12、cessing is 传统流计算处理器Imagine Stream Processor传统流计算处理器Imagine Stream Process传统流计算处理器Imagine 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 fi

13、les (435GB/s). Imagine is able to sustain performance of up to 18.3GOPS on key applications.Imagine is designed to fit on a 2.56cm2 0.18um CMOS chip and to operate at 400MHz. 传统流计算处理器Imagine architecture i传统流计算处理器不同应用在各个存储器层次上的带宽需求：传统流计算处理器不同应用在各个存储器层次上的带宽需求：CUDA流计算模型CUDA的存储器架构与传统的流处理器存储架构的区别？在PTX上面

14、的体现？CUDA流计算模型CUDA的存储器架构与传统的流处理器存储架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 += bStep) / Declaration of the shared memory array As used to / store the sub-matrix of A _shared_ float AsB

15、LOCK_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 thread loads / one element of each matrix AS(ty, tx) = Aa + wA * ty + tx; BS(ty, tx) = Bb + wB *

16、 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.shared.f32 %rd14+0, %f2;.loc28870ld.global.f32 %f3, %rd19+0;st.sha

17、red.f32 %rd15+0, %f3;.loc28900bar.sync 0;Shared Memory / Loop over alTexture 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 = blockIdx.y*bloc

18、kDim.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(t

19、ex, tu, tv);#endif / #ifndef _SIMPLETEXTURE_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,%f14

20、1,%f142, tex,%f132,%f136,%f137,%f138;Texture Cache Access_global_基于AMD CAL器件的流计算模型LDS（Local Data Share） = Memory Read and Write Cache在存储系统上与CUDA架构的区别？基于AMD CAL器件的流计算模型LDS（Local Dat基于Fermi的CUDA 3.0架构如何增加层次化存储器架构的效率？如何改进Cache一致性协议？基于Fermi的CUDA 3.0架构如何增加层次化存储器架Larrabee 上的存储架构 与Fermi GPU的区别？Larrabee 上的存

21、储架构 与Fermi GPU的区别图形硬件与GPU体系结构04修改71张课件存储系统传统GPU中的CacheFixed funtion pipelineIndexed Vertex CachePost-vertex CacheTexture Slot CacheZ/Stencil CacheColor CacheShader UnitConstant CacheTexture Cache存储系统传统GPU中的CacheGPU中的总线哪一级流水线需要最高的优先级？GPU中的总线哪一级流水线需要最高的优先级？3D图形硬件之前的时代Ray-casting技术代表作：德军总部3D3D图形硬件之前的时代

22、Ray-casting技术第一代Voodoo GPU架构 4-channel memory controller第一代Voodoo GPU架构 4-channel memMulti-channel memory controller单一寻址空间的Video Memory伴随着Early-Z/ Z cull技术 总线复杂度的压力Multi-channel memory controlle存储需求的变更 : Anisotropic FilteringMIPMAP的作用MIPMAP技术的缺陷存储需求的变更 : Anisotropic Filterin ,Normal ,Normal ,LOD Bia

23、s = - 3 ,LOD Bias = - 3 ,4X AA ,4X AA ,8X Aniso ,8X Aniso存储需求的变更 : Relief map有什么不同？Normal mapParallax mapRelief map存储需求的变更 : Relief map有什么不同？存储需求的变更 : Relief map凹凸铁图对延迟容忍的要求思考：总线带宽和总线延迟的关系？存储需求的变更 : Relief map凹凸铁图对延迟容忍的 , ,PS2 Memory SystemPS2 Memory SystemGameCube Memory SystemGameCube Memory SystemXbox Memory SystemXbox Memory SystemWii Memory SystemWii Memory SystemXbox360 Memory SystemWii Memory SystemWii Memory SystemWii Memory SystemXbox360 Memory SystemW