compiler optimization for openmp accelerator offloading
play

Compiler Optimization For (OpenMP) Accelerator Offloading Johannes - PowerPoint PPT Presentation

Dec 01, 2022 •202 likes •1.91k views

EuroLLVM April 8, 2019 Brussels, Belgium Leadership Computing Facility Argonne National Laboratory Compiler Optimization For (OpenMP) Accelerator Offloading Johannes Doerfert and Hal Finkel https://www.alcf.anl.gov/ This research was

  1. Masters delegate work for concurrent execution. Masters execute bar concurrently, workers idle. 1 master and N-1 worker teams, worker teams M threads: 7/14 The Compiler Black Box — Behind the Curtain (of Clang) #pragma omp target teams { foo(); #pragma omp parallel work(); // <- Hotspot bar(); } Masters execute foo concurrently, workers idle.

  2. Masters execute bar concurrently, workers idle. 1 master and N-1 worker teams, worker teams M threads: 7/14 The Compiler Black Box — Behind the Curtain (of Clang) #pragma omp target teams { foo(); #pragma omp parallel work(); // <- Hotspot bar(); } Masters execute foo concurrently, workers idle. Masters delegate work for concurrent execution.

  3. 1 master and N-1 worker teams, worker teams M threads: 7/14 The Compiler Black Box — Behind the Curtain (of Clang) #pragma omp target teams { foo(); #pragma omp parallel work(); // <- Hotspot bar(); } Masters execute foo concurrently, workers idle. Masters delegate work for concurrent execution. Masters execute bar concurrently, workers idle.

  4. • a separate master team costs resources • synchronization has overhead • currently impossible to optimization 7/14 Problems: 1 master and N-1 worker teams, worker teams M threads: The Compiler Black Box — Behind the Curtain (of Clang) #pragma omp target teams { foo(); #pragma omp parallel work(); // <- Hotspot bar(); } Masters execute foo concurrently, workers idle. Masters delegate work for concurrent execution. Masters execute bar concurrently, workers idle.

  5. • synchronization has overhead • currently impossible to optimization 7/14 Problems: 1 master and N-1 worker teams, worker teams M threads: The Compiler Black Box — Behind the Curtain (of Clang) #pragma omp target teams { foo(); #pragma omp parallel work(); // <- Hotspot • a separate master team costs resources bar(); } Masters execute foo concurrently, workers idle. Masters delegate work for concurrent execution. Masters execute bar concurrently, workers idle.

  6. • currently impossible to optimization 7/14 Problems: 1 master and N-1 worker teams, worker teams M threads: The Compiler Black Box — Behind the Curtain (of Clang) #pragma omp target teams { foo(); #pragma omp parallel work(); // <- Hotspot • a separate master team costs resources bar(); • synchronization has overhead } Masters execute foo concurrently, workers idle. Masters delegate work for concurrent execution. Masters execute bar concurrently, workers idle.

  7. 1 master and N-1 worker teams, worker teams M threads: 7/14 Problems: The Compiler Black Box — Behind the Curtain (of Clang) #pragma omp target teams { foo(); #pragma omp parallel work(); // <- Hotspot • a separate master team costs resources bar(); • synchronization has overhead } • currently impossible to optimization Masters execute foo concurrently, workers idle. Masters delegate work for concurrent execution. Masters execute bar concurrently, workers idle.

  8. few + Logic +Logic + RT Calls Device RT Device Code Code Host Code IPO Module Cross Fat Binary Opt. 8/14 Host 1. Offmoad-Specific Optimizations on Device Code Opt. Device Gen. Code OpenMP Offload — Overview OpenMP Clang Assembly LLVM-IR LLVM

  9. few + Logic +Logic + RT Calls Device RT Device Code Host Code Module Cross IPO Fat Binary Opt. Host 8/14 1. Offmoad-Specific Optimizations on Device Code Opt. Device Gen. Code OpenMP Offload — Overview OpenMP Clang Assembly LLVM-IR LLVM Code

  10. few + Logic +Logic + RT Calls Device RT Device Code Host Code Module Cross IPO Fat Binary Opt. Host 8/14 1. Offmoad-Specific Optimizations on Device Code Opt. Device Gen. Code OpenMP Offload — Overview OpenMP Clang Assembly LLVM-IR LLVM Code

  11. +Logic few + Logic Device RT + RT Calls IPO Module Cross Fat Binary Opt. Host 1. Offmoad-Specific Optimizations on Device Code 8/14 Opt. Device Gen. Code OpenMP Offload — Overview OpenMP Clang Assembly LLVM-IR LLVM Device Code Code Host Code

  12. +Logic Device RT few + Logic IPO Module Cross Fat Binary Opt. Host 1. Offmoad-Specific Optimizations on Device Code 8/14 Opt. Device Gen. Code OpenMP Offload — Overview OpenMP Clang Assembly LLVM-IR LLVM Device Code+ RT Calls Code Host Code

  13. +Logic Device RT few IPO Module Cross Fat Binary Opt. Host 1. Offmoad-Specific Optimizations on Device Code 8/14 Opt. Device Code Gen. OpenMP Offload — Overview OpenMP Clang Assembly LLVM-IR LLVM Device Code+ RT Calls + Logic Code Host Code

  14. +Logic few Opt. IPO Module Cross Fat Binary Opt. Host 1. Offmoad-Specific Optimizations on Device Code 8/14 Device Gen. Code OpenMP Offload — Overview OpenMP Clang Assembly LLVM-IR LLVM Device RT Device Code+ RT Calls + Logic Code Host Code

  15. +Logic few Opt. IPO Module Cross Fat Binary Opt. Host 1. Offmoad-Specific Optimizations on Device Code 8/14 Device Gen. Code OpenMP Offload — Overview OpenMP Clang Assembly LLVM-IR LLVM Device RT Device Code+ RT Calls + Logic Code Host Code

  16. +Logic few Opt. IPO Module Cross Fat Binary Opt. Host 1. Offmoad-Specific Optimizations on Device Code 8/14 Device Gen. Code OpenMP Offload — Overview OpenMP Clang Assembly LLVM-IR LLVM Device RT Device Code+ RT Calls + Logic Code Host Code

  17. +Logic few Opt. IPO Module Cross Fat Binary Opt. Host 1. Offmoad-Specific Optimizations on Device Code 8/14 Device Gen. Code OpenMP Offload — Overview OpenMP Clang Assembly LLVM-IR LLVM Device RT Device Code+ RT Calls + Logic Code Host Code

  18. +Logic few Opt. IPO Module Cross Fat Binary Opt. Host 1. Offmoad-Specific Optimizations on Device Code 8/14 Device Gen. Code OpenMP Offload — Overview & Directions OpenMP Clang Assembly LLVM-IR LLVM Device RT Device Code+ RT Calls + Logic Code Host Code

  19. few 8/14 1. Offmoad-Specific Optimizations on Device Code IPO Code Gen. Module Cross Fat Binary Opt. Device Opt. Host OpenMP Offload — Overview & Directions OpenMP Clang Assembly LLVM-IR LLVM Device RT+Logic Device Code+ RT Calls + Logic Code Host Code

  20. 8/14 1. Offmoad-Specific Optimizations on Device Code IPO Code Gen. Module Cross Fat Binary Opt. Device Opt. Host OpenMP Offload — Overview & Directions OpenMP Clang Assembly LLVM-IR LLVM Device RT+Logic Device Code+few RT Calls + Logic Code Host Code

  21. Interested? Take a look and contact me :) 1. Offmoad-Specific Optimizations on Device Code Reviewers are needed! allow for more to come! Pending patches “fix” the motivating problem and 8/14 IPO Module Cross Fat Binary Opt. Host Code Gen. Opt. Device OpenMP Offload — Overview & Directions OpenMP Clang Assembly LLVM-IR LLVM Device RT+Logic Device Code+few RT Calls + Logic Code Host Code

  22. Interested? Take a look and contact me :) 1. Offmoad-Specific Optimizations on Device Code Reviewers are needed! allow for more to come! Pending patches “fix” the motivating problem and 8/14 IPO Module Cross Fat Binary Opt. Host Code Gen. Opt. Device OpenMP Offload — Overview & Directions OpenMP Clang Assembly LLVM-IR LLVM Device RT+Logic Device Code+few RT Calls + Logic Code Host Code

  23. Interested? Take a look and contact me :) 1. Offmoad-Specific Optimizations on Device Code Reviewers are needed! allow for more to come! Pending patches “fix” the motivating problem and 8/14 IPO Module Cross Fat Binary Opt. Host Code Gen. Opt. Device OpenMP Offload — Overview & Directions OpenMP Clang Assembly LLVM-IR LLVM Device RT+Logic Device Code+few RT Calls + Logic Code Host Code

  24. 8/14 1. Offmoad-Specific Optimizations on Device Code IPO Code Gen. Module Cross Fat Binary Opt. Device Opt. Host OpenMP Offload — Overview & Directions OpenMP Clang Assembly LLVM-IR LLVM Device RT+Logic Device Code+few RT Calls + Logic Code Host Code

  25. 8/14 2. Optimize Device and Host Code Together IPO Code Gen. Module Cross Fat Binary Opt. Device Opt. Host OpenMP Offload — Overview & Directions OpenMP Clang Assembly LLVM-IR LLVM Device RT+Logic Device Code+few RT Calls + Logic Code Host Code

  26. 8/14 2. Optimize Device and Host Code Together IPO Code Gen. Module Cross Fat Binary Opt. Device Opt. Host OpenMP Offload — Overview & Directions OpenMP Clang Assembly LLVM-IR LLVM Device RT+Logic Device Code+few RT Calls + Logic Code Host Code

  27. 8/14 Device mization Opti- Device AND Host Opt. Opt. Host Fat Binary Cross Module Gen. Code IPO 2. Optimize Device and Host Code Together OpenMP Offload — Overview & Directions OpenMP Clang Assembly LLVM-IR LLVM Device RT+Logic Device Code+few RT Calls + Logic Code Host Code

  28. • A straight-forward #pragma omp target front-end: • Interface exposes information and implementation choices: • Device RT interface & implementation are separated: ⋄ simplified implementation CGOpenMPRuntimeNVPTX.cpp ~5.0k loc ⋄ improved reusability (F18, ...) CGOpenMPRuntimeTRegion.cpp ~0.5k loc ⋄ “smartness” is moved in the compiler middle-end ⋄ simplifies analyses and transformations in LLVM ⋄ simplifies generated LLVM-IR ⋄ most LLVM & Clang parts become target agnostic 9/14 Pending Patches — Target Region Interface

  29. • Interface exposes information and implementation choices: • Device RT interface & implementation are separated: ⋄ simplified implementation CGOpenMPRuntimeNVPTX.cpp ~5.0k loc ⋄ improved reusability (F18, ...) CGOpenMPRuntimeTRegion.cpp ~0.5k loc ⋄ “smartness” is moved in the compiler middle-end ⋄ simplifies analyses and transformations in LLVM ⋄ simplifies generated LLVM-IR ⋄ most LLVM & Clang parts become target agnostic 9/14 Pending Patches — Target Region Interface • A straight-forward #pragma omp target front-end:

  30. • Interface exposes information and implementation choices: • Device RT interface & implementation are separated: ⋄ simplified implementation CGOpenMPRuntimeNVPTX.cpp ~5.0k loc ⋄ improved reusability (F18, ...) CGOpenMPRuntimeTRegion.cpp ~0.5k loc ⋄ “smartness” is moved in the compiler middle-end ⋄ simplifies analyses and transformations in LLVM ⋄ simplifies generated LLVM-IR ⋄ most LLVM & Clang parts become target agnostic 9/14 Pending Patches — Target Region Interface • A straight-forward #pragma omp target front-end:

  31. • Device RT interface & implementation are separated: ⋄ simplified implementation CGOpenMPRuntimeNVPTX.cpp ~5.0k loc ⋄ improved reusability (F18, ...) CGOpenMPRuntimeTRegion.cpp ~0.5k loc ⋄ “smartness” is moved in the compiler middle-end ⋄ simplifies analyses and transformations in LLVM ⋄ simplifies generated LLVM-IR ⋄ most LLVM & Clang parts become target agnostic 9/14 Pending Patches — Target Region Interface • A straight-forward #pragma omp target front-end: • Interface exposes information and implementation choices:

  32. ⋄ simplified implementation CGOpenMPRuntimeNVPTX.cpp ~5.0k loc ⋄ improved reusability (F18, ...) CGOpenMPRuntimeTRegion.cpp ~0.5k loc ⋄ “smartness” is moved in the compiler middle-end ⋄ simplifies analyses and transformations in LLVM ⋄ simplifies generated LLVM-IR ⋄ most LLVM & Clang parts become target agnostic 9/14 Pending Patches — Target Region Interface • A straight-forward #pragma omp target front-end: • Interface exposes information and implementation choices: • Device RT interface & implementation are separated:

  33. • if legal, switch all boolean UseSPMDMode flags to true • currently, no (unknown) global side-efgects allowed outside parallel regions. 10/14 1. Offload-Specific Optimizations — “SPMD-zation” • use inter-procedural reasoning to place minimal guards/synchronization

  34. • currently, no (unknown) global side-efgects allowed outside parallel regions. 10/14 1. Offload-Specific Optimizations — “SPMD-zation” • use inter-procedural reasoning to place minimal guards/synchronization • if legal, switch all boolean UseSPMDMode flags to true

  35. 10/14 1. Offload-Specific Optimizations — “SPMD-zation” • use inter-procedural reasoning to place minimal guards/synchronization • if legal, switch all boolean UseSPMDMode flags to true • currently, no (unknown) global side-efgects allowed outside parallel regions.

  36. • reachability & post-dominance restrict the set of potential next parallel regions • reuse already communicated/shared values if possible • currently, a simple state machine is generated with explicit conditionals for all to work on known parallel regions in the module 11/14 1. Offload-Specific Optimizations — Custom State Machines • use optimized state-machines when unavoidable

  37. • reuse already communicated/shared values if possible • currently, a simple state machine is generated with explicit conditionals for all to work on known parallel regions in the module 11/14 1. Offload-Specific Optimizations — Custom State Machines • use optimized state-machines when unavoidable • reachability & post-dominance restrict the set of potential next parallel regions

  38. • currently, a simple state machine is generated with explicit conditionals for all to work on known parallel regions in the module 11/14 1. Offload-Specific Optimizations — Custom State Machines • use optimized state-machines when unavoidable • reachability & post-dominance restrict the set of potential next parallel regions • reuse already communicated/shared values if possible

  39. known parallel regions in the module to work on 11/14 1. Offload-Specific Optimizations — Custom State Machines • use optimized state-machines when unavoidable • reachability & post-dominance restrict the set of potential next parallel regions • reuse already communicated/shared values if possible • currently, a simple state machine is generated with explicit conditionals for all

  40. TransitiveCallSite AbstractCallSite Passes (IPOs) 12/14 2. Optimize Device and Host Together — Abstract Call Sites CallInst InvokeInst CallSite Passes (IPOs)

  41. Passes (IPOs) 12/14 2. Optimize Device and Host Together — Abstract Call Sites CallInst InvokeInst CallSite TransitiveCallSite AbstractCallSite Passes (IPOs)

  42. Passes (IPOs) 12/14 Functional changes required for Inter-procedural Constant Propagation: 2. Optimize Device and Host Together — Abstract Call Sites CallInst InvokeInst CallSite TransitiveCallSite AbstractCallSite Passes (IPOs)

  43. 13/14 Abstract Call Sites — Performance Results

  44. 13/14 Abstract Call Sites — Performance Results

  45. 13/14 Abstract Call Sites — Performance Results

  46. 13/14 Abstract Call Sites — Performance Results

  47. 14/14 Conclusion

  48. 14/14 Conclusion

  49. 14/14 Conclusion

  50. 14/14 Conclusion

  51. 14/14 Conclusion

  53. I: Attribute Propagation — Bidirectional Information Transfer: read/write-only , restrict / noalias , … II: Variable Privatization — Limit Variable Lifetimes: shared(var) ⟶ firstprivate(var) ⟶ private(var) III: Parallel Region Expansion — Maximize Parallel Contexts: IV: Barrier Elimination — Eliminate Redundant Barriers V: Communication Optimization — Move Computations Around: seq. compute&result comm. vs. operand comm. &par. compute ⟹ reduce start/stop overheads and expose barriers OpenMP - Aware Optimizations (see IWOMP’18)

  54. II: Variable Privatization — Limit Variable Lifetimes: shared(var) ⟶ firstprivate(var) ⟶ private(var) III: Parallel Region Expansion — Maximize Parallel Contexts: IV: Barrier Elimination — Eliminate Redundant Barriers V: Communication Optimization — Move Computations Around: seq. compute&result comm. vs. operand comm. &par. compute ⟹ reduce start/stop overheads and expose barriers OpenMP - Aware Optimizations (see IWOMP’18) I: Attribute Propagation — Bidirectional Information Transfer: read/write-only , restrict / noalias , …

  55. III: Parallel Region Expansion — Maximize Parallel Contexts: IV: Barrier Elimination — Eliminate Redundant Barriers V: Communication Optimization — Move Computations Around: seq. compute&result comm. vs. operand comm. &par. compute ⟹ reduce start/stop overheads and expose barriers OpenMP - Aware Optimizations (see IWOMP’18) I: Attribute Propagation — Bidirectional Information Transfer: read/write-only , restrict / noalias , … II: Variable Privatization — Limit Variable Lifetimes: shared(var) ⟶ firstprivate(var) ⟶ private(var)

  56. IV: Barrier Elimination — Eliminate Redundant Barriers V: Communication Optimization — Move Computations Around: seq. compute&result comm. vs. operand comm. &par. compute ⟹ reduce start/stop overheads and expose barriers OpenMP - Aware Optimizations (see IWOMP’18) I: Attribute Propagation — Bidirectional Information Transfer: read/write-only , restrict / noalias , … II: Variable Privatization — Limit Variable Lifetimes: shared(var) ⟶ firstprivate(var) ⟶ private(var) III: Parallel Region Expansion — Maximize Parallel Contexts:

  57. V: Communication Optimization — Move Computations Around: seq. compute&result comm. vs. operand comm. &par. compute ⟹ reduce start/stop overheads and expose barriers OpenMP - Aware Optimizations (see IWOMP’18) I: Attribute Propagation — Bidirectional Information Transfer: read/write-only , restrict / noalias , … II: Variable Privatization — Limit Variable Lifetimes: shared(var) ⟶ firstprivate(var) ⟶ private(var) III: Parallel Region Expansion — Maximize Parallel Contexts: IV: Barrier Elimination — Eliminate Redundant Barriers

  58. ⟹ reduce start/stop overheads and expose barriers OpenMP - Aware Optimizations (see IWOMP’18) I: Attribute Propagation — Bidirectional Information Transfer: read/write-only , restrict / noalias , … II: Variable Privatization — Limit Variable Lifetimes: shared(var) ⟶ firstprivate(var) ⟶ private(var) III: Parallel Region Expansion — Maximize Parallel Contexts: IV: Barrier Elimination — Eliminate Redundant Barriers V: Communication Optimization — Move Computations Around: seq. compute&result comm. vs. operand comm. &par. compute

  59. ⟹ reduce start/stop overheads and expose barriers OpenMP - Aware Optimizations (see IWOMP’18) I: Attribute Propagation — Bidirectional Information Transfer: read/write-only , restrict / noalias , … II: Variable Privatization — Limit Variable Lifetimes: shared(var) ⟶ firstprivate(var) ⟶ private(var) III: Parallel Region Expansion — Maximize Parallel Contexts: IV: Barrier Elimination — Eliminate Redundant Barriers V: Communication Optimization — Move Computations Around: seq. compute&result comm. vs. operand comm. &par. compute

  60. ⟹ reduce start/stop overheads and expose barriers OpenMP - Aware Optimizations (see IWOMP’18) I: Attribute Propagation — In LLVM: Attribute Deduction (IPO!) read/write-only , restrict / noalias , … II: Variable Privatization — In LLVM: Argument Promotion (IPO!) shared(var) ⟶ firstprivate(var) ⟶ private(var) III: Parallel Region Expansion — Maximize Parallel Contexts: IV: Barrier Elimination — Eliminate Redundant Barriers V: Communication Optimization — Move Computations Around: seq. compute&result comm. vs. operand comm. &par. compute

  61. OpenMP Input: Early Outlining #pragma omp parallel for for ( int i = 0; i < N; i++) Out[i] = In[i] + In[i+N];

  62. OpenMP Input: Early Outlining #pragma omp parallel for for ( int i = 0; i < N; i++) Out[i] = In[i] + In[i+N]; // Parallel region replaced by a runtime call. omp_rt_parallel_for(0, N, &body_fn, &N, &In, &Out);

  63. OpenMP Input: Early Outlining #pragma omp parallel for for ( int i = 0; i < N; i++) Out[i] = In[i] + In[i+N]; // Parallel region replaced by a runtime call. omp_rt_parallel_for(0, N, &body_fn, &N, &In, &Out); // Parallel region outlined in the front-end (clang)! static void body_fn( int tid, int *N, float ** In, float ** Out) { int lb = omp_get_lb(tid), ub = omp_get_ub(tid); for ( int i = lb; i < ub; i++) (*Out)[i] = (*In)[i] + (*In)[i + (*N)] }

  64. OpenMP Input: Early Outlining #pragma omp parallel for for ( int i = 0; i < N; i++) Out[i] = In[i] + In[i+N]; // Parallel region replaced by a runtime call. omp_rt_parallel_for(0, N, &body_fn, &N, &In, &Out); // Parallel region outlined in the front-end (clang)! static void body_fn( int tid, int * N, float ** In, float ** Out) { int lb = omp_get_lb(tid), ub = omp_get_ub(tid); for ( int i = lb; i < ub; i++) (*Out)[i] = (*In)[i] + (*In)[i + (*N)] }

  65. OpenMP Input: An Abstract Parallel IR #pragma omp parallel for for ( int i = 0; i < N; i++) Out[i] = In[i] + In[i+N]; // Parallel region replaced by an annotated loop for /* parallel */ ( int i = 0; i < N; i++) body_fn(i, &N, &In, &Out); // Parallel region outlined in the front-end (clang)! static void body_fn( int i , int * N, float ** In, float ** Out) { (*Out)[i] = (*In)[i] + (*In)[i + (*N)] }

  66. OpenMP Input: Early Outlined + Transitive Calls #pragma omp parallel for for ( int i = 0; i < N; i++) Out[i] = In[i] + In[i+N]; // Parallel region replaced by a runtime call. omp_rt_parallel_for(0, N, &body_fn, &N, &In, &Out); // Model transitive call: body_fn(?, &N, &In, &Out); // Parallel region outlined in the front-end (clang)! static void body_fn( int tid, int * N, float ** In, float ** Out) { int lb = omp_get_lb(tid), ub = omp_get_ub(tid); for ( int i = lb; i < ub; i++) (*Out)[i] = (*In)[i] + (*In)[i + (*N)] }

  67. OpenMP Input: − integration cost per IPO Early Outlined + Transitive Calls #pragma omp parallel for for ( int i = 0; i < N; i++) Out[i] = In[i] + In[i+N]; // Parallel region replaced by a runtime call. omp_rt_parallel_for(0, N, &body_fn, &N, &In, &Out); // Model transitive call: body_fn(?, &N, &In, &Out); // Parallel region outlined in the front-end (clang)! static void body_fn( int tid, int * N, float ** In, float ** Out) { int lb = omp_get_lb(tid), ub = omp_get_ub(tid); + valid and executable IR for ( int i = lb; i < ub; i++) (*Out)[i] = (*In)[i] + (*In)[i + (*N)] } + no unintended interactions

  68. TransitiveCallSite AbstractCallSite Passes (IPOs) IPO in LLVM CallInst InvokeInst CallSite Passes (IPOs)

Recommend

OpenMP Device Offloading to FPGA Accelerators Lukas Sommer, Jens Korinth, Andreas Koch

OpenMP Device Offloading to FPGA Accelerators Lukas Sommer, Jens Korinth, Andreas Koch

OpenMP Device Offloading to FPGA Accelerators Lukas Sommer, Jens Korinth, Andreas Koch Motivation Increasing use of heterogeneous systems to overcome CPU power limitations 2017-07-12 | OpenMP FPGA Device Offloading | L. Sommer, J. Korinth,

655 views • 43 slides
Integrating GPU Support for OpenMP Offloading Directives into Clang Carlo Bertolli , Samuel F.

Integrating GPU Support for OpenMP Offloading Directives into Clang Carlo Bertolli , Samuel F.

Integrating GPU Support for OpenMP Offloading Directives into Clang Carlo Bertolli , Samuel F. Antao, Gheorghe-Teodor Bercea, Arpith C. Jacob, Alexandre E. Eichenberger, Tong Chen, Zehra Sura, Hyojin Sung, Georgios Rokos, David Appelhans, Kevin

597 views • 32 slides
OpenMP Offloading Verification and Validation: Workflow and Road to 5.0 Thomas Huber &amp; Joshua

OpenMP Offloading Verification and Validation: Workflow and Road to 5.0 Thomas Huber &amp; Joshua

OpenMP Offloading Verification and Validation: Workflow and Road to 5.0 Thomas Huber &amp; Joshua Davis (UD) Jose Monsalve Diaz (UD) Swaroop Pophale (ORNL) Sunita Chandrasekaran (UD) 1 Kyle Friedline (UD) Oscar Hernandez (ORNL) David E.

420 views • 22 slides
OpenMP 1 What is OpenMP? An Application Program Interface (API) used to explicitly direct

OpenMP 1 What is OpenMP? An Application Program Interface (API) used to explicitly direct

OpenMP 1 What is OpenMP? An Application Program Interface (API) used to explicitly direct multi-threaded, shared memory parallelism. Comprised of three primary API components: Compiler Directives Runtime Library Routines

899 views • 67 slides
QMT QCD Multi Threading First steps Step 1: General Evaluation OpenMP vs.

QMT QCD Multi Threading First steps Step 1: General Evaluation OpenMP vs.

QMT QCD Multi Threading First steps Step 1: General Evaluation OpenMP vs. Explicit Thread library (Chen) Explicit thread library can do better than OpenMP OpenMP performance is compiler dependent Intel compiler does

388 views • 5 slides
Optimization on one core OpenMP, MPI and hybrid programming An introduction to the de-facto

Optimization on one core OpenMP, MPI and hybrid programming An introduction to the de-facto

Optimization on one core OpenMP, MPI and hybrid programming An introduction to the de-facto industrial standards Vincent Keller Nicolas Richart July 10, 2019 Table of Contents Parallelization with OpenMP Lecture based on specifications ver

566 views • 42 slides
Introduction to OpenMP Lecture 2: OpenMP fundamentals Overview Basic Concepts in OpenMP

Introduction to OpenMP Lecture 2: OpenMP fundamentals Overview Basic Concepts in OpenMP

Introduction to OpenMP Lecture 2: OpenMP fundamentals Overview Basic Concepts in OpenMP History of OpenMP Compiling and running OpenMP programs 2 What is OpenMP? OpenMP is an API designed for programming shared memory parallel

547 views • 17 slides
Recommended Reading A Brief Introduction to OpenMP OpenMP FAQ http://openmp.org/openmp-faq.html

Recommended Reading A Brief Introduction to OpenMP OpenMP FAQ http://openmp.org/openmp-faq.html

Recommended Reading A Brief Introduction to OpenMP OpenMP FAQ http://openmp.org/openmp-faq.html Will Knottenbelt OpenMP on Wikipedia http://en.wikipedia.org/wiki/OpenMP Imperial College London wjk@doc.ic.ac.uk OpenMP Tutorial

253 views • 4 slides
OpenMP 4 - Whats New? SciNet Developer Seminar Ramses van Zon September 25, 2013 Intro to

OpenMP 4 - Whats New? SciNet Developer Seminar Ramses van Zon September 25, 2013 Intro to

OpenMP 4 - Whats New? SciNet Developer Seminar Ramses van Zon September 25, 2013 Intro to OpenMP I For shared memory systems. I Add parallelism to functioning serial code. I For C, C++ and Fortran I http://openmp.org I Compiler/run-time does

686 views • 49 slides
Threaded Programming Lecture 2: OpenMP fundamentals Overview Basic Concepts in OpenMP

Threaded Programming Lecture 2: OpenMP fundamentals Overview Basic Concepts in OpenMP

Threaded Programming Lecture 2: OpenMP fundamentals Overview Basic Concepts in OpenMP History of OpenMP Compiling and running OpenMP programs 2 What is OpenMP? OpenMP is an API designed for programming shared memory parallel

396 views • 14 slides
Advanced OpenMP Lecture 11: OpenMP 4.0 OpenMP 4.0 Version 4.0 was released in July 2013

Advanced OpenMP Lecture 11: OpenMP 4.0 OpenMP 4.0 Version 4.0 was released in July 2013

Advanced OpenMP Lecture 11: OpenMP 4.0 OpenMP 4.0 Version 4.0 was released in July 2013 Starting to make an appearance in production compilers Whats new in 4.0 User defined reductions Construct cancellation Portable SIMD

713 views • 22 slides
OpenMP Paolo Burgio paolo.burgio@unimore.it A history of OpenMP 1997 OpenMP for

OpenMP Paolo Burgio paolo.burgio@unimore.it A history of OpenMP 1997 OpenMP for

An introduction to OpenMP Paolo Burgio paolo.burgio@unimore.it A history of OpenMP 1997 OpenMP for Fortran 1.0 1998 OpenMP for C/C++ 1.0 2000 Regular, loop-based parallelism OpenMP for Fortran 2.0 2002

518 views • 13 slides
OpenMP 4.0 and Beyond! Aidan Chalk, Hartree Centre, STFC What is OpenMP? OpenMP is an API

OpenMP 4.0 and Beyond! Aidan Chalk, Hartree Centre, STFC What is OpenMP? OpenMP is an API

OpenMP 4.0 and Beyond! Aidan Chalk, Hartree Centre, STFC What is OpenMP? OpenMP is an API &amp; standard for shared memory parallel computing. Works with C, C++ and Fortran. It was first released in 1997, and version 4.5 was released

1.14k views • 51 slides
Status of GPU offloading on Wayland Axel Davy FOSDEM 2014 Status of GPU offloading on Wayland

Status of GPU offloading on Wayland Axel Davy FOSDEM 2014 Status of GPU offloading on Wayland

Status of GPU offloading on Wayland Status of GPU offloading on Wayland Axel Davy FOSDEM 2014 Status of GPU offloading on Wayland How to do GPU offloading 1 GPU offloading with X DRI2 2 GPU offloading with Wayland 3 and XWayland? 4

427 views • 29 slides
Automatic Compiler Based FPGA Accelerator for CNN Training Shreyas Venkataramanaiah 1 , Yufei Ma

Automatic Compiler Based FPGA Accelerator for CNN Training Shreyas Venkataramanaiah 1 , Yufei Ma

Automatic Compiler Based FPGA Accelerator for CNN Training Shreyas Venkataramanaiah 1 , Yufei Ma 1 , Shihui Yin 1 , Eriko Nurvithadhi 2 , Aravind Dasu 3 , Yu Cao 1 , Jae-sun Seo 1 1 School of ECEE, Arizona State University, Tempe, AZ, USA 2 Intel

539 views • 21 slides
Parallel Programming with OpenMP CS240A, T. Yang 1 A Programmer s View of OpenMP What

Parallel Programming with OpenMP CS240A, T. Yang 1 A Programmer s View of OpenMP What

Parallel Programming with OpenMP CS240A, T. Yang 1 A Programmer s View of OpenMP What is OpenMP? Open specification for Multi-Processing Standard API for defining multi-threaded shared-memory programs openmp.org

612 views • 35 slides
Aestimo : A Feedback-Directed Optimization Evaluation Tool A Compiler Perspective on Input

Aestimo : A Feedback-Directed Optimization Evaluation Tool A Compiler Perspective on Input

Aestimo : A Feedback-Directed Optimization Evaluation Tool A Compiler Perspective on Input Characterization Paul Berube Compiler Design and Optimization Laboratory University of Alberta March 21, 2006 Paul Berube, ISPASS 2006 1 Outline

910 views • 37 slides
Seamless Offloading of Web App Computations From Mobile Device to Edge Clouds via HTML5 Web

Seamless Offloading of Web App Computations From Mobile Device to Edge Clouds via HTML5 Web

Seamless Offloading of Web App Computations From Mobile Device to Edge Clouds via HTML5 Web Worker Migration Hyuk Jin Jeong Seoul National University SoCC 2019 Virtual Machine &amp; Optimization Laboratory Department of Electrical and

824 views • 29 slides
Mitigating the Compiler Optimization Phase-Ordering Problem using Machine Learning Sameer

Mitigating the Compiler Optimization Phase-Ordering Problem using Machine Learning Sameer

Mitigating the Compiler Optimization Phase-Ordering Problem using Machine Learning Sameer Kulkarni &amp; John Cavazos Intro - Why we need code optimization? - Prog. Language design flaw - e.g. goto statement in C/Cpp - People are evil and

854 views • 19 slides
SHARED MEMORY PROGRAMMING WITH OPENMP Lecture 9: OpenMP Performance 2 A common scenario.....

SHARED MEMORY PROGRAMMING WITH OPENMP Lecture 9: OpenMP Performance 2 A common scenario.....

SHARED MEMORY PROGRAMMING WITH OPENMP Lecture 9: OpenMP Performance 2 A common scenario..... So I wrote my OpenMP program, and I checked it gave the right answers, so I ran some timing tests, and the speedup was, well, a bit disappointing

596 views • 33 slides
A GENERAL-PURPOSE CRN-TO-DSD COMPILER WITH FORMAL VERIFICATION, OPTIMIZATION, AND SIMULATION

A GENERAL-PURPOSE CRN-TO-DSD COMPILER WITH FORMAL VERIFICATION, OPTIMIZATION, AND SIMULATION

A GENERAL-PURPOSE CRN-TO-DSD COMPILER WITH FORMAL VERIFICATION, OPTIMIZATION, AND SIMULATION CAPABILITIES Stefan Badelt , Seung Woo Shin, Robert F. Johnson, Qing Dong, Chris Thachuk, and Erik Winfree DNA and Natural Algorithms (DNA) Group,

320 views • 28 slides
Compiler-assisted Performance Analysis Adam Nemet Apple anemet@apple.com Hotspot User

Compiler-assisted Performance Analysis Adam Nemet Apple anemet@apple.com Hotspot User

Compiler-assisted Performance Analysis Adam Nemet Apple anemet@apple.com Hotspot User Bottleneck Compiler Optimization X, Y 2 Hotspot Hotspot Legality User Compiler Bottleneck Cost Model Compiler Optimization Some Optimizations?

1.19k views • 90 slides
Speeding Up Reactive Transport Code Using OpenMP By Jared McLaughlin OpenMP A standard for

Speeding Up Reactive Transport Code Using OpenMP By Jared McLaughlin OpenMP A standard for

Speeding Up Reactive Transport Code Using OpenMP By Jared McLaughlin OpenMP A standard for parallelizing Fortran and C/C++ on shared memory systems Minimal changes to sequential code required Incremental parallelization OpenMP

338 views • 13 slides
Automatic Compiler-Based Optimization of Graph Analytics for the GPU Sreepathi Pai The

Automatic Compiler-Based Optimization of Graph Analytics for the GPU Sreepathi Pai The

Automatic Compiler-Based Optimization of Graph Analytics for the GPU Sreepathi Pai The University of Texas at Austin May 8, 2017 NVIDIA GTC Parallel Graph Processing is not easy 299ms HD-BFS 84ms USA Road Network LiveJournal Social

601 views • 48 slides

More recommend