<html><head><meta http-equiv="Content-Type" content="text/html charset=utf-8"></head><body style="word-wrap: break-word; -webkit-nbsp-mode: space; -webkit-line-break: after-white-space;" class=""><br class=""><div><blockquote type="cite" class=""><div class="">On Jan 11, 2017, at 3:51 PM, Reid Kleckner via llvm-dev <<a href="mailto:llvm-dev@lists.llvm.org" class="">llvm-dev@lists.llvm.org</a>> wrote:</div><br class="Apple-interchange-newline"><div class=""><div dir="ltr" class="">+1, tokens are the current True Way to create single-entry multi-exit regions. Your example for an annotated loop would look like:<div class=""><br class=""></div><div class=""><span style="font-size:12.8px" class="">%region = call token @llvm.openmp.regionstart(metadata ...) ; whatever parameters you need here</span><br style="font-size:12.8px" class=""><span style="font-size:12.8px" class=""> loop</span><br style="font-size:12.8px" class=""><span style="font-size:12.8px" class="">call void @llvm.openmp.regionend(token %region)</span><br class=""></div><div class=""><span style="font-size:12.8px" class=""><br class=""></span></div><div class=""><span style="font-size:12.8px" class="">If you use tokens, I would recommend proposal (c), where you introduce new intrinsics for every new kind of region, instead of adding one overly generic set of region intrinsics.</span></div></div></div></blockquote><div><br class=""></div><div>Can you elaborate why? I’m curious.</div><div><br class=""></div><div>Thanks,</div><div><br class=""></div><div>— </div><div>Mehdi</div><div><br class=""></div><br class=""><blockquote type="cite" class=""><div class=""><div dir="ltr" class=""><div class=""><span style="font-size:12.8px" class=""><br class=""></span></div><div class=""><span style="font-size:12.8px" class="">We already have a way to form regions with real barriers, and it's tokens.</span></div></div><div class="gmail_extra"><br class=""><div class="gmail_quote">On Wed, Jan 11, 2017 at 2:17 PM, David Majnemer via llvm-dev <span dir="ltr" class=""><<a href="mailto:llvm-dev@lists.llvm.org" target="_blank" class="">llvm-dev@lists.llvm.org</a>></span> wrote:<br class=""><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex"><div dir="ltr" class="">FWIW, we needed to maintain single entry-multiple exit regions for WinEH and we accomplished it via a different mechanism.<div class=""><br class=""></div><div class="">We had an instruction which produces a value of type Token (<a href="http://llvm.org/docs/LangRef.html#token-type" target="_blank" class="">http://llvm.org/docs/LangRef.<wbr class="">html#token-type</a>) which let us establish the region and another instruction to exit the region by consuming it. The dominance rules allowed us to avoid situations where the compiler might trash the regions in weird ways and made sure that regions would be left unharmed.</div><div class=""><br class=""></div><div class="">AFAIK, a similar approach using Token could work here. I think it would reduce the amount of stuff you'd need LLVM to maintain.</div><div class=""><br class=""></div></div><div class="HOEnZb"><div class="h5"><div class="gmail_extra"><br class=""><div class="gmail_quote">On Wed, Jan 11, 2017 at 2:02 PM, Hal Finkel via llvm-dev <span dir="ltr" class=""><<a href="mailto:llvm-dev@lists.llvm.org" target="_blank" class="">llvm-dev@lists.llvm.org</a>></span> wrote:<br class=""><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">A Proposal for adding an experimental IR-level region-annotation infrastructure<br class="">
==============================<wbr class="">==============================<wbr class="">================= <br class="">
Hal Finkel (ANL) and Xinmin Tian (Intel)<br class="">
<br class="">
This is a proposal for adding an experimental infrastructure to support<br class="">
annotating regions in LLVM IR, making use of intrinsics and metadata, and<br class="">
a generic analysis to allow transformations to easily make use of these<br class="">
annotated regions. This infrastructure is flexible enough to support<br class="">
representation of directives for parallelization, vectorization, and<br class="">
offloading of both loops and more-general code regions. Under this scheme,<br class="">
the conceptual distance between source-level directives and the region<br class="">
annotations need not be significant, making the incremental cost of<br class="">
supporting new directives and modifiers often small. It is not, however,<br class="">
specific to those use cases.<br class="">
<br class="">
Problem Statement<br class="">
=================<br class="">
There are a series of discussions on LLVM IR extensions for representing region<br class="">
and loop annotations for parallelism, and other user-guided transformations,<br class="">
among both industrial and academic members of the LLVM community. Increasing<br class="">
the quality of our OpenMP implementation is an important motivating use case,<br class="">
but certainly not the only one. For OpenMP in particular, we've discussed<br class="">
having an IR representation for years. Presently, all OpenMP pragmas are<br class="">
transformed directly into runtime-library calls in Clang, and outlining (i.e.<br class="">
extracting parallel regions into their own functions to be invoked by the<br class="">
runtime library) is done in Clang as well. Our implementation does not further<br class="">
optimize OpenMP constructs, and a lot of thought has been put into how we might<br class="">
improve this. For some optimizations, such as redundant barrier removal, we<br class="">
could use a TargetLibraryInfo-like mechanism to recognize frontend-generated<br class="">
runtime calls and proceed from there. Dealing with cases where we lose<br class="">
pointer-aliasing information, information on loop bounds, etc. we could improve<br class="">
by improving our inter-procedural-analysis capabilities. We should do that<br class="">
regardless. However, there are important cases where the underlying scheme we<br class="">
want to use to lower the various parallelism constructs, especially when<br class="">
targeting accelerators, changes depending on what is in the parallel region.<br class="">
In important cases where we can see everything (i.e. there aren't arbitrary<br class="">
external calls), code generation should proceed in a way that is very different<br class="">
from the general case. To have a sensible implementation, this must be done<br class="">
after inlining. When using LTO, this should be done during the link-time phase.<br class="">
As a result, we must move away from our purely-front-end based lowering scheme.<br class="">
The question is what to do instead, and how to do it in a way that is generally<br class="">
useful to the entire community.<br class="">
<br class="">
Designs previously discussed can be classified into four categories:<br class="">
<br class="">
(a) Add a large number of new kinds of LLVM metadata, and use them to annotate<br class="">
each necessary instruction for parallelism, data attributes, etc.<br class="">
(b) Add several new LLVM instructions such as, for parallelism, fork, spawn,<br class="">
join, barrier, etc.<br class="">
(c) Add a large number of LLVM intrinsics for directives and clauses, each<br class="">
intrinsic representing a directive or a clause.<br class="">
(d) Add a small number of LLVM intrinsics for region or loop annotations,<br class="">
represent the directive/clause names using metadata and the remaining<br class="">
information using arguments.<br class="">
<br class="">
Here we're proposing (d), and below is a brief pros and cons analysis based on<br class="">
these discussions and our own experiences of supporting region/loop annotations<br class="">
in LLVM-based compilers. The table below shows a short summary of our analysis.<br class="">
<br class="">
Various commercial compilers (e.g. from Intel, IBM, Cray, PGI), and GCC [1,2],<br class="">
have IR-level representations for parallelism constructs. Based on experience<br class="">
from these previous developments, we'd like a solution for LLVM that maximizes<br class="">
optimization enablement while minimizing the maintenance costs and complexity<br class="">
increase experienced by the community as a whole.<br class="">
<br class="">
Representing the desired information in the LLVM IR is just the first step. The<br class="">
challenge is to maintain the desired semantics without blocking useful<br class="">
optimizations. With options (c) and (d), dependencies can be preserved mainly<br class="">
based on the use/def chain of the arguments of each intrinsic, and a manageable<br class="">
set LLVM analysis and transformations can be made aware of certain kinds of<br class="">
annotations in order to enable specific optimizations. In this regard,<br class="">
options (c) and (d) are close with respect to maintenance efforts. However,<br class="">
based on our experiences, option (d) is preferable because it is easier to<br class="">
extend to support new directives and clauses in the future without the need to<br class="">
add new intrinsics as required by option (c).<br class="">
<br class="">
Table 1. Pros/cons summary of LLVM IR experimental extension options<br class="">
<br class="">
--------+---------------------<wbr class="">-+----------------------------<wbr class="">------------------- <br class="">
Options | Pros | Cons<br class="">
--------+---------------------<wbr class="">-+----------------------------<wbr class="">------------------- <br class="">
(a) | No need to add new | LLVM passes do not always maintain metadata.<br class="">
| instructions or | Need to educate many passes (if not all) to<br class="">
| new intrinsics | understand and handle them.<br class="">
--------+---------------------<wbr class="">-+----------------------------<wbr class="">------------------- <br class="">
(b) | Parallelism becomes | Huge effort for extending all LLVM passes and<br class="">
| first class citizen | code generation to support new instructions.<br class="">
| | A large set of information still needs to be<br class="">
| | represented using other means.<br class="">
--------+---------------------<wbr class="">-+----------------------------<wbr class="">------------------- <br class="">
(c) | Less impact on the | A large number of intrinsics must be added.<br class="">
| exist LLVM passes. | Some of the optimizations need to be<br class="">
| Fewer requirements | educated to understand them.<br class="">
| for passes to |<br class="">
| maintain metadata. |<br class="">
--------+---------------------<wbr class="">-+----------------------------<wbr class="">------------------- <br class="">
(d) | Minimal impact on | Some of the optimizations need to be<br class="">
| existing LLVM | educated to understand them.<br class="">
| optimizations passes.| No requirements for all passes to maintain<br class="">
| directive and clause | large set of metadata with values.<br class="">
| names use metadata |<br class="">
| strings. |<br class="">
--------+---------------------<wbr class="">-+----------------------------<wbr class="">------------------- <br class="">
<br class="">
Regarding (a), LLVM already uses metadata for certain loop information (e.g.<br class="">
annotations directing loop transformations and assertions about loop-carried<br class="">
dependencies), but there is no natural or consistent way to extend this scheme<br class="">
to represent necessary data-movement or region information.<br class="">
<br class="">
<br class="">
New Intrinsics for Region and Value Annotations<br class="">
==============================<wbr class="">================<br class="">
The following new (experimental) intrinsics are proposed which allow:<br class="">
<br class="">
a) Annotating a code region marked with directives / pragmas,<br class="">
b) Annotating values associated with the region (or loops), that is, those<br class="">
values associated with directives / pragmas.<br class="">
c) Providing information on LLVM IR transformations needed for the annotated<br class="">
code regions (or loops).<br class="">
<br class="">
These can be used both by frontends and also by transformation passes (e.g.<br class="">
automated parallelization). The names used here are similar to those used by<br class="">
our internal prototype, but obviously we expect a community bikeshed<br class="">
discussion.<br class="">
<br class="">
def int_experimental_directive : Intrinsic<[], [llvm_metadata_ty],<br class="">
[IntrArgMemOnly],<br class="">
"llvm.experimental.directive"><wbr class="">;<br class="">
<br class="">
def int_experimental_dir_qual : Intrinsic<[], [llvm_metadata_ty],<br class="">
[IntrArgMemOnly],<br class="">
"llvm.experimental.dir.qual">;<br class="">
<br class="">
def int_experimental_dir_qual_opnd : Intrinsic<[],<br class="">
[llvm_metadata_ty, llvm_any_ty],<br class="">
[IntrArgMemOnly],<br class="">
"llvm.experimental.dir.qual.op<wbr class="">nd">;<br class="">
<br class="">
def int_experimental_dir_qual_opnd<wbr class="">list : Intrinsic<<br class="">
[],<br class="">
[llvm_metadata_ty, llvm_vararg_ty],<br class="">
[IntrArgMemOnly],<br class="">
"llvm.experimental.dir.qual.op<wbr class="">ndlist">;<br class="">
<br class="">
Note that calls to these intrinsics might need to be annotated with the<br class="">
convergent attribute when they represent fork/join operations, barriers, and<br class="">
similar.<br class="">
<br class="">
Usage Examples<br class="">
==============<br class="">
<br class="">
This section shows a few examples using these experimental intrinsics.<br class="">
LLVM developers who will use these intrinsics can defined their own MDstring.<br class="">
All details of using these intrinsics on representing OpenMP 4.5 constructs are described in [1][3].<br class="">
<br class="">
<br class="">
Example I: An OpenMP combined construct<br class="">
<br class="">
#pragma omp target teams distribute parallel for simd<br class="">
loop<br class="">
<br class="">
LLVM IR<br class="">
-------<br class="">
call void @llvm.experimental.directive(m<wbr class="">etadata !0)<br class="">
call void @llvm.experimental.directive(m<wbr class="">etadata !1)<br class="">
call void @llvm.experimental.directive(m<wbr class="">etadata !2)<br class="">
call void @llvm.experimental.directive(m<wbr class="">etadata !3)<br class="">
loop<br class="">
call void @llvm.experimental.directive(m<wbr class="">etadata !6)<br class="">
call void @llvm.experimental.directive(m<wbr class="">etadata !5)<br class="">
call void @llvm.experimental.directive(m<wbr class="">etadata !4)<br class="">
<br class="">
!0 = metadata !{metadata !DIR.OMP.TARGET}<br class="">
!1 = metadata !{metadata !DIR.OMP.TEAMS}<br class="">
!2 = metadata !{metadata !<a href="http://dir.omp.distribute.parloop.si/" target="_blank" class="">DIR.OMP.DISTRIBUTE.PARLOOP.SI</a><wbr class="">MD}<br class="">
<br class="">
!6 = metadata !{metadata !DIR.OMP.END.DISTRIBUTE.PARLOO<wbr class="">P.SIMD}<br class="">
!5 = metadata !{metadata !DIR.OMP.END.TEAMS}<br class="">
!4 = metadata !{metadata !DIR.OMP.END.TARGET}<br class="">
<br class="">
Example II: Assume x,y,z are int variables, and s is a non-POD variable.<br class="">
Then, lastprivate(x,y,s,z) is represented as:<br class="">
<br class="">
LLVM IR<br class="">
-------<br class="">
call void @llvm.experimental.dir.qual.op<wbr class="">ndlist(<br class="">
metadata !1, %x, %y, metadata !2, %a, %ctor, %dtor, %z)<br class="">
<br class="">
!1 = metadata !{metadata !QUAL.OMP.PRIVATE}<br class="">
!2 = metadata !{metadata !QUAL.OPND.NONPOD}<br class="">
<br class="">
Example III: A prefetch pragma example<br class="">
<br class="">
// issue vprefetch1 for xp with a distance of 20 vectorized iterations ahead<br class="">
// issue vprefetch0 for yp with a distance of 10 vectorized iterations ahead<br class="">
#pragma prefetch x:1:20 y:0:10<br class="">
for (i=0; i<2*N; i++) { xp[i*m + j] = -1; yp[i*n +j] = -2; }<br class="">
<br class="">
LLVM IR<br class="">
-------<br class="">
call void @llvm.experimental.directive(m<wbr class="">etadata !0)<br class="">
call void @llvm.experimental.dir.qual.op<wbr class="">nslist(metadata !1, %xp, 1, 20,<br class="">
metadata !1, %yp, 0, 10)<br class="">
loop<br class="">
call void @llvm.experimental.directive(m<wbr class="">etadata !3)<br class="">
<br class="">
References<br class="">
==========<br class="">
<br class="">
[1] LLVM Framework and IR extensions for Parallelization, SIMD Vectorization<br class="">
and Offloading Support. SC'2016 LLVM-HPC3 Workshop. (Xinmin Tian <a href="http://et.al/" rel="noreferrer" target="_blank" class="">et.al</a>.)<br class="">
Saltlake City, Utah.<br class="">
<br class="">
[2] Extending LoopVectorizer towards supporting OpenMP4.5 SIMD and outer loop<br class="">
auto-vectorization. (Hideki Saito, <a href="http://et.al/" rel="noreferrer" target="_blank" class="">et.al</a>.) LLVM Developers' Meeting 2016,<br class="">
San Jose.<br class="">
<br class="">
[3] Intrinsics, Metadata, and Attributes: The Story continues! (Hal Finkel)<br class="">
LLVM Developers' Meeting, 2016. San Jose<br class="">
<br class="">
[4] LLVM Intrinsic Function and Metadata String Interface for Directive (or<br class="">
Pragmas) Representation. Specification Draft v0.9, Intel Corporation, 2016.<br class="">
<br class="">
<br class="">
Acknowledgements<br class="">
================<br class="">
We would like to thank Chandler Carruth (Google), Johannes Doerfert (Saarland<br class="">
Univ.), Yaoqing Gao (HuaWei), Michael Wong (Codeplay), Ettore Tiotto,<br class="">
Carlo Bertolli, Bardia Mahjour (IBM), and all other LLVM-HPC IR Extensions WG<br class="">
members for their constructive feedback on the LLVM framework and IR extension<br class="">
proposal.<br class="">
<br class="">
Proposed Implementation<br class="">
=======================<br class="">
<br class="">
Two sets of patches of supporting these experimental intrinsics and demonstrate<br class="">
the usage are ready for community review.<br class="">
<br class="">
a) Clang patches that support core OpenMP pragmas using this approach.<br class="">
b) W-Region framework patches: CFG restructuring to form single-entry-<br class="">
single-exit work region (W-Region) based on annotations, Demand-driven<br class="">
intrinsic parsing, and WRegionInfo collection and analysis passes,<br class="">
Dump functions of WRegionInfo.<br class="">
<br class="">
On top of this functionality, we will provide the transformation patches for<br class="">
core OpenMP constructs (e.g. start with "#pragma omp parallel for" loop for<br class="">
lowering and outlining, and "#pragma omp simd" to hook it up with<br class="">
LoopVectorize.cpp). We have internal implementations for many constructs now.<br class="">
We will break this functionality up to create a series of patches for<br class="">
community review.<span class="m_8178307607774293309HOEnZb"><font color="#888888" class=""><br class="">
<br class="">
-- <br class="">
Hal Finkel<br class="">
Lead, Compiler Technology and Programming Languages<br class="">
Leadership Computing Facility<br class="">
Argonne National Laboratory<br class="">
<br class="">
______________________________<wbr class="">_________________<br class="">
LLVM Developers mailing list<br class="">
<a href="mailto:llvm-dev@lists.llvm.org" target="_blank" class="">llvm-dev@lists.llvm.org</a><br class="">
<a href="http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev" rel="noreferrer" target="_blank" class="">http://lists.llvm.org/cgi-bin/<wbr class="">mailman/listinfo/llvm-dev</a><br class="">
</font></span></blockquote></div><br class=""></div>
</div></div><br class="">______________________________<wbr class="">_________________<br class="">
LLVM Developers mailing list<br class="">
<a href="mailto:llvm-dev@lists.llvm.org" class="">llvm-dev@lists.llvm.org</a><br class="">
<a href="http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev" rel="noreferrer" target="_blank" class="">http://lists.llvm.org/cgi-bin/<wbr class="">mailman/listinfo/llvm-dev</a><br class="">
<br class=""></blockquote></div><br class=""></div>
_______________________________________________<br class="">LLVM Developers mailing list<br class=""><a href="mailto:llvm-dev@lists.llvm.org" class="">llvm-dev@lists.llvm.org</a><br class="">http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev<br class=""></div></blockquote></div><br class=""></body></html>