图形硬件与GPU体系结构201004修改.pptVIP

工业界的TBR GPU 桌面领域 Larrabee（纯软件图形流水线 ） 嵌入式领域 视频游戏机 Nintendo GameCube : Flipper (ATI) Nintendo Wii Sony PS2 : Xbox360 : Xeons (ATI R500) Mobile GPU ARM Mali55 Imagination PowerVR SGX Vivante GC – ScalarMorphic 三星电子 FIMG-3DSE v1.5 三菱 Z3D 展望 TSMC 80mm2 (300mm Wafer) TSMC 140mm2 (300mm Wafer) TSMC 260mm2 (300mm Wafer) TSMC 450mm2 (300mm Wafer) 展望包括嵌入式以及可重构体系结构。 以及Memory/Power/ILP Wall * 展望包括嵌入式以及可重构体系结构。 以及Memory/Power/ILP Wall * 这种改进主要是为了增加分支能力 * 传统流计算处理器 不同应用在各个存储器层次上的带宽需求： 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 += bStep) { // Declaration of the shared memory array As used to // store the sub-matrix of A __shared__ float As[BLOCK_SIZE][BLOCK_SIZE]; // Declaration of the shared memory array Bs used to // store the sub-matrix of B __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE]; // Load the matrices from device memory // to shared memory; each thread loads // one element of each matrix AS(ty, tx) = A[a + wA * ty + tx]; BS(ty, tx) = B[b + wB * ty + tx]; // Synchronize to make sure the matrices are loaded __syncthreads(); 从Device memory到Shared memory CUDA PTX $Lt_0_2818: //loop Loop body line 71, nesting depth: 1, estimated iterations: unknown .loc 28 86 0 ld.global.f32 %f2, [%rd23+0]; st.shared.f32 [%rd14+0], %f2; .loc 28 87 0 ld.global.f32 %f3, [%rd19+0]; st.shared.f32 [%rd15+0], %f3; .loc 28 90 0 bar.sync 0; Texture Cache Access __global__ void transformKernel( 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*blockDim.y