similar to: [RFC] IR-level Region Annotations

David, one quick question, is there a way to preserve and associate a set of “properties, value info/attr ” to the given region using Token? Thanks, Xinmin From: llvm-dev [mailto:llvm-dev-bounces at lists.llvm.org] On Behalf Of David Majnemer via llvm-dev Sent: Wednesday, January 11, 2017 2:18 PM To: Hal Finkel <hfinkel at anl.gov> Cc: llvm-dev <llvm-dev at lists.llvm.org> Subject:

+1, tokens are the current True Way to create single-entry multi-exit regions. Your example for an annotated loop would look like: %region = call token @llvm.openmp.regionstart(metadata ...) ; whatever parameters you need here loop call void @llvm.openmp.regionend(token %region) If you use tokens, I would recommend proposal (c), where you introduce new intrinsics for every new kind of region,

Would you send us the LLVM IR for below example using token and OpBundle. So, we can understand better. Thanks. #pragma omp target teams distribute parallel for simd shared(xp, yp) linear(i) firstprivate(m, n) map(m, n) for (i=0; i<2*N; i++) { xp[i*m + j] = -1; yp[i*n +j] = -2; } #pragma prefetch x:1:20 y:0:10 for (i=0; i<2*N; i++) { xp[i*m + j] = -1; yp[i*n +j] = -2; } From: Hongbin

And “map” and “firstprivate” … are represented as MDString, right? Thanks. From: Hongbin Zheng [mailto:etherzhhb at gmail.com] Sent: Wednesday, January 11, 2017 3:58 PM To: Tian, Xinmin <xinmin.tian at intel.com> Cc: David Majnemer <david.majnemer at gmail.com>; Hal Finkel <hfinkel at anl.gov>; llvm-dev at lists.llvm.org Subject: Re: [llvm-dev] [RFC] IR-level Region

Mehdi, thanks for good questions. >>>>>Something isn’t clear to me about how do you preserve the validity of the region annotations since regular passes don’t know about the attached semantic? There are some small changes we have to make in some optimizations to make sure the optimizations does not validation attached annotation semantics. 1) provide hand-shaking / query utils for

Interesting, this is similar to what we have. One more question, these stuff in the yellow, are they represented as LLVM VALUEs? In other words, does the LLVM optimizer update them? ,E.g. %m is re-named %m.1 in the loop, is the “m” in the token @..... is updated as well? In the RFC, the “m” is argument of intrinsic call, all use-def info are used by optimizer, and optimizer updates them

> (d) Add a small number of LLVM intrinsics for region or loop annotations, > represent the directive/clause names using metadata and the remaining > information using arguments. > > Here we're proposing (d), I think this would serve the goal of communicating source-level directives and annotations down to LLVM passes and back-ends, while deferring inlining and

On 01/11, Daniel Berlin via llvm-dev wrote: > > > > def int_experimental_directive : Intrinsic<[], [llvm_metadata_ty], > > [IntrArgMemOnly], > > "llvm.experimental.directive">; > > > > def int_experimental_dir_qual : Intrinsic<[], [llvm_metadata_ty], > > [IntrArgMemOnly], > >

Hi Sanjoy, Yes, that's exactly what we have been looking at recently here, but the region tags seem to make it possible to express the control flow as well, so I think we could start with reg ions+metadata, as Hal and Xinmin proposed, and then figure out what needs to be first class instructions. --Vikram Adve > On Jan 19, 2017, at 11:03 PM, Sanjoy Das <sanjoy at

> On Jan 20, 2017, at 10:44 AM, Tian, Xinmin via llvm-dev <llvm-dev at lists.llvm.org> wrote: > > Sanjoy, the IR would be like something below. It is ok to hoist alloca instruction outside the region. There are some small changes in optimizer to understand region-annotation intrinsic. > > { void main() { > i32* val = alloca i32 > tok =

Hi, I'm going to club together some responses. I agree that outlining function sub-bodies and passing in the function pointers to said outlined bodies to OpenMP helpers lets us correctly implement the semantics we need. However, unless I severely misunderstood the thread, I thought the key idea was to move *away* from that representation and towards a representation that _allows_

> On Jan 20, 2017, at 11:17 AM, Tian, Xinmin <xinmin.tian at intel.com> wrote: > >>>>> This means that the optimizer has to be aware of it, I’m missing the magic here? > > This is one option. > > The another option is that, as I mentioned in our LLVM-HPC paper in our implementation. We have a "prepare phase for pre-privatization" can be invoked

> On Jan 31, 2017, at 5:38 PM, Tian, Xinmin <xinmin.tian at intel.com> wrote: > >>>>> Ok, but this looks like a “workaround" for your specific use-case, I don’t see how it can scale as a model-agnostic and general-purpose region semantic. > > I would say it is a design trade-off. I’m not sure if we’re talking about the same thing here: my understanding at

From: mehdi.amini at apple.com [mailto:mehdi.amini at apple.com] Sent: Tuesday, January 31, 2017 9:03 PM To: Tian, Xinmin <xinmin.tian at intel.com> Cc: Sanjoy Das <sanjoy at playingwithpointers.com>; Adve, Vikram Sadanand <vadve at illinois.edu>; llvm-dev at lists.llvm.org; llvm-dev-request at lists.llvm.org Subject: Re: [llvm-dev] [RFC] IR-level Region Annotations On Jan 31,

[XT] Back from Biz trips, trying to catch up with the discussion. >>>>I agree that outlining function sub-bodies and passing in the function pointers to said outlined bodies to OpenMP helpers lets us correctly implement the semantics we need. However, unless I severely misunderstood the thread, I thought the key idea was to move *away* from that representation and towards a

>>>>Ok, but this looks like a “workaround" for your specific use-case, I don’t see how it can scale as a model-agnostic and general-purpose region semantic. I would say it is a design trade-off. Regardless it is a new instruction or an intrinsics with token/tag, it will consist of model-agnostic part and model-non-agnostic part. The package comes with a framework for parsing

In this case, inliner is educated to add all local variables to the tag of enclosing parallel region, if there is enclosing parallel region. In our icc implementation, it is even simple, as we have routine level symbol table, the inliner adds ”private” attribute to those local variables w/o checking enclosing scope, the parallelizer does check and use it. Xinmin From: mehdi.amini at apple.com

> On Jan 31, 2017, at 10:59 PM, Tian, Xinmin <xinmin.tian at intel.com> wrote: > > >   <> > From: mehdi.amini at apple.com <mailto:mehdi.amini at apple.com> [mailto:mehdi.amini at apple.com <mailto:mehdi.amini at apple.com>] > Sent: Tuesday, January 31, 2017 9:03 PM > To: Tian, Xinmin <xinmin.tian at intel.com <mailto:xinmin.tian at

Let me try this. You can simply consider the prepare-phase (e.g. pre-privatization) were done in FE (actually a library can be used by multiple FEs at LLVM IR level), the region is run with 1 thread, region annotation (scope, single-entry-single-exit) as memory barrier conservatively for now (instead of checking individual memory dependency, aliasing via tags which is the actual

> On Jan 31, 2017, at 6:48 PM, Tian, Xinmin <xinmin.tian at intel.com> wrote: > > Let me try this. > > You can simply consider the prepare-phase (e.g. pre-privatization) were done in FE (actually a library can be used by multiple FEs at LLVM IR level), the region is run with 1 thread, region annotation (scope, single-entry-single-exit) as memory barrier conservatively