Module 4.3 - Memory Model and Locality Tiled Matrix Multiplication - PowerPoint PPT Presentation

GPU Teaching Kit Accelerated Computing Module 4.3 - Memory Model and Locality Tiled Matrix Multiplication

Objective – To understand the design of a tiled parallel algorithm for matrix multiplication – Loading a tile – Phased execution – Barrier Synchronization 2

Matrix Multiplication – Data access pattern N – Each thread - a row of M and a column of N – Each thread block – a strip of M and a WIDTH strip of N M P BLOCK_WIDTHE WIDTH Row BLOCK_WIDTH WIDTH WIDTH Col 3

Tiled Matrix Multiplication – Break up the execution of each N thread into phases – so that the data accesses by the thread block in each phase are WIDTH focused on one tile of M and one tile of N – The tile is of BLOCK_SIZE elements in each dimension M P BLOCK_WIDTHE WIDTH Row BLOCK_WIDTH WIDTH WIDTH Col 4

Loading a Tile – All threads in a block participate – Each thread loads one M element and one N element in tiled code 5 5

Phase 0 Load for Block (0,0) N 0,0 N 0,1 N 0,2 N 0,3 N 0,0 N 0,1 Shared Memory N 1,0 N 1,1 N 1,2 N 1,3 N 1,0 N 1,1 N 2,0 N 2,1 N 2,2 N 2,3 N 3,0 N 3,1 N 3,2 N 3,3 Shared Memory P 0,0 P 0,1 P 0,2 P 0,3 M 0,0 M 0,1 M 0,2 M 0,3 M 0,0 M 0,1 P 1,0 P 1,1 P 1,2 P 1,3 M 1,0 M 1,1 M 1,2 M 1,3 M 1,0 M 1,1 P 2,0 P 2,1 P 2,2 P 2,3 M 2,0 M 2,1 M 2,2 M 2,3 P 3,0 P 3,1 P 3,2 P 3,3 M 3,0 M 3,1 M 3,2 M 3,3 6

Phase 0 Use for Block (0,0) (iteration 0) N 0,0 N 0,1 N 0,2 N 0,3 N 0,0 N 0,1 Shared Memory N 1,0 N 1,1 N 1,2 N 1,3 N 1,0 N 1,1 N 2,0 N 2,1 N 2,2 N 2,3 N 3,0 N 3,1 N 3,2 N 3,3 Shared Memory P 0,0 P 0,1 P 0,2 P 0,3 M 0,0 M 0,1 M 0,2 M 0,3 M 0,0 M 0,1 P 1,0 P 1,1 P 1,2 P 1,3 M 1,0 M 1,1 M 1,2 M 1,3 M 1,0 M 1,1 P 2,0 P 2,1 P 2,2 P 2,3 M 2,0 M 2,1 M 2,2 M 2,3 P 3,0 P 3,1 P 3,2 P 3,3 M 3,0 M 3,1 M 3,2 M 3,3 7

Phase 0 Use for Block (0,0) (iteration 1) N 0,0 N 0,1 N 0,2 N 0,3 N 0,0 N 0,1 Shared Memory N 1,0 N 1,1 N 1,2 N 1,3 N 1,0 N 1,1 N 2,0 N 2,1 N 2,2 N 2,3 N 3,0 N 3,1 N 3,2 N 3,3 Shared Memory P 0,0 P 0,1 P 0,2 P 0,3 M 0,0 M 0,1 M 0,2 M 0,3 M 0,0 M 0,1 P 1,0 P 1,1 P 1,2 P 1,3 M 1,0 M 1,1 M 1,2 M 1,3 M 1,0 M 1,1 P 2,0 P 2,1 P 2,2 P 2,3 M 2,0 M 2,1 M 2,2 M 2,3 P 3,0 P 3,1 P 3,2 P 3,3 M 3,0 M 3,1 M 3,2 M 3,3 8

Phase 1 Load for Block (0,0) N 0,0 N 0,1 N 0,2 N 0,3 N 1,0 N 1,1 N 1,2 N 1,3 N 2,0 N 2,1 N 2,2 N 2,3 N 2,0 N 2,1 Shared Memory N 3,0 N 3,1 N 3,2 N 3,3 N 3,0 N 3,1 Shared Memory P 0,0 P 0,1 P 0,2 P 0,3 M 0,0 M 0,1 M 0,2 M 0,3 M 0,2 M 0,3 P 1,0 P 1,1 P 1,2 P 1,3 M 1,0 M 1,1 M 1,2 M 1,3 M 1,2 M 1,3 P 2,0 P 2,1 P 2,2 P 2,3 M 2,0 M 2,1 M 2,2 M 2,3 P 3,0 P 3,1 P 3,2 P 3,3 M 3,0 M 3,1 M 3,2 M 3,3 9

Phase 1 Use for Block (0,0) (iteration 0) N 0,0 N 0,1 N 0,2 N 0,3 N 1,0 N 1,1 N 1,2 N 1,3 N 2,0 N 2,1 N 2,2 N 2,3 N 2,0 N 2,1 Shared Memory N 3,0 N 3,1 N 3,2 N 3,3 N 3,0 N 3,1 Shared Memory P 0,0 P 0,1 P 0,2 P 0,3 M 0,0 M 0,1 M 0,2 M 0,3 M 0,2 M 0,3 P 1,0 P 1,1 P 1,2 P 1,3 M 1,0 M 1,1 M 1,2 M 1,3 M 1,2 M 1,3 P 2,0 P 2,1 P 2,2 P 2,3 M 2,0 M 2,1 M 2,2 M 2,3 P 3,0 P 3,1 P 3,2 P 3,3 M 3,0 M 3,1 M 3,2 M 3,3 10

Phase 1 Use for Block (0,0) (iteration 1) N 0,0 N 0,1 N 0,2 N 0,3 N 1,0 N 1,1 N 1,2 N 1,3 N 2,0 N 2,1 N 2,2 N 2,3 N 2,0 N 2,1 Shared Memory N 3,0 N 3,1 N 3,2 N 3,3 N 3,0 N 3,1 Shared Memory P 0,0 P 0,1 P 0,2 P 0,3 M 0,0 M 0,1 M 0,2 M 0,3 M 0,2 M 0,3 P 1,0 P 1,1 P 1,2 P 1,3 M 1,0 M 1,1 M 1,2 M 1,3 M 1,2 M 1,3 P 2,0 P 2,1 P 2,2 P 2,3 M 2,0 M 2,1 M 2,2 M 2,3 P 3,0 P 3,1 P 3,2 P 3,3 M 3,0 M 3,1 M 3,2 M 3,3 11

Execution Phases of Toy Example 12

Execution Phases of Toy Example (cont.) Shared memory allows each value to be accessed by multiple threads 13

Barrier Synchronization – Synchronize all threads in a block – __syncthreads() – All threads in the same block must reach the __syncthreads() before any of the them can move on – Best used to coordinate the phased execution tiled algorithms – To ensure that all elements of a tile are loaded at the beginning of a phase – To ensure that all elements of a tile are consumed at the end of a phase 14

GPU Teaching Kit Accelerated Computing The GPU Teaching Kit is licensed by NVIDIA and the University of Illinois under the Creative Commons Attribution-NonCommercial 4.0 International License.

Module 4.3 - Memory Model and Locality Tiled Matrix Multiplication - PowerPoint PPT Presentation

GPU Teaching Kit Accelerated Computing Module 4.3 - Memory Model and Locality Tiled Matrix Multiplication Objective To understand the design of a tiled parallel algorithm for matrix multiplication Loading a tile Phased execution

Module 4.1 Memory and Data Locality CUDA Memories Objective To learn to effectively use

Module 4.5 - Memory and Data Locality Handling Arbitrary Matrix Sizes in Tiled Algorithms

Lecture 13 The C++ Memory model Synchronization variables Implementing synchronization

Cache Management Improving Memory Locality and Reducing Memory Latency Introduction Memory

Memory Hierarchy: Caching CSE 141, S2'06 Jeff Brown The memory subsystem Computer Control

The Memory Hierarchy Today Storage technologies and trends Locality of reference

Cache Memory Raul Queiroz Feitosa Content Memory Hierarchy Principle of Locality Some

1 Locate set Cache Read Example: Direct Mapped Cache (E = 1) Check if any line in set

Abstractions for Practical Systems Caching and the memory hierarchy Operating systems and the

Module 2: Loops Thomas Schwarz, SJ Repetition Computational model for kindergardeners a=3

Lightning Talk: The C++11 Memory Model Meeting C++ 2012 Neuss, Germany Presented by Marc Mutz

1 Memory Read Transaction (1) Memory Read Transaction (2) CPU places address A on the memory

1 Memory Read Transaction (1) Memory Read Transaction (2) CPU places address A on the memory

Lecture 4: Locality and parallelism in simulation I David Bindel 6 Sep 2011 Logistics

Virtual Memory ! Overview / Motivation ! Simple Approach: Overlays ! Locality of

Chapter 5 Large and Fast: Exploiting Memory Hierarchy 5.1 Introduction Principle of Locality

Towards Exploiting Data Locality for Irregular Applications on Shared-Memory Multicore

Memory Hierarchy and Direct Map Caches Lecture 11 CDA 3103 06-25-2014 5.1 Introduction

Real clocks: a toy model for non-locality luis j. garay Universidad Complutense de Madrid

CONTEXT LOCALITY LOCALITY LOCALITY LOCALITY LAYOUTS M E E R L U S T R O A D PICK

Function Pointers Refined Memory Model 1 The C0 Memory Model so far Local Memory Allocated

Memory Hierarchy: Cache Memory hierarchy Cache basics Locality Cache organization Cache-aware

Memory Hierarchy Design Chapter 5 and Appendix C 1 Overview Problem CPU vs Memory

Introduction to OpenMP Lecture 8: Memory model Why do we need a memory model? On modern