图形硬件与GPU体系结构201004修改3说课讲解.ppt

传统流计算处理器 不同应用在各个存储器层次上的带宽需求： 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 + 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_odata[y*width + x] = tex2D(tex, tu, tv); } #endif // #ifndef _SIMPLETEXTURE_KERNEL_H_ .tex .u64 tex; mov.f32 %f132, %f80; mul.f32 %f133, %f8