University:
Boston University
Course:
MET CS 201 | Introduction to Programming
Academic year:

2022
Views:

240

Pages:

24
Author:

Haiden Dudley

Blocking/Tiling for Locality 1 Objective • Reuse each data accessed from the global memory multiple times – Across threads – shared memory blocking – Within a thread - register tiling • Register tiling is also often used to re-use computation results for increased efficiency. 2 Shared Memory Blocking Basic Idea in Global Memory Thread 1 Thread 2 … Global Memory in On-chip Memory Thread 1 Thread 2 … 3 Basic Concept of Blocking/Tiling • In a congested traffic system, significant reduction of vehicles can greatly improve the delay seen by all vehicles – Carpooling for commuters – Blocking/Tiling for global memory accesses • drivers = threads, • cars = data 4 Some computations are more challenging to block/tile than others. • Some carpools may be easier than others – More efficient if neighbors are also classmates or coworkers – Some vehicles may be more suitable for carpooling • Similar variations exist in blocking/tiling 5 Carpools need synchronization. • Good – when people have similar schedule Worker A Time Worker B sleep sleep work work dinner dinner • Bad – when people have very different schedule Worker A party time Worker B sleep sleep work work dinner 6 Same with Blocking/Tiling • Good – when threads have similar access timing Thread 1 Time Thread 2 … Thread 1 time Thread 2 • Bad – when threads have very different timing 7 Outline of Technique • Identify a block/tile of global memory content that are accessed by multiple threads • Load the block/tile from global memory into onchip memory • Have the multiple threads to access their data from the on-chip memory • Move on to the next block/tile 8 Tiled Matrix Multiply tx Nd WIDTH 0 1 2 TILE_WIDTH-1 Pd Pdsub WIDTH by 1 ty 0 1 2 2 TILE_WIDTH TILE_WIDTH – An entire row may need to be in on-chip memory – Not enough on-chip memory for large input matrices Md Thread 2 Thread 1 0 bx 1 TILE_WIDTHE • Each row of Md is accessed by multiple threads • Problem: some threads can be much further along than others 0 TILE_WIDTH-1 TILE_WIDTH TILE_WIDTH 2 WIDTH TILE_WIDTH WIDTH 9 A Small Example • Can we use two onchip memory locations to reduce the number of M accesses by the two threads? , – Not if the two threads can have very different timing! , , , ,0 Pd2,1Pd3,1 Pd0,2Pd1,2Pd2,2Pd3,2 Pd0,3Pd1,3Pd2,3Pd3,3 10 Every M and N Element is used exactly twice in generating a 2X2 tile of P P0,0 thread0,0 Access order P1,0 thread1,0 P0,1 thread0,1 P1,1 thread1,1 M0,0 * N0,0 M0,0 * N1,0 M0,1 * N0,0 M0,1 * N1,0 M1,0 * N0,1 M1,0 * N1,1 M1,1 * N0,1 M1,1 * N1,1 M2,0 * N0,2 M2,0 * N1,2 M2,1 * N0,2 M2,1 * N1,2 M3,0 * N0,3 M3,0 * N1,3 M3,1 * N0,3 M3,1 * N1,3 11 Breaking Md and Nd into Tiles Nd0,0Nd1,0 Phase 1 Nd0,1Nd1,1 Nd0,2Nd1,2 Nd0,3Nd1,3 Md0,0Md1,0Md2,0Md3,0 Md0,1Md1,1Md2,1Md3,1 Pd0,0Pd1,0Pd2,0Pd3,0 Pd0,1Pd1,1Pd2,1Pd3,1 Pd0,2Pd1,2Pd2,2Pd3,2 Pd0,3Pd1,3Pd2,3Pd3,3 12 Breaking Md and Nd into Tiles (cont.) Phase 2 Nd0,0 Nd1,0 Nd0,1 Nd1,1 Nd0,2 Nd1,2 Nd0,3 Nd1,3 Md0,0 Md1,0 Md2,0 Md3,0 Pd0,0Pd1,0Pd2,0Pd3,0 Md0,1 Md1,1 Md2,1 Md3,1 Pd0,1Pd1,1Pd2,1Pd3,1 Pd0,2Pd1,2Pd2,2Pd3,2 Pd0,3Pd1,3Pd2,3Pd3,3 13 Each phase uses one tile from Md and one from Nd Phase 1 StepPhase 4 Step 2 5 Step 6 T0,0 Md0,0 ↓ Mds0,0 Nd0,0 ↓ Nds0,0 PValue0,0 += Mds0,0*Nds0,0 + Mds1,0*Nds0,1 Md2,0 ↓ Mds0,0 Nd0,2 ↓ Nds0,0 PValue0,0 += Mds0,0*Nds0,0 + Mds1,0*Nds0,1 T1,0 Md1,0 ↓ Mds1,0 Nd1,0 ↓ Nds1,0 PValue1,0 += Mds0,0*Nds1,0 + Mds1,0*Nds1,1 Md3,0 ↓ Mds1,0 Nd1,2 ↓ Nds1,0 PValue1,0 += Mds0,0*Nds1,0 + Mds1,0*Nds1,1 T0,1 Md0,1 ↓ Mds0,1 Nd0,1 ↓ Nds0,1 PdValue0,1 += Mds0,1*Nds0,0 + Mds1,1*Nds0,1 Md2,1 ↓ Mds0,1 Nd0,3 ↓ Nds0,1 PdValue0,1 += Mds0,1*Nds0,0 + Mds1,1*Nds0,1 T1,1 Md1,1 ↓ Mds1,1 Nd1,1 ↓ Nds1,1 PdValue1,1 += Mds0,1*Nds1,0 + Mds1,1*Nds1,1 Md3,1 ↓ Mds1,1 Nd1,3 ↓ Nds1,1 PdValue1,1 += Mds0,1*Nds1,0 + Mds1,1*Nds1,1 time time 14 bx Tiled Multiply – Large Matrices 0 1 2 tx Md Nd WIDTH 0 1 2 TILE_WIDTH-1 TILE_WIDTH TILE_WIDTH • Make sure that tiles are all loaded in vertical patters from the global memory • Md data can then be accessed from shared memory in horizontal direction Pd 1 ty Pdsub WIDTH by 0 1 2 TILE_WIDTHE 0 TILE_WIDTH-1 TILE_WIDTH TILE_WIDTH 2 WIDTH TILE_WIDTH WIDTH 15 First-order Size Considerations • Assume – TILE_WIDTH of 16 gives 16*16 = 256 threads – A 1024*1024 Pd gives 64*64 = 4096 Thread Blocks • Each thread block perform 2*256 = 512 float loads from global memory for 256 * (2*16) = 8,192 mul/add operations. – Memory bandwidth no longer a limiting factor – Could use thread coarsening to further reduce traffic • Each thread block can have up to 1024 threads – Can use 32*32 tiles to further reduce traffic 16 Memory Access Pattern (Corner Turning) Md Nd WIDTH Original Access Pattern WIDTH Copy into scratchpad memory Md Tiled Access Pattern Nd Perform multiplication with scratchpad values 17 Memory Layout of a Matrix in C Access direction in Kernel code M0,0 M1,0 M2,0 M3,0 M0,1 M1,1 M2,1 M3,1 M0,2 M1,2 M2,2 M3,2 M0,3 M1,3 M2,3 M3,3 Time Period 1 T1 T2 T3 T4 Time Period 2 T1 T2 T3 T4 … M M0,0 M1,0 M2,0 M3,0 M0,1 M1,1 M2,1 M3,1 M0,2 M1,2 M2,2 M3,2 M0,3 M1,3 M2,3 M3,3 18 Memory Layout of a Matrix in C M0,0 M1,0 M2,0 M3,0 Access direction in Kernel code M0,1 M1,1 M2,1 M3,1 M0,2 M1,2 M2,2 M3,2 M0,3 M1,3 M2,3 M3,3 … Time Period 2 T1 T2 T3 T4 Time Period 1 T1 T2 T3 T4 M M0,0 M1,0 M2,0 M3,0 M0,1 M1,1 M2,1 M3,1 M0,2 M1,2 M2,2 M3,2 M0,3 M1,3 M2,3 M3,3 19 Loading a Tile • All threads in a block participate – Each thread loads one Md element and one Nd element in based tiled code • Assign the loaded element to each thread such that the accesses within each warp is coalesced 20 CUDA Code – Kernel Execution Configuration // Setup the execution configuration dim3 dimBlock(TILE_WIDTH, TILE_WIDTH); dim3 dimGrid(Width / TILE_WIDTH, Width / TILE_WIDTH); 21 Tiled Multiply Md tx m bx Pd by Pdsub k WIDTH ty 0 1 2 TILE_WIDTHE by k WIDTH Nd m 0 1 2 0 1 2 TILE_WIDTH-1 TILE_WIDTH • Each thread computes one element of Pdsub bx 1 TILE_WIDTH TILE_WIDTH • Each block computes one square sub-matrix Pdsub of size 0 TILE_WIDTH-1 TILE_WIDTH TILE_WIDTH 2 WIDTH TILE_WIDTH WIDTH 22 Tiled Matrix Multiplication Kernel __global__ void MatrixMulKernel(float* Md, float* Nd, float* Pd, int Width) { 1. 2. __shared__ float Mds[TILE_WIDTH][TILE_WIDTH]; __shared__ float Nds[TILE_WIDTH][TILE_WIDTH]; 3. 4. int bx = blockIdx.x; int by = blockIdx.y; int tx = threadIdx.x; int ty = threadIdx.y; // Identify the row and column of the Pd element to work on 5. int Row = by * TILE_WIDTH + ty; 6. int Col = bx * TILE_WIDTH + tx; 7. float Pvalue = 0; // Loop over the Md and Nd tiles required to compute the Pd element 8. for (int m = 0; m < Width/TILE_WIDTH; ++m) { // Coolaborative loading of Md and Nd tiles into shared memory 9. Mds[tx][ty] = Md[Row*Width + m*TILE_WIDTH + tx]; 10. Nds[tx][ty] = Nd[(m*TILE_WIDTH + ty) * Width + Col)]; 11. __syncthreads(); 12. for (int k = 0; k < TILE_WIDTH; ++k) 13. Pvalue += Mds[tx][k] * Nds[k][ty]; 14. __syncthreads(); 15.} ©Wen-mei W. Hwu and David Kirk/NVIDIA, 16. Pd[Row*Width+Col] = Pvalue; Chile, January 5-7, 2011 23 Shared Memory and Threading • Each SM in Fermi has 64KB on-chip SRAM, partitioned into 48KB L1 cache and 16KB shared memory, or vice versa – SM shared memory size is implementation dependent! – For TILE_WIDTH = 16, each thread block uses 2*256*4B = 2KB of shared memory. – Can potentially have up to 8 Thread Blocks actively executing • This allows up to 8*512 = 4,096 pending loads. (2 per thread, 256 threads per block) – The next TILE_WIDTH 32 would lead to 2*32*32*4B= 8KB shared memory usage per thread block, allowing 2 or 6 thread blocks active at the same time (Problem with earlier GPUs!) • Using 16x16 tiling, we reduce the accesses to the global memory by a factor of 16 – A 150GB/s bandwidth can now support (150/4)*16 = 600 GFLOPS! ©Wen-mei W. Hwu and David Kirk/NVIDIA, Chile, January 5-7, 2011 24