STRUCTURAL AGNOSTIC SPMV ADAPTING CSR ADAPTIVE FOR IRREGULAR

  • Slides: 32
Download presentation
STRUCTURAL AGNOSTIC SPMV: ADAPTING CSR -ADAPTIVE FOR IRREGULAR MATRICES MAYANK DAGA AND JOSEPH L.

STRUCTURAL AGNOSTIC SPMV: ADAPTING CSR -ADAPTIVE FOR IRREGULAR MATRICES MAYANK DAGA AND JOSEPH L. GREATHOUSE AMD RESEARCH ADVANCED MICRO DEVICES, INC.

THIS TALK IN ONE SLIDE Demonstrate how to save space and time by using

THIS TALK IN ONE SLIDE Demonstrate how to save space and time by using the CSR format for GPU-based Sp. MV Dynamic GPU-based Sp. MV algorithm that is efficient for regular and irregular matrices 2 x faster than CSR-Adaptive 36% faster than CSR 5 at 10 x less overheads 2 | STRUCTURAL AGNOSTIC SPMV: ADAPTING CSR-ADAPTIVE FOR IRREGULAR MATRICES | DECEMBER 17, 2015 | 2015 HIPC | BENGALURU, INDIA |

SPARSE MATRIX-VECTOR MULTIPLICATION (SPMV) 1. 0 - 2. 0 - 3. 0 1 22

SPARSE MATRIX-VECTOR MULTIPLICATION (SPMV) 1. 0 - 2. 0 - 3. 0 1 22 - 4. 0 - 5. 0 - 2 28 - - 6. 0 - - 7. 0 - - 8. 0 - - 9. 0 - - - × 3 4 5 = 18 39 18 1*1 + 2*3 + 3*5 Traditionally Bandwidth Limited 3 | STRUCTURAL AGNOSTIC SPMV: ADAPTING CSR-ADAPTIVE FOR IRREGULAR MATRICES | DECEMBER 17, 2015 | 2015 HIPC | BENGALURU, INDIA |

Performance of Sp. MV is dominated by how the sparse matrix is stored in

Performance of Sp. MV is dominated by how the sparse matrix is stored in memory 4 | STRUCTURAL AGNOSTIC SPMV: ADAPTING CSR-ADAPTIVE FOR IRREGULAR MATRICES | DECEMBER 17, 2015 | 2015 HIPC | BENGALURU, INDIA |

COMPRESSED SPARSE ROW (CSR) 1. 0 - 2. 0 - 3. 0 - 4.

COMPRESSED SPARSE ROW (CSR) 1. 0 - 2. 0 - 3. 0 - 4. 0 - 5. 0 - - - 6. 0 - - 7. 0 - - 8. 0 - - 9. 0 - - - row delimiters: 0 3 5 6 8 9 columns: 0 2 4 1 3 2 0 3 1 values: 1. 0 2. 0 3. 0 4. 0 5. 0 6. 0 7. 0 8. 0 9. 0 5 | STRUCTURAL AGNOSTIC SPMV: ADAPTING CSR-ADAPTIVE FOR IRREGULAR MATRICES | DECEMBER 17, 2015 | 2015 HIPC | BENGALURU, INDIA |

CSR-SCALAR One thread/work-item per matrix row Workgroup #1 Workgroup #2 Workgroup #3 Uncoalesced accesses

CSR-SCALAR One thread/work-item per matrix row Workgroup #1 Workgroup #2 Workgroup #3 Uncoalesced accesses 6 | STRUCTURAL AGNOSTIC SPMV: ADAPTING CSR-ADAPTIVE FOR IRREGULAR MATRICES | DECEMBER 17, 2015 | 2015 HIPC | BENGALURU, INDIA |

CSR-VECTOR One workgroup (multiple threads) per matrix row Workgroup #1 Workgroup #2 . .

CSR-VECTOR One workgroup (multiple threads) per matrix row Workgroup #1 Workgroup #2 . . . Workgroup #12 Bad: poor parallelism Good: coalesced accesses 7 | STRUCTURAL AGNOSTIC SPMV: ADAPTING CSR-ADAPTIVE FOR IRREGULAR MATRICES | DECEMBER 17, 2015 | 2015 HIPC | BENGALURU, INDIA |

MAKING SPMV FAST ON GPUS Neither CSR-scalar nor CSR-vector offer ideal performance Traditional solution:

MAKING SPMV FAST ON GPUS Neither CSR-scalar nor CSR-vector offer ideal performance Traditional solution: new matrix storage format + new algorithms ‒ More than fifty new storage formats have been proposed rewrite software space and time overhead Can we find a good GPU algorithm that does not modify the original CSR data? 8 | STRUCTURAL AGNOSTIC SPMV: ADAPTING CSR-ADAPTIVE FOR IRREGULAR MATRICES | DECEMBER 17, 2015 | 2015 HIPC | BENGALURU, INDIA |

Automatically adapts to different problems in different domains CSR-Adaptive

Automatically adapts to different problems in different domains CSR-Adaptive

INCREASING BANDWIDTH EFFICIENCY CSR-vector coalesces long rows by streaming into the LDS 10 |

INCREASING BANDWIDTH EFFICIENCY CSR-vector coalesces long rows by streaming into the LDS 10 | STRUCTURAL AGNOSTIC SPMV: ADAPTING CSR-ADAPTIVE FOR IRREGULAR MATRICES | DECEMBER 17, 2015 | 2015 HIPC | BENGALURU, INDIA |

INCREASING BANDWIDTH EFFICIENCY How can we get good bandwidth for other “short” rows? Load

INCREASING BANDWIDTH EFFICIENCY How can we get good bandwidth for other “short” rows? Load these into the LDS 11 | STRUCTURAL AGNOSTIC SPMV: ADAPTING CSR-ADAPTIVE FOR IRREGULAR MATRICES | DECEMBER 17, 2015 | 2015 HIPC | BENGALURU, INDIA |

CSR-STREAM: STREAM MATRIX INTO THE LDS Local Data Share Fast accesses everywhere Block 1

CSR-STREAM: STREAM MATRIX INTO THE LDS Local Data Share Fast accesses everywhere Block 1 Block 2 Block 3 Block 4 12 | STRUCTURAL AGNOSTIC SPMV: ADAPTING CSR-ADAPTIVE FOR IRREGULAR MATRICES | DECEMBER 17, 2015 | 2015 HIPC | BENGALURU, INDIA |

How can we get good bandwidth for “medium-sized” rows? Improved reduction technique Efficient

How can we get good bandwidth for “medium-sized” rows? Improved reduction technique Efficient use of memory subsystem ✗ Block 1 Block 2 Block 3 CSR-Stream CSR-Vector Block 4 13 | STRUCTURAL AGNOSTIC SPMV: ADAPTING CSR-ADAPTIVE FOR IRREGULAR MATRICES | DECEMBER 17, 2015 | 2015 HIPC | BENGALURU, INDIA |

CSR-ADAPTIVE: SHORT AND MEDIUM ROWS Block 1 Short rows Medium-sized rows CSR-Stream CSR-Vector Block

CSR-ADAPTIVE: SHORT AND MEDIUM ROWS Block 1 Short rows Medium-sized rows CSR-Stream CSR-Vector Block 2 Block 3 Block 4 14 | STRUCTURAL AGNOSTIC SPMV: ADAPTING CSR-ADAPTIVE FOR IRREGULAR MATRICES | DECEMBER 17, 2015 | 2015 HIPC | BENGALURU, INDIA |

WHAT ABOUT VERY LONG ROWS? CSR-Vector: One workgroup can become a performance bottle-neck Multiple

WHAT ABOUT VERY LONG ROWS? CSR-Vector: One workgroup can become a performance bottle-neck Multiple workgroups compute the long row ✗ Block 1 Block 2 Block 3 CSR-Vector. L CSR-Vector Block 4 Block 5 15 | STRUCTURAL AGNOSTIC SPMV: ADAPTING CSR-ADAPTIVE FOR IRREGULAR MATRICES | DECEMBER 17, 2015 | 2015 HIPC | BENGALURU, INDIA |

CSR-VECTORL (LONG ROWS) Global Memory Local Data Share Block 1 Block 2 + atomic

CSR-VECTORL (LONG ROWS) Global Memory Local Data Share Block 1 Block 2 + atomic Block 3 Local Data Share Block 4 Block 5 16 | STRUCTURAL AGNOSTIC SPMV: ADAPTING CSR-ADAPTIVE FOR IRREGULAR MATRICES | DECEMBER 17, 2015 | 2015 HIPC | BENGALURU, INDIA |

OPTIMIZED CSR-STREAM: LOGARITHMIC REDUCTION 4 rows assigned to a workgroup with 8 threads –

OPTIMIZED CSR-STREAM: LOGARITHMIC REDUCTION 4 rows assigned to a workgroup with 8 threads – 2 threads / row 17 | STRUCTURAL AGNOSTIC SPMV: ADAPTING CSR-ADAPTIVE FOR IRREGULAR MATRICES | DECEMBER 17, 2015 | 2015 HIPC | BENGALURU, INDIA |

CSR-ADAPTIVE A complete Sp. MV solution CSR-Adaptive CSR-Stream CSR-Vector. L Short rows Medium-sized rows

CSR-ADAPTIVE A complete Sp. MV solution CSR-Adaptive CSR-Stream CSR-Vector. L Short rows Medium-sized rows Long rows Block 1 Block 2 Block 3 Block 4 Block 5 18 | STRUCTURAL AGNOSTIC SPMV: ADAPTING CSR-ADAPTIVE FOR IRREGULAR MATRICES | DECEMBER 17, 2015 | 2015 HIPC | BENGALURU, INDIA |

EXPERIMENTAL SETUP CPU: AMD A 10 -7850 K ‒ 32 GB DDR 3 -2133

EXPERIMENTAL SETUP CPU: AMD A 10 -7850 K ‒ 32 GB DDR 3 -2133 SDRAM GPU: AMD Fire. Pro. TM W 9100 ‒ 44 parallel 64 -wide compute units (CUs) – (2816 ALUs) ‒ 5. 2 single-precision TFLOPs / 2. 6 double-precision TFLOPS ‒ 320 GB/s 32 different input matrices used in previous works AMD APP SDK v 2. 9 AMD Fire. Pro driver v 14. 20 Beta on Ubuntu 14. 04. 2 19 | STRUCTURAL AGNOSTIC SPMV: ADAPTING CSR-ADAPTIVE FOR IRREGULAR MATRICES | DECEMBER 17, 2015 | 2015 HIPC | BENGALURU, INDIA |

Baseline Opt_CSR_Stream_Vector CSR-Adaptive (+Long. Rows) 20 | STRUCTURAL AGNOSTIC SPMV: ADAPTING CSR-ADAPTIVE FOR IRREGULAR

Baseline Opt_CSR_Stream_Vector CSR-Adaptive (+Long. Rows) 20 | STRUCTURAL AGNOSTIC SPMV: ADAPTING CSR-ADAPTIVE FOR IRREGULAR MATRICES | DECEMBER 17, 2015 | 2015 HIPC | BENGALURU, INDIA | transient sls boyd 2 Rucci 1 rajat 31 ldoor crankseg_2 bone 010 ins 2 Full. Chip dc 2 20 ASIC_680 k Si 41 Ge 41 H 72 mip 1 50 in-2004 Ga 41 As 41 H 72 eu-2005 circuit 5 M LP Webbase Circuit FEM/Accelerator Epidemiology Economics FEM/Ship QCD FEM/Harbor Wind Tunnel FEM/Cantilever FEM/Spheres Protein Dense 2 GFLOPS (Single Precision) CSR-ADAPTIVE OPTIMIZATIONS AMD FIREPRO W 9100 GPU 70 60 2 x 40 30 9 x 10 0

Lose out on 1/32 matrices 50 21 | STRUCTURAL AGNOSTIC SPMV: ADAPTING CSR-ADAPTIVE FOR

Lose out on 1/32 matrices 50 21 | STRUCTURAL AGNOSTIC SPMV: ADAPTING CSR-ADAPTIVE FOR IRREGULAR MATRICES | DECEMBER 17, 2015 | 2015 HIPC | BENGALURU, INDIA | Har Mean transient sls boyd 2 Rucci 1 rajat 31 ldoor crankseg_2 bone 010 ins 2 Full. Chip dc 2 30 ASIC_680 k Si 41 Ge 41 H 72 mip 1 in-2004 Ga 41 As 41 H 72 CSR 5 eu-2005 circuit 5 M 70 LP Webbase Circuit FEM/Accelerator Epidemiology Economics FEM/Ship QCD FEM/Harbor Wind Tunnel FEM/Cantilever FEM/Spheres Protein Dense 2 GFLOPS (Single Precision) CSR-ADAPTIVE VS CSR 5 AMD FIREPRO W 9100 GPU (SINGLE PRECISION) CSR-Adaptive 60 28% 40 2 x 20 10 0

OVERHEAD: CSR-ADAPTIVE VS CSR 5 22 | STRUCTURAL AGNOSTIC SPMV: ADAPTING CSR-ADAPTIVE FOR IRREGULAR

OVERHEAD: CSR-ADAPTIVE VS CSR 5 22 | STRUCTURAL AGNOSTIC SPMV: ADAPTING CSR-ADAPTIVE FOR IRREGULAR MATRICES | DECEMBER 17, 2015 | 2015 HIPC | BENGALURU, INDIA |

BANDWIDTH: CSR-ADAPTIVE 23 | STRUCTURAL AGNOSTIC SPMV: ADAPTING CSR-ADAPTIVE FOR IRREGULAR MATRICES | DECEMBER

BANDWIDTH: CSR-ADAPTIVE 23 | STRUCTURAL AGNOSTIC SPMV: ADAPTING CSR-ADAPTIVE FOR IRREGULAR MATRICES | DECEMBER 17, 2015 | 2015 HIPC | BENGALURU, INDIA |

CONCLUSION Optimized GPU-based Sp. MV algorithm that beats the competition by 2 x Operates

CONCLUSION Optimized GPU-based Sp. MV algorithm that beats the competition by 2 x Operates within 15% of GPU’s achievable b/w Minimal overhead for data structure generation Shipping now as part of AMD’s cl. SPARSE library https: //github. com/cl. Math. Libraries/cl. SPARSE 24 | STRUCTURAL AGNOSTIC SPMV: ADAPTING CSR-ADAPTIVE FOR IRREGULAR MATRICES | DECEMBER 17, 2015 | 2015 HIPC | BENGALURU, INDIA |

DISCLAIMER & ATTRIBUTION The information presented in this document is for informational purposes only

DISCLAIMER & ATTRIBUTION The information presented in this document is for informational purposes only and may contain technical inaccuracies, omissions and typographical errors. The information contained herein is subject to change and may be rendered inaccurate for many reasons, including but not limited to product and roadmap changes, component and motherboard version changes, new model and/or product releases, product differences between differing manufacturers, software changes, BIOS flashes, firmware upgrades, or the like. AMD assumes no obligation to update or otherwise correct or revise this information. However, AMD reserves the right to revise this information and to make changes from time to the content hereof without obligation of AMD to notify any person of such revisions or changes. AMD MAKES NO REPRESENTATIONS OR WARRANTIES WITH RESPECT TO THE CONTENTS HEREOF AND ASSUMES NO RESPONSIBILITY FOR ANY INACCURACIES, ERRORS OR OMISSIONS THAT MAY APPEAR IN THIS INFORMATION. AMD SPECIFICALLY DISCLAIMS ANY IMPLIED WARRANTIES OF MERCHANTABILITY OR FITNESS FOR ANY PARTICULAR PURPOSE. IN NO EVENT WILL AMD BE LIABLE TO ANY PERSON FOR ANY DIRECT, INDIRECT, SPECIAL OR OTHER CONSEQUENTIAL DAMAGES ARISING FROM THE USE OF ANY INFORMATION CONTAINED HEREIN, EVEN IF AMD IS EXPRESSLY ADVISED OF THE POSSIBILITY OF SUCH DAMAGES. ATTRIBUTION © 2015 Advanced Micro Devices, Inc. All rights reserved. AMD, the AMD Arrow logo, AMD Fire. Pro and combinations thereof are trademarks of Advanced Micro Devices, Inc. in the United States and/or other jurisdictions. Other names are for informational purposes only and may be trademarks of their respective owners. 25 | STRUCTURAL AGNOSTIC SPMV: ADAPTING CSR-ADAPTIVE FOR IRREGULAR MATRICES | DECEMBER 17, 2015 | 2015 HIPC | BENGALURU, INDIA |

Backup Slides

Backup Slides

FINDING BLOCKS FOR CSR-ADAPTIVE DONE ONCE ON THE CPU Assuming 8 LDS entries per

FINDING BLOCKS FOR CSR-ADAPTIVE DONE ONCE ON THE CPU Assuming 8 LDS entries per workgroup row blocks: 0 3 5 6 6 7 8 10 12 CSR Structure Unchanged row delimiters: 0 3 6 19 32 columns: values: Block 1 Block 2 Block 3 Block 4 Block 5 27 | STRUCTURAL AGNOSTIC SPMV: ADAPTING CSR-ADAPTIVE FOR IRREGULAR MATRICES | DECEMBER 17, 2015 | 2015 HIPC | BENGALURU, INDIA |

A DECENT SPMV FOR THE CPU IS SIMPLE #pragma omp parallel for 1. for

A DECENT SPMV FOR THE CPU IS SIMPLE #pragma omp parallel for 1. for (i = 0; i < n. Rows; i++) { 2. temp = 0; 3. for (j = row. Ds[i]; j < row. Ds[i+1]; j++) { 4. temp += vals[j] * vec[cols[j]]; 5. } 6. output[i] = temp; 7. } Only seven lines of code provides decent Sp. MV performance on the CPU 28 | STRUCTURAL AGNOSTIC SPMV: ADAPTING CSR-ADAPTIVE FOR IRREGULAR MATRICES | DECEMBER 17, 2015 | 2015 HIPC | BENGALURU, INDIA |

CSR-VECTOR USING THE LDS Local Data Share + + Fast uncoalesced accesses! Fast accesses

CSR-VECTOR USING THE LDS Local Data Share + + Fast uncoalesced accesses! Fast accesses everywhere 29 | STRUCTURAL AGNOSTIC SPMV: ADAPTING CSR-ADAPTIVE FOR IRREGULAR MATRICES | DECEMBER 17, 2015 | 2015 HIPC | BENGALURU, INDIA |

CSR-VECTOR USING THE LDS Local Data Share + Block 1 Block 2 + +

CSR-VECTOR USING THE LDS Local Data Share + Block 1 Block 2 + + + Block 3 Block 4 30 | STRUCTURAL AGNOSTIC SPMV: ADAPTING CSR-ADAPTIVE FOR IRREGULAR MATRICES | DECEMBER 17, 2015 | 2015 HIPC | BENGALURU, INDIA |

31 | STRUCTURAL AGNOSTIC SPMV: ADAPTING CSR-ADAPTIVE FOR IRREGULAR MATRICES | DECEMBER 17, 2015

31 | STRUCTURAL AGNOSTIC SPMV: ADAPTING CSR-ADAPTIVE FOR IRREGULAR MATRICES | DECEMBER 17, 2015 | 2015 HIPC | BENGALURU, INDIA | Har Mean transient sls boyd 2 Rucci 1 rajat 31 ldoor crankseg_2 bone 010 ins 2 Full. Chip dc 2 ASIC_680 k Si 41 Ge 41 H 72 mip 1 in-2004 Ga 41 As 41 H 72 CSR 5 eu-2005 circuit 5 M 50 LP Webbase Circuit FEM/Accelerator Epidemiology Economics FEM/Ship QCD FEM/Harbor Wind Tunnel FEM/Cantilever FEM/Spheres Protein Dense 2 GFLOPS (Double Precision) CSR-ADAPTIVE VS CSR 5 (BEST IN ACADEMIA) AMD FIREPRO W 9100 GPU (DOUBLE PRECISION) CSR-Adaptive 40 30 20 10 0

LOCAL DATA SHARE (LDS) MEMORY Local Data Share (64 KB) Scratchpad (software-addressed) on-chip cache

LOCAL DATA SHARE (LDS) MEMORY Local Data Share (64 KB) Scratchpad (software-addressed) on-chip cache Statically divided between all workgroups in a CU Highly ported to allow scatter/gather (uncoalesced) accesses 32 | STRUCTURAL AGNOSTIC SPMV: ADAPTING CSR-ADAPTIVE FOR IRREGULAR MATRICES | DECEMBER 17, 2015 | 2015 HIPC | BENGALURU, INDIA |

eligion agnostic atheism agnostic theism apathetic or pragmaticeligion agnostic atheism agnostic theism apathetic or pragmatic
Jeopardy Irregular Yo Regular Endings Irregular Preterites IrregularJeopardy Irregular Yo Regular Endings Irregular Preterites Irregular
Irregular Preterite Verbs These irregular verbs are irregularIrregular Preterite Verbs These irregular verbs are irregular
No Adaptive Adaptive Adaptive UI Adaptive UI UWPNo Adaptive Adaptive Adaptive UI Adaptive UI UWP
Adaptive Immunity Overview of Adaptive Immunity Adaptive immunityAdaptive Immunity Overview of Adaptive Immunity Adaptive immunity
Corporate Responsibility Corporate Social Responsibility CSR CSR policyCorporate Responsibility Corporate Social Responsibility CSR CSR policy
OVERVIEW OF CSR GLOBAL PROFESSIONAL OPPORTUNITIES IN CSROVERVIEW OF CSR GLOBAL PROFESSIONAL OPPORTUNITIES IN CSR
CSR i jenhjde Dansk Erhvervs seminar CSR iCSR i jenhjde Dansk Erhvervs seminar CSR i
CSR Mohiuddin Babar What is CSR Why shouldCSR Mohiuddin Babar What is CSR Why should
CSR in COLT Italia Piano operativo CSR pianoCSR in COLT Italia Piano operativo CSR piano
Corporate Social Responsibility CSR Spoleensk odpovdnost podnik CSRCorporate Social Responsibility CSR Spoleensk odpovdnost podnik CSR
Discussion on CSR CSR should be mandatory CourseDiscussion on CSR CSR should be mandatory Course
Spoleensk odpovdnost firmy CSR 2011 CSR v KovohutchSpoleensk odpovdnost firmy CSR 2011 CSR v Kovohutch
CSR Magyarorszgon Szigeti Ceclia A CSR viszonylag jCSR Magyarorszgon Szigeti Ceclia A CSR viszonylag j
CSR environment and safety CSR milj og arbejdsmiljCSR environment and safety CSR milj og arbejdsmilj
CSR Collaborative Strategic Reading What is CSR nCSR Collaborative Strategic Reading What is CSR n
Corporate Social Responsibility CSR Bidang Pertanian Konseptualisasi CSRCorporate Social Responsibility CSR Bidang Pertanian Konseptualisasi CSR
CSR 2 0 Management Perspectives of CSR PracticesCSR 2 0 Management Perspectives of CSR Practices
Progetto CSR Veneto Scuole CSR Training sulla ResponsabilitProgetto CSR Veneto Scuole CSR Training sulla Responsabilit
Corporate Social Responsibility CSR Bidang Pertanian Perkembangan CSRCorporate Social Responsibility CSR Bidang Pertanian Perkembangan CSR
CSR INDEX index spoleensk odpovdnosti 2812020 CSR naCSR INDEX index spoleensk odpovdnosti 2812020 CSR na
Corporate Social Responsibility CSR Bidang Pertanian CSR bidangCorporate Social Responsibility CSR Bidang Pertanian CSR bidang
CSR Standards Navigation Tool Extractive Sector CSR CounsellorsCSR Standards Navigation Tool Extractive Sector CSR Counsellors
Transformational CSR Simon Cooper CSR Consultancy Ltd SustainabilityTransformational CSR Simon Cooper CSR Consultancy Ltd Sustainability