<div dir="ltr">Hi Jonas,<div><br></div><div>Sure, we are trying to do so. The public lists often reject my emails because it is large and I cannot fill all people in that mailing list here, though.</div><div><br></div><div>Thanks,</div><div>Lingda Li</div></div><br><div class="gmail_quote"><div dir="ltr" class="gmail_attr">On Sat, Jun 29, 2019 at 8:39 AM Jonas Hahnfeld <<a href="mailto:hahnjo@hahnjo.de">hahnjo@hahnjo.de</a>> wrote:<br></div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex">Hi Lingda,<br>
<br>
may I ask to start discussions about important decisions related to <br>
Clang's OpenMP support on the public mailing list instead of having <br>
private conversations? That would help to get feedback from people not <br>
being part of the selected circle participating in the "OpenMP / HPC in <br>
Clang / LLVM Multi-company Telecom".<br>
<br>
Thanks,<br>
Jonas<br>
<br>
On 2019-06-28 15:59, Lingda Li via cfe-dev wrote:<br>
> On Fri, Jun 28, 2019 at 9:49 AM Li, Lingda <<a href="mailto:lli@bnl.gov" target="_blank">lli@bnl.gov</a>> wrote:<br>
> <br>
>> I don't think we can have the buffer allocated within the mapper<br>
>> function. It has to be done in the runtime, because of nested<br>
>> mappers.<br>
>> First, all mapper functions are born in the same way. We cannot<br>
>> make the outer most mapper function allocate memory, whether the<br>
>> inner one doesn't and has to use what is allocated by the outer most<br>
>> mapper function.<br>
>> I suppose we still need to allocate memory in the runtime, so the<br>
>> runtime can pass the pointer and size to the mapper function, and<br>
>> the outer mapper function can then pass them into inner ones.<br>
>> Again, this is just like the current implementation, except that we<br>
>> don't use vecter::push_back(), instead we use something like a<br>
>> manual implementation of vector::push_back() (because we need to use<br>
>> the pointer and the current index)<br>
>> <br>
>> I believe the key question here is whether it is true that (the<br>
>> overhead of push_back() > the overhead of precalculating the total<br>
>> number + the memory allocation overhead + directly memory write).<br>
>> This will decide whether this change is necessary. Any opinions?<br>
>> <br>
>> Thanks,<br>
>> Lingda Li<br>
>> <br>
>> -------------------------<br>
>> <br>
>> FROM: Alexey Bataev <<a href="mailto:Alexey.Bataev@ibm.com" target="_blank">Alexey.Bataev@ibm.com</a>><br>
>> SENT: Thursday, June 27, 2019 5:05 PM<br>
>> TO: Li, Lingda<br>
>> CC: Alexandre Eichenberger; Chapman, Barbara (Contact); Kevin K<br>
>> O'Brien; Carlo Bertolli; Deepak Eachempati; Denny, Joel E.; David<br>
>> Oehmke; Ettore Tiotto; <a href="mailto:fraggamuffin@gmail.com" target="_blank">fraggamuffin@gmail.com</a>; Rokos, Georgios;<br>
>> Gheorghe-Teod Bercea; <a href="mailto:gregory.rodgers@amd.com" target="_blank">gregory.rodgers@amd.com</a>; Hal Finkel; Sharif,<br>
>> Hashim; Cownie, James H; Sjodin, Jan; <a href="mailto:jbeyer@nvidia.com" target="_blank">jbeyer@nvidia.com</a>; Doerfert,<br>
>> Johannes Rudolf; Jones, Jeff C; <a href="mailto:josem@udel.edu" target="_blank">josem@udel.edu</a>; Robichaux, Joseph;<br>
>> Jeff Heath; <a href="mailto:khaldi.dounia@gmail.com" target="_blank">khaldi.dounia@gmail.com</a>; Kelvin Li; Bobrovsky,<br>
>> Konstantin S; Kotsifakou, Maria; <a href="mailto:lopezmg@ornl.org" target="_blank">lopezmg@ornl.org</a>; Lopez, Matthew<br>
>> Graham; Menard, Lorri; Martin Kong; Sarah McNamara; Rice, Michael P;<br>
>> Matt Martineau; <a href="mailto:oscar@ornl.gov" target="_blank">oscar@ornl.gov</a>; Jeeva Paudel; Rao, Premanand M;<br>
>> Krishnaiyer, Rakesh; Narayanaswamy, Ravi; Monteleone, Robert;<br>
>> Lieberman, Ron; Samuel Antao; Jeffrey Sandoval; Sunita<br>
>> Chandrasekaran; <a href="mailto:sergey.y.ostanevich@gmail.com" target="_blank">sergey.y.ostanevich@gmail.com</a>; Sergio Pino Gallardo;<br>
>> Dmitriev, Serguei N; Chan, SiuChi; Sunil Shrestha; Wilmarth, Terry<br>
>> L; Tianyi Zhang; <a href="mailto:vadve@illinois.edu" target="_blank">vadve@illinois.edu</a>; Wang Chen; Wael Yehia; Tian,<br>
>> Xinmin<br>
>> SUBJECT: Re: Re: Re: RE: Comparison of 2 schemes to implement OpenMP<br>
>> 5.0 declare mapper codegen<br>
>> <br>
>> Yes, we need 2 functions, but thw first one can be optimized very<br>
>> effectively. After the optimizations and inlining it will end up<br>
>> with just return s1+s2+s3... I think, inost cases those sizes will<br>
>> be constant, since the mapper maps constant number of elements. And,<br>
>> thus, this expression will be optimized to just a constant value.<br>
>> You don't need to pass these functions to runtime. We can call the<br>
>> directly from the compiler.<br>
>> 1st call: get number of elements.<br>
>> 2nd: allocate the buffer<br>
>> 3rd call: call mapper with this preallocated buffer that fills this<br>
>> buffer without any calls of the runtime functions.<br>
>> 4th call: call the runtime to pass the buffer to the runtime.<br>
>> <br>
>> Best regards,<br>
>> Alexey Bataev<br>
>> <br>
>> 27 июня 2019 г., в 16:53, Li, Lingda <<a href="mailto:lli@bnl.gov" target="_blank">lli@bnl.gov</a>><br>
>> написал(а):<br>
>> <br>
>>> If we precalculate the size, first, it means we need to generate<br>
>>> 2 functions for each mapper, rather than 1 now. One for mapping<br>
>>> information filling as we have, the other for size calculation<br>
>>> (This will not return constant values, because size depends on how<br>
>>> many instances we are mapping). Both these 2 functions will need<br>
>>> to be passed to the runtime. The runtime will need to precalculate<br>
>>> the number of components first, then allocate memory, then call<br>
>>> the mapper function to fill it up.<br>
>>> <br>
>>> Compared with the scheme 1, the differences are:<br>
>>> 1) An extra call to calculate the total number, while scheme 1<br>
>>> does not;<br>
>>> 2) A preallocated buffer, whose pointer and the current number<br>
>>> should be passed to the mapper function, then the mapper function<br>
>>> uses them to fill components, while scheme 1 uses push_back() to<br>
>>> do the same thing.<br>
>>> <br>
>>> Is there really a benefit doing this? push_back() should be<br>
>>> efficient enough compared with directly writing to memory.<br>
>>> <br>
>>> If people here think that, the overhead of push_back() > the<br>
>>> overhead of precalculating the total number + the memory<br>
>>> allocation overhead + directly memory write, then we can consider<br>
>>> this scheme.<br>
>>> <br>
>>> Thanks,<br>
>>> Lingda Li<br>
>>> <br>
>>> -------------------------<br>
>>> <br>
>>> FROM: Alexey Bataev <<a href="mailto:Alexey.Bataev@ibm.com" target="_blank">Alexey.Bataev@ibm.com</a>><br>
>>> SENT: Thursday, June 27, 2019 4:26 PM<br>
>>> TO: Li, Lingda<br>
>>> CC: Alexandre Eichenberger; Chapman, Barbara (Contact); Kevin K<br>
>>> O'Brien; Carlo Bertolli; Deepak Eachempati; Denny, Joel E.; David<br>
>>> Oehmke; Ettore Tiotto; <a href="mailto:fraggamuffin@gmail.com" target="_blank">fraggamuffin@gmail.com</a>; Rokos, Georgios;<br>
>>> Gheorghe-Teod Bercea; <a href="mailto:gregory.rodgers@amd.com" target="_blank">gregory.rodgers@amd.com</a>; Hal Finkel; Sharif,<br>
>>> Hashim; Cownie, James H; Sjodin, Jan; <a href="mailto:jbeyer@nvidia.com" target="_blank">jbeyer@nvidia.com</a>; Doerfert,<br>
>>> Johannes Rudolf; Jones, Jeff C; <a href="mailto:josem@udel.edu" target="_blank">josem@udel.edu</a>; Robichaux, Joseph;<br>
>>> Jeff Heath; <a href="mailto:khaldi.dounia@gmail.com" target="_blank">khaldi.dounia@gmail.com</a>; Kelvin Li; Bobrovsky,<br>
>>> Konstantin S; Kotsifakou, Maria; <a href="mailto:lopezmg@ornl.org" target="_blank">lopezmg@ornl.org</a>; Lopez, Matthew<br>
>>> Graham; Menard, Lorri; Martin Kong; Sarah McNamara; Rice, Michael<br>
>>> P; Matt Martineau; <a href="mailto:oscar@ornl.gov" target="_blank">oscar@ornl.gov</a>; Jeeva Paudel; Rao, Premanand M;<br>
>>> Krishnaiyer, Rakesh; Narayanaswamy, Ravi; Monteleone, Robert;<br>
>>> Lieberman, Ron; Samuel Antao; Jeffrey Sandoval; Sunita<br>
>>> Chandrasekaran; <a href="mailto:sergey.y.ostanevich@gmail.com" target="_blank">sergey.y.ostanevich@gmail.com</a>; Sergio Pino<br>
>>> Gallardo; Dmitriev, Serguei N; Chan, SiuChi; Sunil Shrestha;<br>
>>> Wilmarth, Terry L; Tianyi Zhang; <a href="mailto:vadve@illinois.edu" target="_blank">vadve@illinois.edu</a>; Wang Chen;<br>
>>> Wael Yehia; Tian, Xinmin<br>
>>> SUBJECT: Re: Re: RE: Comparison of 2 schemes to implement OpenMP<br>
>>> 5.0 declare mapper codegen<br>
>>> <br>
>>> If the functions are inlined (the ines, intended for size<br>
>>> precalculation). They can be optimized out very effectively since<br>
>>> in most cases they will return constant values.<br>
>>> If we could do this, we won't need vectors and oush_backs, we can<br>
>>> use preallocated memory and internal counter.<br>
>>> --------------<br>
>>> Best regards,<br>
>>> Alexey Bataev<br>
>>> <br>
>>> <graycol.gif>"Li, Lingda" ---06/27/2019 04:13:03 PM---Hi Alexey, I<br>
>>> think that's why we choose to use variable size storage like<br>
>>> std::vector to store the m<br>
>>> <br>
>>> From: "Li, Lingda" <<a href="mailto:lli@bnl.gov" target="_blank">lli@bnl.gov</a>><br>
>>> To: Alexey Bataev <<a href="mailto:Alexey.Bataev@ibm.com" target="_blank">Alexey.Bataev@ibm.com</a>>, Deepak Eachempati<br>
>>> <<a href="mailto:deachempat@cray.com" target="_blank">deachempat@cray.com</a>><br>
>>> Cc: "Narayanaswamy, Ravi" <<a href="mailto:ravi.narayanaswamy@intel.com" target="_blank">ravi.narayanaswamy@intel.com</a>>,<br>
>>> "Alexandre Eichenberger" <<a href="mailto:alexe@us.ibm.com" target="_blank">alexe@us.ibm.com</a>>, "Chapman, Barbara<br>
>>> (Contact)" <<a href="mailto:barbara.chapman@stonybrook.edu" target="_blank">barbara.chapman@stonybrook.edu</a>>, "Bobrovsky,<br>
>>> Konstantin S" <<a href="mailto:konstantin.s.bobrovsky@intel.com" target="_blank">konstantin.s.bobrovsky@intel.com</a>>, Carlo Bertolli<br>
>>> <<a href="mailto:cbertol@us.ibm.com" target="_blank">cbertol@us.ibm.com</a>>, "Chan, SiuChi" <<a href="mailto:siuchi.chan@amd.com" target="_blank">siuchi.chan@amd.com</a>>,<br>
>>> "Cownie, James H" <<a href="mailto:james.h.cownie@intel.com" target="_blank">james.h.cownie@intel.com</a>>, David Oehmke<br>
>>> <<a href="mailto:doehmke@cray.com" target="_blank">doehmke@cray.com</a>>, "Denny, Joel E." <<a href="mailto:dennyje@ornl.gov" target="_blank">dennyje@ornl.gov</a>>,<br>
>>> "Dmitriev, Serguei N" <<a href="mailto:serguei.n.dmitriev@intel.com" target="_blank">serguei.n.dmitriev@intel.com</a>>, "Doerfert,<br>
>>> Johannes Rudolf" <<a href="mailto:jdoerfert@anl.gov" target="_blank">jdoerfert@anl.gov</a>>, Ettore Tiotto<br>
>>> <<a href="mailto:etiotto@ca.ibm.com" target="_blank">etiotto@ca.ibm.com</a>>, "<a href="mailto:fraggamuffin@gmail.com" target="_blank">fraggamuffin@gmail.com</a>"<br>
>>> <<a href="mailto:fraggamuffin@gmail.com" target="_blank">fraggamuffin@gmail.com</a>>, Gheorghe-Teod Bercea<br>
>>> <<a href="mailto:Gheorghe-Teod.Bercea@ibm.com" target="_blank">Gheorghe-Teod.Bercea@ibm.com</a>>, Hal Finkel <<a href="mailto:hfinkel@anl.gov" target="_blank">hfinkel@anl.gov</a>>,<br>
>>> "<a href="mailto:jbeyer@nvidia.com" target="_blank">jbeyer@nvidia.com</a>" <<a href="mailto:jbeyer@nvidia.com" target="_blank">jbeyer@nvidia.com</a>>, Jeeva Paudel<br>
>>> <<a href="mailto:pjeeva01@ca.ibm.com" target="_blank">pjeeva01@ca.ibm.com</a>>, Jeff Heath <<a href="mailto:jrheath@ca.ibm.com" target="_blank">jrheath@ca.ibm.com</a>>, Jeffrey<br>
>>> Sandoval <<a href="mailto:sandoval@cray.com" target="_blank">sandoval@cray.com</a>>, "Jones, Jeff C"<br>
>>> <<a href="mailto:jeff.c.jones@intel.com" target="_blank">jeff.c.jones@intel.com</a>>, "<a href="mailto:josem@udel.edu" target="_blank">josem@udel.edu</a>" <<a href="mailto:josem@udel.edu" target="_blank">josem@udel.edu</a>>,<br>
>>> Kelvin Li <<a href="mailto:kli@ca.ibm.com" target="_blank">kli@ca.ibm.com</a>>, "Kevin K O'Brien"<br>
>>> <<a href="mailto:caomhin@us.ibm.com" target="_blank">caomhin@us.ibm.com</a>>, "<a href="mailto:khaldi.dounia@gmail.com" target="_blank">khaldi.dounia@gmail.com</a>"<br>
>>> <<a href="mailto:khaldi.dounia@gmail.com" target="_blank">khaldi.dounia@gmail.com</a>>, "Kotsifakou, Maria"<br>
>>> <<a href="mailto:kotsifa2@illinois.edu" target="_blank">kotsifa2@illinois.edu</a>>, "Krishnaiyer, Rakesh"<br>
>>> <<a href="mailto:rakesh.krishnaiyer@intel.com" target="_blank">rakesh.krishnaiyer@intel.com</a>>, "Lieberman, Ron"<br>
>>> <<a href="mailto:Ron.Lieberman@amd.com" target="_blank">Ron.Lieberman@amd.com</a>>, "Lopez, Matthew Graham"<br>
>>> <<a href="mailto:lopezmg@ornl.gov" target="_blank">lopezmg@ornl.gov</a>>, "<a href="mailto:lopezmg@ornl.org" target="_blank">lopezmg@ornl.org</a>" <<a href="mailto:lopezmg@ornl.org" target="_blank">lopezmg@ornl.org</a>>, Martin<br>
>>> Kong <<a href="mailto:martin.richard.kong@gmail.com" target="_blank">martin.richard.kong@gmail.com</a>>, Matt Martineau<br>
>>> <<a href="mailto:m.martineau@bristol.ac.uk" target="_blank">m.martineau@bristol.ac.uk</a>>, "Menard, Lorri"<br>
>>> <<a href="mailto:lorri.menard@intel.com" target="_blank">lorri.menard@intel.com</a>>, "Monteleone, Robert"<br>
>>> <<a href="mailto:robert.monteleone@intel.com" target="_blank">robert.monteleone@intel.com</a>>, "<a href="mailto:oscar@ornl.gov" target="_blank">oscar@ornl.gov</a>" <<a href="mailto:oscar@ornl.gov" target="_blank">oscar@ornl.gov</a>>,<br>
>>> "Rao, Premanand M" <<a href="mailto:premanand.m.rao@intel.com" target="_blank">premanand.m.rao@intel.com</a>>, "Rice, Michael P"<br>
>>> <<a href="mailto:michael.p.rice@intel.com" target="_blank">michael.p.rice@intel.com</a>>, "Robichaux, Joseph"<br>
>>> <<a href="mailto:joseph.robichaux@intel.com" target="_blank">joseph.robichaux@intel.com</a>>, "<a href="mailto:gregory.rodgers@amd.com" target="_blank">gregory.rodgers@amd.com</a>"<br>
>>> <<a href="mailto:gregory.rodgers@amd.com" target="_blank">gregory.rodgers@amd.com</a>>, "Rokos, Georgios"<br>
>>> <<a href="mailto:georgios.rokos@intel.com" target="_blank">georgios.rokos@intel.com</a>>, Samuel Antao <<a href="mailto:Samuel.Antao@ibm.com" target="_blank">Samuel.Antao@ibm.com</a>>,<br>
>>> "Sarah McNamara" <<a href="mailto:mcnamara@ca.ibm.com" target="_blank">mcnamara@ca.ibm.com</a>>,<br>
>>> "<a href="mailto:sergey.y.ostanevich@gmail.com" target="_blank">sergey.y.ostanevich@gmail.com</a>" <<a href="mailto:sergey.y.ostanevich@gmail.com" target="_blank">sergey.y.ostanevich@gmail.com</a>>,<br>
>>> Sergio Pino Gallardo <<a href="mailto:sergiop@udel.edu" target="_blank">sergiop@udel.edu</a>>, "Sharif, Hashim"<br>
>>> <<a href="mailto:hsharif3@illinois.edu" target="_blank">hsharif3@illinois.edu</a>>, "Sjodin, Jan" <<a href="mailto:Jan.Sjodin@amd.com" target="_blank">Jan.Sjodin@amd.com</a>>, Sunil<br>
>>> Shrestha <<a href="mailto:sshrestha@cray.com" target="_blank">sshrestha@cray.com</a>>, Sunita Chandrasekaran<br>
>>> <<a href="mailto:schandra@udel.edu" target="_blank">schandra@udel.edu</a>>, "Tian, Xinmin" <<a href="mailto:xinmin.tian@intel.com" target="_blank">xinmin.tian@intel.com</a>>,<br>
>>> Tianyi Zhang <<a href="mailto:tzhan18@lsu.edu" target="_blank">tzhan18@lsu.edu</a>>, "<a href="mailto:vadve@illinois.edu" target="_blank">vadve@illinois.edu</a>"<br>
>>> <<a href="mailto:vadve@illinois.edu" target="_blank">vadve@illinois.edu</a>>, Wael Yehia <<a href="mailto:wyehia@ca.ibm.com" target="_blank">wyehia@ca.ibm.com</a>>, Wang Chen<br>
>>> <<a href="mailto:wdchen@ca.ibm.com" target="_blank">wdchen@ca.ibm.com</a>>, "Wilmarth, Terry L"<br>
>>> <<a href="mailto:terry.l.wilmarth@intel.com" target="_blank">terry.l.wilmarth@intel.com</a>><br>
>>> Date: 06/27/2019 04:13 PM<br>
>>> Subject: [EXTERNAL] Re: RE: Comparison of 2 schemes to implement<br>
>>> OpenMP 5.0 declare mapper codegen<br>
>>> <br>
>>> -------------------------<br>
>>> <br>
>>> Hi Alexey,<br>
>>> <br>
>>> I think that's why we choose to use variable size storage like<br>
>>> std::vector to store the mapping information at the first place,<br>
>>> right? It'll be costly to precalculate the total number of<br>
>>> components, especially in the presence of nested mappers. Besides,<br>
>>> a runtime function call is just a std::vector::push, so I think<br>
>>> it's okay to have multiple function calls.<br>
>>> <br>
>>> Thanks,<br>
>>> Lingda Li<br>
>>> <br>
>>> -------------------------<br>
>>> <br>
>>> FROM: Alexey Bataev <<a href="mailto:Alexey.Bataev@ibm.com" target="_blank">Alexey.Bataev@ibm.com</a>><br>
>>> Sent: Thursday, June 27, 2019 3:52 PM<br>
>>> To: Deepak Eachempati<br>
>>> Cc: Li, Lingda; Narayanaswamy, Ravi; Alexandre Eichenberger;<br>
>>> Chapman, Barbara (Contact); Bobrovsky, Konstantin S; Carlo<br>
>>> Bertolli; Chan, SiuChi; Cownie, James H; David Oehmke; Denny, Joel<br>
>>> E.; Dmitriev, Serguei N; Doerfert, Johannes Rudolf; Ettore Tiotto;<br>
>>> <a href="mailto:fraggamuffin@gmail.com" target="_blank">fraggamuffin@gmail.com</a>; Gheorghe-Teod Bercea; Hal Finkel;<br>
>>> <a href="mailto:jbeyer@nvidia.com" target="_blank">jbeyer@nvidia.com</a>; Jeeva Paudel; Jeff Heath; Jeffrey Sandoval;<br>
>>> Jones, Jeff C; <a href="mailto:josem@udel.edu" target="_blank">josem@udel.edu</a>; Kelvin Li; Kevin K O'Brien;<br>
>>> <a href="mailto:khaldi.dounia@gmail.com" target="_blank">khaldi.dounia@gmail.com</a>; Kotsifakou, Maria; Krishnaiyer, Rakesh;<br>
>>> Lieberman, Ron; Lopez, Matthew Graham; <a href="mailto:lopezmg@ornl.org" target="_blank">lopezmg@ornl.org</a>; Martin<br>
>>> Kong; Matt Martineau; Menard, Lorri; Monteleone, Robert;<br>
>>> <a href="mailto:oscar@ornl.gov" target="_blank">oscar@ornl.gov</a>; Rao, Premanand M; Rice, Michael P; Robichaux,<br>
>>> Joseph; <a href="mailto:gregory.rodgers@amd.com" target="_blank">gregory.rodgers@amd.com</a>; Rokos, Georgios; Samuel Antao;<br>
>>> Sarah McNamara; <a href="mailto:sergey.y.ostanevich@gmail.com" target="_blank">sergey.y.ostanevich@gmail.com</a>; Sergio Pino<br>
>>> Gallardo; Sharif, Hashim; Sjodin, Jan; Sunil Shrestha; Sunita<br>
>>> Chandrasekaran; Tian, Xinmin; Tianyi Zhang; <a href="mailto:vadve@illinois.edu" target="_blank">vadve@illinois.edu</a>;<br>
>>> Wael Yehia; Wang Chen; Wilmarth, Terry L<br>
>>> Subject: Re: RE: Comparison of 2 schemes to implement OpenMP 5.0<br>
>>> declare mapper codegen<br>
>>> <br>
>>> Lingda, can we in scheme 1 precalculate the total number of<br>
>>> components, allocate memory for these precalculate number of<br>
>>> elements, then fill it with mappers and only after that call the<br>
>>> runtime function (only once!) to transfer the mappings to the<br>
>>> runtime?<br>
>>> <br>
>>> Best regards,<br>
>>> Alexey Bataev<br>
>>> <br>
>>> 27 июня 2019 г., в 15:44, Deepak Eachempati<br>
>>> <<a href="mailto:deachempat@cray.com" target="_blank">deachempat@cray.com</a>> написал(а):<br>
>>> <br>
>>> Got it. Thanks.<br>
>>> <br>
>>> -- Deepak<br>
>>> <br>
>>> FROM: Li, Lingda [mailto:<a href="mailto:lli@bnl.gov" target="_blank">lli@bnl.gov</a>]<br>
>>> Sent: Thursday, June 27, 2019 2:41 PM<br>
>>> To: Deepak Eachempati <<a href="mailto:deachempat@cray.com" target="_blank">deachempat@cray.com</a>>; Narayanaswamy, Ravi<br>
>>> <<a href="mailto:ravi.narayanaswamy@intel.com" target="_blank">ravi.narayanaswamy@intel.com</a>>; 'Alexandre Eichenberger'<br>
>>> <<a href="mailto:alexe@us.ibm.com" target="_blank">alexe@us.ibm.com</a>>; 'Alexey Bataev' <<a href="mailto:Alexey.Bataev@ibm.com" target="_blank">Alexey.Bataev@ibm.com</a>>;<br>
>>> Chapman, Barbara (Contact) <<a href="mailto:barbara.chapman@stonybrook.edu" target="_blank">barbara.chapman@stonybrook.edu</a>>;<br>
>>> Bobrovsky, Konstantin S <<a href="mailto:konstantin.s.bobrovsky@intel.com" target="_blank">konstantin.s.bobrovsky@intel.com</a>>; 'Carlo<br>
>>> Bertolli' <<a href="mailto:cbertol@us.ibm.com" target="_blank">cbertol@us.ibm.com</a>>; 'Chan, SiuChi'<br>
>>> <<a href="mailto:siuchi.chan@amd.com" target="_blank">siuchi.chan@amd.com</a>>; Cownie, James H <<a href="mailto:james.h.cownie@intel.com" target="_blank">james.h.cownie@intel.com</a>>;<br>
>>> David Oehmke <<a href="mailto:doehmke@cray.com" target="_blank">doehmke@cray.com</a>>; 'Denny, Joel E.'<br>
>>> <<a href="mailto:dennyje@ornl.gov" target="_blank">dennyje@ornl.gov</a>>; Dmitriev, Serguei N<br>
>>> <<a href="mailto:serguei.n.dmitriev@intel.com" target="_blank">serguei.n.dmitriev@intel.com</a>>; Doerfert, Johannes Rudolf<br>
>>> <<a href="mailto:jdoerfert@anl.gov" target="_blank">jdoerfert@anl.gov</a>>; 'Ettore Tiotto' <<a href="mailto:etiotto@ca.ibm.com" target="_blank">etiotto@ca.ibm.com</a>>;<br>
>>> '<a href="mailto:fraggamuffin@gmail.com" target="_blank">fraggamuffin@gmail.com</a>' <<a href="mailto:fraggamuffin@gmail.com" target="_blank">fraggamuffin@gmail.com</a>>; 'Gheorghe-Teod<br>
>>> Bercea' <<a href="mailto:Gheorghe-Teod.Bercea@ibm.com" target="_blank">Gheorghe-Teod.Bercea@ibm.com</a>>; Hal Finkel<br>
>>> <<a href="mailto:hfinkel@anl.gov" target="_blank">hfinkel@anl.gov</a>>; '<a href="mailto:jbeyer@nvidia.com" target="_blank">jbeyer@nvidia.com</a>' <<a href="mailto:jbeyer@nvidia.com" target="_blank">jbeyer@nvidia.com</a>>; 'Jeeva<br>
>>> Paudel' <<a href="mailto:pjeeva01@ca.ibm.com" target="_blank">pjeeva01@ca.ibm.com</a>>; 'Jeff Heath' <<a href="mailto:jrheath@ca.ibm.com" target="_blank">jrheath@ca.ibm.com</a>>;<br>
>>> Jeffrey Sandoval <<a href="mailto:sandoval@cray.com" target="_blank">sandoval@cray.com</a>>; Jones, Jeff C<br>
>>> <<a href="mailto:jeff.c.jones@intel.com" target="_blank">jeff.c.jones@intel.com</a>>; '<a href="mailto:josem@udel.edu" target="_blank">josem@udel.edu</a>' <<a href="mailto:josem@udel.edu" target="_blank">josem@udel.edu</a>>;<br>
>>> 'Kelvin Li' <<a href="mailto:kli@ca.ibm.com" target="_blank">kli@ca.ibm.com</a>>; 'Kevin K O'Brien'<br>
>>> <<a href="mailto:caomhin@us.ibm.com" target="_blank">caomhin@us.ibm.com</a>>; '<a href="mailto:khaldi.dounia@gmail.com" target="_blank">khaldi.dounia@gmail.com</a>'<br>
>>> <<a href="mailto:khaldi.dounia@gmail.com" target="_blank">khaldi.dounia@gmail.com</a>>; 'Kotsifakou, Maria'<br>
>>> <<a href="mailto:kotsifa2@illinois.edu" target="_blank">kotsifa2@illinois.edu</a>>; Krishnaiyer, Rakesh<br>
>>> <<a href="mailto:rakesh.krishnaiyer@intel.com" target="_blank">rakesh.krishnaiyer@intel.com</a>>; Lieberman, Ron<br>
>>> <<a href="mailto:Ron.Lieberman@amd.com" target="_blank">Ron.Lieberman@amd.com</a>>; 'Lopez, Matthew Graham'<br>
>>> <<a href="mailto:lopezmg@ornl.gov" target="_blank">lopezmg@ornl.gov</a>>; '<a href="mailto:lopezmg@ornl.org" target="_blank">lopezmg@ornl.org</a>' <<a href="mailto:lopezmg@ornl.org" target="_blank">lopezmg@ornl.org</a>>; 'Martin<br>
>>> Kong' <<a href="mailto:martin.richard.kong@gmail.com" target="_blank">martin.richard.kong@gmail.com</a>>; 'Matt Martineau'<br>
>>> <<a href="mailto:m.martineau@bristol.ac.uk" target="_blank">m.martineau@bristol.ac.uk</a>>; Menard, Lorri<br>
>>> <<a href="mailto:lorri.menard@intel.com" target="_blank">lorri.menard@intel.com</a>>; Monteleone, Robert<br>
>>> <<a href="mailto:robert.monteleone@intel.com" target="_blank">robert.monteleone@intel.com</a>>; <a href="mailto:oscar@ornl.gov" target="_blank">oscar@ornl.gov</a>; Rao, Premanand M<br>
>>> <<a href="mailto:premanand.m.rao@intel.com" target="_blank">premanand.m.rao@intel.com</a>>; Rice, Michael P<br>
>>> <<a href="mailto:michael.p.rice@intel.com" target="_blank">michael.p.rice@intel.com</a>>; Robichaux, Joseph<br>
>>> <<a href="mailto:joseph.robichaux@intel.com" target="_blank">joseph.robichaux@intel.com</a>>; <a href="mailto:gregory.rodgers@amd.com" target="_blank">gregory.rodgers@amd.com</a>; Rokos,<br>
>>> Georgios <<a href="mailto:georgios.rokos@intel.com" target="_blank">georgios.rokos@intel.com</a>>; '<a href="mailto:samuel.antao@ibm.com" target="_blank">samuel.antao@ibm.com</a>'<br>
>>> <<a href="mailto:samuel.antao@ibm.com" target="_blank">samuel.antao@ibm.com</a>>; 'Sarah McNamara' <<a href="mailto:mcnamara@ca.ibm.com" target="_blank">mcnamara@ca.ibm.com</a>>;<br>
>>> '<a href="mailto:sergey.y.ostanevich@gmail.com" target="_blank">sergey.y.ostanevich@gmail.com</a>' <<a href="mailto:sergey.y.ostanevich@gmail.com" target="_blank">sergey.y.ostanevich@gmail.com</a>>;<br>
>>> 'Sergio Pino Gallardo' <<a href="mailto:sergiop@udel.edu" target="_blank">sergiop@udel.edu</a>>; 'Sharif, Hashim'<br>
>>> <<a href="mailto:hsharif3@illinois.edu" target="_blank">hsharif3@illinois.edu</a>>; Sjodin, Jan <<a href="mailto:Jan.Sjodin@amd.com" target="_blank">Jan.Sjodin@amd.com</a>>; Sunil<br>
>>> Shrestha <<a href="mailto:sshrestha@cray.com" target="_blank">sshrestha@cray.com</a>>; 'Sunita Chandrasekaran'<br>
>>> <<a href="mailto:schandra@udel.edu" target="_blank">schandra@udel.edu</a>>; Tian, Xinmin <<a href="mailto:xinmin.tian@intel.com" target="_blank">xinmin.tian@intel.com</a>>; Tianyi<br>
>>> Zhang <<a href="mailto:tzhan18@lsu.edu" target="_blank">tzhan18@lsu.edu</a>>; '<a href="mailto:vadve@illinois.edu" target="_blank">vadve@illinois.edu</a>'<br>
>>> <<a href="mailto:vadve@illinois.edu" target="_blank">vadve@illinois.edu</a>>; 'Wael Yehia' <<a href="mailto:wyehia@ca.ibm.com" target="_blank">wyehia@ca.ibm.com</a>>; 'Wang<br>
>>> Chen' <<a href="mailto:wdchen@ca.ibm.com" target="_blank">wdchen@ca.ibm.com</a>>; Wilmarth, Terry L<br>
>>> <<a href="mailto:terry.l.wilmarth@intel.com" target="_blank">terry.l.wilmarth@intel.com</a>><br>
>>> Subject: Re: Comparison of 2 schemes to implement OpenMP 5.0<br>
>>> declare mapper codegen<br>
>>> <br>
>>> In the current scheme, all mappings within a mapper function is<br>
>>> done atomically by one thread. In the mapper function of the<br>
>>> example in the original email, <push> will just push the mapping<br>
>>> information into an internal data structure. Once all mapping<br>
>>> information is available, the runtime will do the real mapping<br>
>>> together. For your example, the behavior is the same as the code<br>
>>> below:<br>
>>> <br>
>>> ...<br>
>>> #pragma omp parallel num_threads(2)<br>
>>> {<br>
>>> if (omp_get_thread_num() == 0) {<br>
>>> #pragma omp target map(s.x, s.p[0:s.x])<br>
>>> {<br>
>>> for (int i = 0; i < s.x; i++) s.p[i] = i;<br>
>>> }<br>
>>> } else {<br>
>>> #pragma omp target map(other_data)<br>
>>> {<br>
>>> // work on other_data<br>
>>> }<br>
>>> }<br>
>>> ...<br>
>>> <br>
>>> -------------------------<br>
>>> FROM: Deepak Eachempati <<a href="mailto:deachempat@cray.com" target="_blank">deachempat@cray.com</a>><br>
>>> Sent: Thursday, June 27, 2019 3:34 PM<br>
>>> To: Li, Lingda; Narayanaswamy, Ravi; 'Alexandre Eichenberger';<br>
>>> 'Alexey Bataev'; Chapman, Barbara (Contact); Bobrovsky, Konstantin<br>
>>> S; 'Carlo Bertolli'; 'Chan, SiuChi'; Cownie, James H; David<br>
>>> Oehmke; 'Denny, Joel E.'; Dmitriev, Serguei N; Doerfert, Johannes<br>
>>> Rudolf ; 'Ettore Tiotto'; '<a href="mailto:fraggamuffin@gmail.com" target="_blank">fraggamuffin@gmail.com</a>'; 'Gheorghe-Teod<br>
>>> Bercea'; Hal Finkel; '<a href="mailto:jbeyer@nvidia.com" target="_blank">jbeyer@nvidia.com</a>'; 'Jeeva Paudel'; 'Jeff<br>
>>> Heath'; Jeffrey Sandoval; Jones, Jeff C; '<a href="mailto:josem@udel.edu" target="_blank">josem@udel.edu</a>'; 'Kelvin<br>
>>> Li'; 'Kevin K O'Brien'; '<a href="mailto:khaldi.dounia@gmail.com" target="_blank">khaldi.dounia@gmail.com</a>'; 'Kotsifakou,<br>
>>> Maria'; Krishnaiyer, Rakesh; Lieberman, Ron ; 'Lopez, Matthew<br>
>>> Graham'; '<a href="mailto:lopezmg@ornl.org" target="_blank">lopezmg@ornl.org</a>'; 'Martin Kong'; 'Matt Martineau';<br>
>>> Menard, Lorri; Monteleone, Robert; <a href="mailto:oscar@ornl.gov" target="_blank">oscar@ornl.gov</a>; Rao, Premanand<br>
>>> M; Rice, Michael P; Robichaux, Joseph; <a href="mailto:gregory.rodgers@amd.com" target="_blank">gregory.rodgers@amd.com</a>;<br>
>>> Rokos, Georgios; '<a href="mailto:samuel.antao@ibm.com" target="_blank">samuel.antao@ibm.com</a>'; 'Sarah McNamara';<br>
>>> '<a href="mailto:sergey.y.ostanevich@gmail.com" target="_blank">sergey.y.ostanevich@gmail.com</a>'; 'Sergio Pino Gallardo'; 'Sharif,<br>
>>> Hashim'; Sjodin, Jan ; Sunil Shrestha; 'Sunita Chandrasekaran';<br>
>>> Tian, Xinmin; Tianyi Zhang; '<a href="mailto:vadve@illinois.edu" target="_blank">vadve@illinois.edu</a>'; 'Wael Yehia';<br>
>>> 'Wang Chen'; Wilmarth, Terry L<br>
>>> Subject: RE: Comparison of 2 schemes to implement OpenMP 5.0<br>
>>> declare mapper codegen<br>
>>> <br>
>>> I was referring to something like this, where another thread is<br>
>>> not trying to map the same data:<br>
>>> <br>
>>> #pragma omp declare mapper(S s) map(s.x) map(s.p[0:s.x])<br>
>>> S s;<br>
>>> ...<br>
>>> #pragma omp parallel num_threads(2)<br>
>>> {<br>
>>> if (omp_get_thread_num() == 0) {<br>
>>> #pragma omp target map(s)<br>
>>> {<br>
>>> for (int i = 0; i < s.x; i++) s.p[i] = i;<br>
>>> }<br>
>>> } else {<br>
>>> #pragma omp target map(other_data)<br>
>>> {<br>
>>> // work on other_data<br>
>>> }<br>
>>> }<br>
>>> ...<br>
>>> <br>
>>> Since I believe you are mapping s.x and s.p as separate map<br>
>>> operations, it is possible that another thread could map<br>
>>> ‘other_data’ in between those two maps. If this happens, will<br>
>>> your implementation still ensure that s.x and s.p are positioned<br>
>>> at the right offsets with respect to the same base address (&s)?<br>
>>> <br>
>>> -- Deepak<br>
>>> <br>
>>> FROM: Li, Lingda [mailto:<a href="mailto:lli@bnl.gov" target="_blank">lli@bnl.gov</a>]<br>
>>> Sent: Thursday, June 27, 2019 2:26 PM<br>
>>> To: Deepak Eachempati <<a href="mailto:deachempat@cray.com" target="_blank">deachempat@cray.com</a>>; Narayanaswamy, Ravi<br>
>>> <<a href="mailto:ravi.narayanaswamy@intel.com" target="_blank">ravi.narayanaswamy@intel.com</a>>; 'Alexandre Eichenberger'<br>
>>> <<a href="mailto:alexe@us.ibm.com" target="_blank">alexe@us.ibm.com</a>>; 'Alexey Bataev' <<a href="mailto:Alexey.Bataev@ibm.com" target="_blank">Alexey.Bataev@ibm.com</a>>;<br>
>>> Chapman, Barbara (Contact) <<a href="mailto:barbara.chapman@stonybrook.edu" target="_blank">barbara.chapman@stonybrook.edu</a>>;<br>
>>> Bobrovsky, Konstantin S <<a href="mailto:konstantin.s.bobrovsky@intel.com" target="_blank">konstantin.s.bobrovsky@intel.com</a>>; 'Carlo<br>
>>> Bertolli' <<a href="mailto:cbertol@us.ibm.com" target="_blank">cbertol@us.ibm.com</a>>; 'Chan, SiuChi'<br>
>>> <<a href="mailto:siuchi.chan@amd.com" target="_blank">siuchi.chan@amd.com</a>>; Cownie, James H <<a href="mailto:james.h.cownie@intel.com" target="_blank">james.h.cownie@intel.com</a>>;<br>
>>> David Oehmke <<a href="mailto:doehmke@cray.com" target="_blank">doehmke@cray.com</a>>; 'Denny, Joel E.'<br>
>>> <<a href="mailto:dennyje@ornl.gov" target="_blank">dennyje@ornl.gov</a>>; Dmitriev, Serguei N<br>
>>> <<a href="mailto:serguei.n.dmitriev@intel.com" target="_blank">serguei.n.dmitriev@intel.com</a>>; Doerfert, Johannes Rudolf<br>
>>> <<a href="mailto:jdoerfert@anl.gov" target="_blank">jdoerfert@anl.gov</a>>; 'Ettore Tiotto' <<a href="mailto:etiotto@ca.ibm.com" target="_blank">etiotto@ca.ibm.com</a>>;<br>
>>> '<a href="mailto:fraggamuffin@gmail.com" target="_blank">fraggamuffin@gmail.com</a>' <<a href="mailto:fraggamuffin@gmail.com" target="_blank">fraggamuffin@gmail.com</a>>; 'Gheorghe-Teod<br>
>>> Bercea' <<a href="mailto:Gheorghe-Teod.Bercea@ibm.com" target="_blank">Gheorghe-Teod.Bercea@ibm.com</a>>; Hal Finkel<br>
>>> <<a href="mailto:hfinkel@anl.gov" target="_blank">hfinkel@anl.gov</a>>; '<a href="mailto:jbeyer@nvidia.com" target="_blank">jbeyer@nvidia.com</a>' <<a href="mailto:jbeyer@nvidia.com" target="_blank">jbeyer@nvidia.com</a>>; 'Jeeva<br>
>>> Paudel' <<a href="mailto:pjeeva01@ca.ibm.com" target="_blank">pjeeva01@ca.ibm.com</a>>; 'Jeff Heath' <<a href="mailto:jrheath@ca.ibm.com" target="_blank">jrheath@ca.ibm.com</a>>;<br>
>>> Jeffrey Sandoval <<a href="mailto:sandoval@cray.com" target="_blank">sandoval@cray.com</a>>; Jones, Jeff C<br>
>>> <<a href="mailto:jeff.c.jones@intel.com" target="_blank">jeff.c.jones@intel.com</a>>; '<a href="mailto:josem@udel.edu" target="_blank">josem@udel.edu</a>' <<a href="mailto:josem@udel.edu" target="_blank">josem@udel.edu</a>>;<br>
>>> 'Kelvin Li' <<a href="mailto:kli@ca.ibm.com" target="_blank">kli@ca.ibm.com</a>>; 'Kevin K O'Brien'<br>
>>> <<a href="mailto:caomhin@us.ibm.com" target="_blank">caomhin@us.ibm.com</a>>; '<a href="mailto:khaldi.dounia@gmail.com" target="_blank">khaldi.dounia@gmail.com</a>'<br>
>>> <<a href="mailto:khaldi.dounia@gmail.com" target="_blank">khaldi.dounia@gmail.com</a>>; 'Kotsifakou, Maria'<br>
>>> <<a href="mailto:kotsifa2@illinois.edu" target="_blank">kotsifa2@illinois.edu</a>>; Krishnaiyer, Rakesh<br>
>>> <<a href="mailto:rakesh.krishnaiyer@intel.com" target="_blank">rakesh.krishnaiyer@intel.com</a>>; Lieberman, Ron<br>
>>> <<a href="mailto:Ron.Lieberman@amd.com" target="_blank">Ron.Lieberman@amd.com</a>>; 'Lopez, Matthew Graham'<br>
>>> <<a href="mailto:lopezmg@ornl.gov" target="_blank">lopezmg@ornl.gov</a>>; '<a href="mailto:lopezmg@ornl.org" target="_blank">lopezmg@ornl.org</a>' <<a href="mailto:lopezmg@ornl.org" target="_blank">lopezmg@ornl.org</a>>; 'Martin<br>
>>> Kong' <<a href="mailto:martin.richard.kong@gmail.com" target="_blank">martin.richard.kong@gmail.com</a>>; 'Matt Martineau'<br>
>>> <<a href="mailto:m.martineau@bristol.ac.uk" target="_blank">m.martineau@bristol.ac.uk</a>>; Menard, Lorri<br>
>>> <<a href="mailto:lorri.menard@intel.com" target="_blank">lorri.menard@intel.com</a>>; Monteleone, Robert<br>
>>> <<a href="mailto:robert.monteleone@intel.com" target="_blank">robert.monteleone@intel.com</a>>; <a href="mailto:oscar@ornl.gov" target="_blank">oscar@ornl.gov</a>; Rao, Premanand M<br>
>>> <<a href="mailto:premanand.m.rao@intel.com" target="_blank">premanand.m.rao@intel.com</a>>; Rice, Michael P<br>
>>> <<a href="mailto:michael.p.rice@intel.com" target="_blank">michael.p.rice@intel.com</a>>; Robichaux, Joseph<br>
>>> <<a href="mailto:joseph.robichaux@intel.com" target="_blank">joseph.robichaux@intel.com</a>>; <a href="mailto:gregory.rodgers@amd.com" target="_blank">gregory.rodgers@amd.com</a>; Rokos,<br>
>>> Georgios <<a href="mailto:georgios.rokos@intel.com" target="_blank">georgios.rokos@intel.com</a>>; '<a href="mailto:samuel.antao@ibm.com" target="_blank">samuel.antao@ibm.com</a>'<br>
>>> <<a href="mailto:samuel.antao@ibm.com" target="_blank">samuel.antao@ibm.com</a>>; 'Sarah McNamara' <<a href="mailto:mcnamara@ca.ibm.com" target="_blank">mcnamara@ca.ibm.com</a>>;<br>
>>> '<a href="mailto:sergey.y.ostanevich@gmail.com" target="_blank">sergey.y.ostanevich@gmail.com</a>' <<a href="mailto:sergey.y.ostanevich@gmail.com" target="_blank">sergey.y.ostanevich@gmail.com</a>>;<br>
>>> 'Sergio Pino Gallardo' <<a href="mailto:sergiop@udel.edu" target="_blank">sergiop@udel.edu</a>>; 'Sharif, Hashim'<br>
>>> <<a href="mailto:hsharif3@illinois.edu" target="_blank">hsharif3@illinois.edu</a>>; Sjodin, Jan <<a href="mailto:Jan.Sjodin@amd.com" target="_blank">Jan.Sjodin@amd.com</a>>; Sunil<br>
>>> Shrestha <<a href="mailto:sshrestha@cray.com" target="_blank">sshrestha@cray.com</a>>; 'Sunita Chandrasekaran'<br>
>>> <<a href="mailto:schandra@udel.edu" target="_blank">schandra@udel.edu</a>>; Tian, Xinmin <<a href="mailto:xinmin.tian@intel.com" target="_blank">xinmin.tian@intel.com</a>>; Tianyi<br>
>>> Zhang <<a href="mailto:tzhan18@lsu.edu" target="_blank">tzhan18@lsu.edu</a>>; '<a href="mailto:vadve@illinois.edu" target="_blank">vadve@illinois.edu</a>'<br>
>>> <<a href="mailto:vadve@illinois.edu" target="_blank">vadve@illinois.edu</a>>; 'Wael Yehia' <<a href="mailto:wyehia@ca.ibm.com" target="_blank">wyehia@ca.ibm.com</a>>; 'Wang<br>
>>> Chen' <<a href="mailto:wdchen@ca.ibm.com" target="_blank">wdchen@ca.ibm.com</a>>; Wilmarth, Terry L<br>
>>> <<a href="mailto:terry.l.wilmarth@intel.com" target="_blank">terry.l.wilmarth@intel.com</a>><br>
>>> Subject: Re: Comparison of 2 schemes to implement OpenMP 5.0<br>
>>> declare mapper codegen<br>
>>> <br>
>>> When 2 threads try to concurrently map the same data, it behaves<br>
>>> the same as when 2 threads concurrently map the same data using<br>
>>> map clauses, and mappers don't introduce extra considerations<br>
>>> here. For instance, both threads use #omp target enter data<br>
>>> concurrently.<br>
>>> <br>
>>> When 2 threads concurrently maps the same data, my understanding<br>
>>> based on the current code is, it will create 2 copies of the same<br>
>>> data, either copy is correctly to use. It may have a problem when<br>
>>> both copies are mapped back if not synchronized correctly, but<br>
>>> this is a programming issue, not the responsibility of OpenMP.<br>
>>> <br>
>>> Thanks,<br>
>>> Lingda Li<br>
>>> <br>
>>> -------------------------<br>
>>> FROM: Deepak Eachempati <<a href="mailto:deachempat@cray.com" target="_blank">deachempat@cray.com</a>><br>
>>> Sent: Thursday, June 27, 2019 3:17 PM<br>
>>> To: Li, Lingda; Narayanaswamy, Ravi; 'Alexandre Eichenberger';<br>
>>> 'Alexey Bataev'; Chapman, Barbara (Contact); Bobrovsky, Konstantin<br>
>>> S; 'Carlo Bertolli'; 'Chan, SiuChi'; Cownie, James H; David<br>
>>> Oehmke; 'Denny, Joel E.'; Dmitriev, Serguei N; Doerfert, Johannes<br>
>>> Rudolf ; 'Ettore Tiotto'; '<a href="mailto:fraggamuffin@gmail.com" target="_blank">fraggamuffin@gmail.com</a>'; 'Gheorghe-Teod<br>
>>> Bercea'; Hal Finkel; '<a href="mailto:jbeyer@nvidia.com" target="_blank">jbeyer@nvidia.com</a>'; 'Jeeva Paudel'; 'Jeff<br>
>>> Heath'; Jeffrey Sandoval; Jones, Jeff C; '<a href="mailto:josem@udel.edu" target="_blank">josem@udel.edu</a>'; 'Kelvin<br>
>>> Li'; 'Kevin K O'Brien'; '<a href="mailto:khaldi.dounia@gmail.com" target="_blank">khaldi.dounia@gmail.com</a>'; 'Kotsifakou,<br>
>>> Maria'; Krishnaiyer, Rakesh; Lieberman, Ron ; 'Lopez, Matthew<br>
>>> Graham'; '<a href="mailto:lopezmg@ornl.org" target="_blank">lopezmg@ornl.org</a>'; 'Martin Kong'; 'Matt Martineau';<br>
>>> Menard, Lorri; Monteleone, Robert; <a href="mailto:oscar@ornl.gov" target="_blank">oscar@ornl.gov</a>; Rao, Premanand<br>
>>> M; Rice, Michael P; Robichaux, Joseph; <a href="mailto:gregory.rodgers@amd.com" target="_blank">gregory.rodgers@amd.com</a>;<br>
>>> Rokos, Georgios; '<a href="mailto:samuel.antao@ibm.com" target="_blank">samuel.antao@ibm.com</a>'; 'Sarah McNamara';<br>
>>> '<a href="mailto:sergey.y.ostanevich@gmail.com" target="_blank">sergey.y.ostanevich@gmail.com</a>'; 'Sergio Pino Gallardo'; 'Sharif,<br>
>>> Hashim'; Sjodin, Jan ; Sunil Shrestha; 'Sunita Chandrasekaran';<br>
>>> Tian, Xinmin; Tianyi Zhang; '<a href="mailto:vadve@illinois.edu" target="_blank">vadve@illinois.edu</a>'; 'Wael Yehia';<br>
>>> 'Wang Chen'; Wilmarth, Terry L<br>
>>> Subject: RE: Comparison of 2 schemes to implement OpenMP 5.0<br>
>>> declare mapper codegen<br>
>>> <br>
>>> Thanks.<br>
>>> <br>
>>> Is it possible for another thread to be concurrently mapped<br>
>>> something else while the maps from the mapper function are taking<br>
>>> place? If so, how do you guarantee that the allocation for each<br>
>>> component will get you the right addresses in device memory? Sorry<br>
>>> if this was covered before and I missed it.<br>
>>> <br>
>>> -- Deepak<br>
>>> <br>
>>> FROM: Li, Lingda [mailto:<a href="mailto:lli@bnl.gov" target="_blank">lli@bnl.gov</a>]<br>
>>> Sent: Thursday, June 27, 2019 2:08 PM<br>
>>> To: Deepak Eachempati <<a href="mailto:deachempat@cray.com" target="_blank">deachempat@cray.com</a>>; Narayanaswamy, Ravi<br>
>>> <<a href="mailto:ravi.narayanaswamy@intel.com" target="_blank">ravi.narayanaswamy@intel.com</a>>; 'Alexandre Eichenberger'<br>
>>> <<a href="mailto:alexe@us.ibm.com" target="_blank">alexe@us.ibm.com</a>>; 'Alexey Bataev' <<a href="mailto:Alexey.Bataev@ibm.com" target="_blank">Alexey.Bataev@ibm.com</a>>;<br>
>>> Chapman, Barbara (Contact) <<a href="mailto:barbara.chapman@stonybrook.edu" target="_blank">barbara.chapman@stonybrook.edu</a>>;<br>
>>> Bobrovsky, Konstantin S <<a href="mailto:konstantin.s.bobrovsky@intel.com" target="_blank">konstantin.s.bobrovsky@intel.com</a>>; 'Carlo<br>
>>> Bertolli' <<a href="mailto:cbertol@us.ibm.com" target="_blank">cbertol@us.ibm.com</a>>; 'Chan, SiuChi'<br>
>>> <<a href="mailto:siuchi.chan@amd.com" target="_blank">siuchi.chan@amd.com</a>>; Cownie, James H <<a href="mailto:james.h.cownie@intel.com" target="_blank">james.h.cownie@intel.com</a>>;<br>
>>> David Oehmke <<a href="mailto:doehmke@cray.com" target="_blank">doehmke@cray.com</a>>; 'Denny, Joel E.'<br>
>>> <<a href="mailto:dennyje@ornl.gov" target="_blank">dennyje@ornl.gov</a>>; Dmitriev, Serguei N<br>
>>> <<a href="mailto:serguei.n.dmitriev@intel.com" target="_blank">serguei.n.dmitriev@intel.com</a>>; Doerfert, Johannes Rudolf<br>
>>> <<a href="mailto:jdoerfert@anl.gov" target="_blank">jdoerfert@anl.gov</a>>; 'Ettore Tiotto' <<a href="mailto:etiotto@ca.ibm.com" target="_blank">etiotto@ca.ibm.com</a>>;<br>
>>> '<a href="mailto:fraggamuffin@gmail.com" target="_blank">fraggamuffin@gmail.com</a>' <<a href="mailto:fraggamuffin@gmail.com" target="_blank">fraggamuffin@gmail.com</a>>; 'Gheorghe-Teod<br>
>>> Bercea' <<a href="mailto:Gheorghe-Teod.Bercea@ibm.com" target="_blank">Gheorghe-Teod.Bercea@ibm.com</a>>; Hal Finkel<br>
>>> <<a href="mailto:hfinkel@anl.gov" target="_blank">hfinkel@anl.gov</a>>; '<a href="mailto:jbeyer@nvidia.com" target="_blank">jbeyer@nvidia.com</a>' <<a href="mailto:jbeyer@nvidia.com" target="_blank">jbeyer@nvidia.com</a>>; 'Jeeva<br>
>>> Paudel' <<a href="mailto:pjeeva01@ca.ibm.com" target="_blank">pjeeva01@ca.ibm.com</a>>; 'Jeff Heath' <<a href="mailto:jrheath@ca.ibm.com" target="_blank">jrheath@ca.ibm.com</a>>;<br>
>>> Jeffrey Sandoval <<a href="mailto:sandoval@cray.com" target="_blank">sandoval@cray.com</a>>; Jones, Jeff C<br>
>>> <<a href="mailto:jeff.c.jones@intel.com" target="_blank">jeff.c.jones@intel.com</a>>; '<a href="mailto:josem@udel.edu" target="_blank">josem@udel.edu</a>' <<a href="mailto:josem@udel.edu" target="_blank">josem@udel.edu</a>>;<br>
>>> 'Kelvin Li' <<a href="mailto:kli@ca.ibm.com" target="_blank">kli@ca.ibm.com</a>>; 'Kevin K O'Brien'<br>
>>> <<a href="mailto:caomhin@us.ibm.com" target="_blank">caomhin@us.ibm.com</a>>; '<a href="mailto:khaldi.dounia@gmail.com" target="_blank">khaldi.dounia@gmail.com</a>'<br>
>>> <<a href="mailto:khaldi.dounia@gmail.com" target="_blank">khaldi.dounia@gmail.com</a>>; 'Kotsifakou, Maria'<br>
>>> <<a href="mailto:kotsifa2@illinois.edu" target="_blank">kotsifa2@illinois.edu</a>>; Krishnaiyer, Rakesh<br>
>>> <<a href="mailto:rakesh.krishnaiyer@intel.com" target="_blank">rakesh.krishnaiyer@intel.com</a>>; Lieberman, Ron<br>
>>> <<a href="mailto:Ron.Lieberman@amd.com" target="_blank">Ron.Lieberman@amd.com</a>>; 'Lopez, Matthew Graham'<br>
>>> <<a href="mailto:lopezmg@ornl.gov" target="_blank">lopezmg@ornl.gov</a>>; '<a href="mailto:lopezmg@ornl.org" target="_blank">lopezmg@ornl.org</a>' <<a href="mailto:lopezmg@ornl.org" target="_blank">lopezmg@ornl.org</a>>; 'Martin<br>
>>> Kong' <<a href="mailto:martin.richard.kong@gmail.com" target="_blank">martin.richard.kong@gmail.com</a>>; 'Matt Martineau'<br>
>>> <<a href="mailto:m.martineau@bristol.ac.uk" target="_blank">m.martineau@bristol.ac.uk</a>>; Menard, Lorri<br>
>>> <<a href="mailto:lorri.menard@intel.com" target="_blank">lorri.menard@intel.com</a>>; Monteleone, Robert<br>
>>> <<a href="mailto:robert.monteleone@intel.com" target="_blank">robert.monteleone@intel.com</a>>; <a href="mailto:oscar@ornl.gov" target="_blank">oscar@ornl.gov</a>; Rao, Premanand M<br>
>>> <<a href="mailto:premanand.m.rao@intel.com" target="_blank">premanand.m.rao@intel.com</a>>; Rice, Michael P<br>
>>> <<a href="mailto:michael.p.rice@intel.com" target="_blank">michael.p.rice@intel.com</a>>; Robichaux, Joseph<br>
>>> <<a href="mailto:joseph.robichaux@intel.com" target="_blank">joseph.robichaux@intel.com</a>>; <a href="mailto:gregory.rodgers@amd.com" target="_blank">gregory.rodgers@amd.com</a>; Rokos,<br>
>>> Georgios <<a href="mailto:georgios.rokos@intel.com" target="_blank">georgios.rokos@intel.com</a>>; '<a href="mailto:samuel.antao@ibm.com" target="_blank">samuel.antao@ibm.com</a>'<br>
>>> <<a href="mailto:samuel.antao@ibm.com" target="_blank">samuel.antao@ibm.com</a>>; 'Sarah McNamara' <<a href="mailto:mcnamara@ca.ibm.com" target="_blank">mcnamara@ca.ibm.com</a>>;<br>
>>> '<a href="mailto:sergey.y.ostanevich@gmail.com" target="_blank">sergey.y.ostanevich@gmail.com</a>' <<a href="mailto:sergey.y.ostanevich@gmail.com" target="_blank">sergey.y.ostanevich@gmail.com</a>>;<br>
>>> 'Sergio Pino Gallardo' <<a href="mailto:sergiop@udel.edu" target="_blank">sergiop@udel.edu</a>>; 'Sharif, Hashim'<br>
>>> <<a href="mailto:hsharif3@illinois.edu" target="_blank">hsharif3@illinois.edu</a>>; Sjodin, Jan <<a href="mailto:Jan.Sjodin@amd.com" target="_blank">Jan.Sjodin@amd.com</a>>; Sunil<br>
>>> Shrestha <<a href="mailto:sshrestha@cray.com" target="_blank">sshrestha@cray.com</a>>; 'Sunita Chandrasekaran'<br>
>>> <<a href="mailto:schandra@udel.edu" target="_blank">schandra@udel.edu</a>>; Tian, Xinmin <<a href="mailto:xinmin.tian@intel.com" target="_blank">xinmin.tian@intel.com</a>>; Tianyi<br>
>>> Zhang <<a href="mailto:tzhan18@lsu.edu" target="_blank">tzhan18@lsu.edu</a>>; '<a href="mailto:vadve@illinois.edu" target="_blank">vadve@illinois.edu</a>'<br>
>>> <<a href="mailto:vadve@illinois.edu" target="_blank">vadve@illinois.edu</a>>; 'Wael Yehia' <<a href="mailto:wyehia@ca.ibm.com" target="_blank">wyehia@ca.ibm.com</a>>; 'Wang<br>
>>> Chen' <<a href="mailto:wdchen@ca.ibm.com" target="_blank">wdchen@ca.ibm.com</a>>; Wilmarth, Terry L<br>
>>> <<a href="mailto:terry.l.wilmarth@intel.com" target="_blank">terry.l.wilmarth@intel.com</a>><br>
>>> Subject: Re: Comparison of 2 schemes to implement OpenMP 5.0<br>
>>> declare mapper codegen<br>
>>> <br>
>>> Hi Deepak,<br>
>>> <br>
>>> Yes, it handles this case. The first part of mapper function<br>
>>> (initially allocate space for the whole array) is just an<br>
>>> optimization, not required for correctness, as suggested by you in<br>
>>> an early discussion.<br>
>>> <br>
>>> In your example, s.x and s.p will be allocated separately (not in<br>
>>> a single allocation). But Clang guarantees that their addresses<br>
>>> will be correct because s.x and s.p share the same base address,<br>
>>> which is &s.<br>
>>> <br>
>>> Thanks,<br>
>>> Lingda Li<br>
>>> <br>
>>> -------------------------<br>
>>> FROM: Deepak Eachempati <<a href="mailto:deachempat@cray.com" target="_blank">deachempat@cray.com</a>><br>
>>> Sent: Thursday, June 27, 2019 2:49 PM<br>
>>> To: Li, Lingda; Narayanaswamy, Ravi; 'Alexandre Eichenberger';<br>
>>> 'Alexey Bataev'; Chapman, Barbara (Contact); Bobrovsky, Konstantin<br>
>>> S; 'Carlo Bertolli'; 'Chan, SiuChi'; Cownie, James H; David<br>
>>> Oehmke; 'Denny, Joel E.'; Dmitriev, Serguei N; Doerfert, Johannes<br>
>>> Rudolf ; '<a href="mailto:estotzer@ti.com" target="_blank">estotzer@ti.com</a>'; 'Ettore Tiotto';<br>
>>> '<a href="mailto:fraggamuffin@gmail.com" target="_blank">fraggamuffin@gmail.com</a>'; 'Gheorghe-Teod Bercea'; Hal Finkel;<br>
>>> '<a href="mailto:jbeyer@nvidia.com" target="_blank">jbeyer@nvidia.com</a>'; 'Jeeva Paudel'; 'Jeff Heath'; Jeffrey<br>
>>> Sandoval; Jones, Jeff C; '<a href="mailto:josem@udel.edu" target="_blank">josem@udel.edu</a>'; 'Kelvin Li'; 'Kevin K<br>
>>> O'Brien'; '<a href="mailto:khaldi.dounia@gmail.com" target="_blank">khaldi.dounia@gmail.com</a>'; 'Kotsifakou, Maria';<br>
>>> Krishnaiyer, Rakesh; Lieberman, Ron ; 'Lopez, Matthew Graham';<br>
>>> '<a href="mailto:lopezmg@ornl.org" target="_blank">lopezmg@ornl.org</a>'; 'Martin Kong'; 'Matt Martineau'; Menard,<br>
>>> Lorri; Monteleone, Robert; <a href="mailto:oscar@ornl.gov" target="_blank">oscar@ornl.gov</a>; Rao, Premanand M; Rice,<br>
>>> Michael P; Robichaux, Joseph; <a href="mailto:gregory.rodgers@amd.com" target="_blank">gregory.rodgers@amd.com</a>; Rokos,<br>
>>> Georgios; '<a href="mailto:samuel.antao@ibm.com" target="_blank">samuel.antao@ibm.com</a>'; 'Sarah McNamara';<br>
>>> '<a href="mailto:sergey.y.ostanevich@gmail.com" target="_blank">sergey.y.ostanevich@gmail.com</a>'; 'Sergio Pino Gallardo'; 'Sharif,<br>
>>> Hashim'; Sjodin, Jan ; Sunil Shrestha; 'Sunita Chandrasekaran';<br>
>>> Tian, Xinmin; Tianyi Zhang; '<a href="mailto:vadve@illinois.edu" target="_blank">vadve@illinois.edu</a>'; 'Wael Yehia';<br>
>>> 'Wang Chen'; Wilmarth, Terry L<br>
>>> Subject: RE: Comparison of 2 schemes to implement OpenMP 5.0<br>
>>> declare mapper codegen<br>
>>> <br>
>>> For Scheme 1, it looks like you are doing separate maps for each<br>
>>> component when size == 1. It seems like the first and last if<br>
>>> statements should have “size >= 1” rather than “size > 1”.<br>
>>> <br>
>>> If the mapper is declared like this:<br>
>>> <br>
>>> struct S {<br>
>>> int x;<br>
>>> ... // other stuff<br>
>>> int *p;<br>
>>> };<br>
>>> <br>
>>> #pragma omp declare mapper(S s) map(s.x) map(s.p[0:s.x])<br>
>>> <br>
>>> And you have:<br>
>>> <br>
>>> S s;<br>
>>> ...<br>
>>> #pragma omp target map(s)<br>
>>> {<br>
>>> for (int i = 0; i < s.x; i++) s.p[i] = i;<br>
>>> }<br>
>>> <br>
>>> Since the target construct is just mapping a single structure of<br>
>>> type S, there should be one map that takes care of mapping storage<br>
>>> for s.x and s.p with a single allocation, and a separate map for<br>
>>> the array section s.p[0:s.x], and finally the pointer attachment<br>
>>> of s.p to s.p[0:s.x]. Does Scheme 1 handle this?<br>
>>> <br>
>>> -- Deepak<br>
>>> <br>
>>> FROM: Li, Lingda [mailto:<a href="mailto:lli@bnl.gov" target="_blank">lli@bnl.gov</a>]<br>
>>> Sent: Thursday, June 27, 2019 1:07 PM<br>
>>> To: Narayanaswamy, Ravi <<a href="mailto:ravi.narayanaswamy@intel.com" target="_blank">ravi.narayanaswamy@intel.com</a>>; 'Alexandre<br>
>>> Eichenberger' <<a href="mailto:alexe@us.ibm.com" target="_blank">alexe@us.ibm.com</a>>; 'Alexey Bataev'<br>
>>> <<a href="mailto:Alexey.Bataev@ibm.com" target="_blank">Alexey.Bataev@ibm.com</a>>; Chapman, Barbara (Contact)<br>
>>> <<a href="mailto:barbara.chapman@stonybrook.edu" target="_blank">barbara.chapman@stonybrook.edu</a>>; Bobrovsky, Konstantin S<br>
>>> <<a href="mailto:konstantin.s.bobrovsky@intel.com" target="_blank">konstantin.s.bobrovsky@intel.com</a>>; 'Carlo Bertolli'<br>
>>> <<a href="mailto:cbertol@us.ibm.com" target="_blank">cbertol@us.ibm.com</a>>; 'Chan, SiuChi' <<a href="mailto:siuchi.chan@amd.com" target="_blank">siuchi.chan@amd.com</a>>;<br>
>>> Cownie, James H <<a href="mailto:james.h.cownie@intel.com" target="_blank">james.h.cownie@intel.com</a>>; David Oehmke<br>
>>> <<a href="mailto:doehmke@cray.com" target="_blank">doehmke@cray.com</a>>; Deepak Eachempati <<a href="mailto:deachempat@cray.com" target="_blank">deachempat@cray.com</a>>;<br>
>>> 'Denny, Joel E.' <<a href="mailto:dennyje@ornl.gov" target="_blank">dennyje@ornl.gov</a>>; Dmitriev, Serguei N<br>
>>> <<a href="mailto:serguei.n.dmitriev@intel.com" target="_blank">serguei.n.dmitriev@intel.com</a>>; Doerfert, Johannes Rudolf<br>
>>> <<a href="mailto:jdoerfert@anl.gov" target="_blank">jdoerfert@anl.gov</a>>; '<a href="mailto:estotzer@ti.com" target="_blank">estotzer@ti.com</a>' <<a href="mailto:estotzer@ti.com" target="_blank">estotzer@ti.com</a>>; 'Ettore<br>
>>> Tiotto' <<a href="mailto:etiotto@ca.ibm.com" target="_blank">etiotto@ca.ibm.com</a>>; '<a href="mailto:fraggamuffin@gmail.com" target="_blank">fraggamuffin@gmail.com</a>'<br>
>>> <<a href="mailto:fraggamuffin@gmail.com" target="_blank">fraggamuffin@gmail.com</a>>; 'Gheorghe-Teod Bercea'<br>
>>> <<a href="mailto:Gheorghe-Teod.Bercea@ibm.com" target="_blank">Gheorghe-Teod.Bercea@ibm.com</a>>; Hal Finkel <<a href="mailto:hfinkel@anl.gov" target="_blank">hfinkel@anl.gov</a>>;<br>
>>> '<a href="mailto:jbeyer@nvidia.com" target="_blank">jbeyer@nvidia.com</a>' <<a href="mailto:jbeyer@nvidia.com" target="_blank">jbeyer@nvidia.com</a>>; 'Jeeva Paudel'<br>
>>> <<a href="mailto:pjeeva01@ca.ibm.com" target="_blank">pjeeva01@ca.ibm.com</a>>; 'Jeff Heath' <<a href="mailto:jrheath@ca.ibm.com" target="_blank">jrheath@ca.ibm.com</a>>; Jeffrey<br>
>>> Sandoval <<a href="mailto:sandoval@cray.com" target="_blank">sandoval@cray.com</a>>; Jones, Jeff C<br>
>>> <<a href="mailto:jeff.c.jones@intel.com" target="_blank">jeff.c.jones@intel.com</a>>; '<a href="mailto:josem@udel.edu" target="_blank">josem@udel.edu</a>' <<a href="mailto:josem@udel.edu" target="_blank">josem@udel.edu</a>>;<br>
>>> 'Kelvin Li' <<a href="mailto:kli@ca.ibm.com" target="_blank">kli@ca.ibm.com</a>>; 'Kevin K O'Brien'<br>
>>> <<a href="mailto:caomhin@us.ibm.com" target="_blank">caomhin@us.ibm.com</a>>; '<a href="mailto:khaldi.dounia@gmail.com" target="_blank">khaldi.dounia@gmail.com</a>'<br>
>>> <<a href="mailto:khaldi.dounia@gmail.com" target="_blank">khaldi.dounia@gmail.com</a>>; 'Kotsifakou, Maria'<br>
>>> <<a href="mailto:kotsifa2@illinois.edu" target="_blank">kotsifa2@illinois.edu</a>>; Krishnaiyer, Rakesh<br>
>>> <<a href="mailto:rakesh.krishnaiyer@intel.com" target="_blank">rakesh.krishnaiyer@intel.com</a>>; Lieberman, Ron<br>
>>> <<a href="mailto:Ron.Lieberman@amd.com" target="_blank">Ron.Lieberman@amd.com</a>>; Li, Lingda <<a href="mailto:lli@bnl.gov" target="_blank">lli@bnl.gov</a>>; 'Lopez, Matthew<br>
>>> Graham' <<a href="mailto:lopezmg@ornl.gov" target="_blank">lopezmg@ornl.gov</a>>; '<a href="mailto:lopezmg@ornl.org" target="_blank">lopezmg@ornl.org</a>' <<a href="mailto:lopezmg@ornl.org" target="_blank">lopezmg@ornl.org</a>>;<br>
>>> 'Martin Kong' <<a href="mailto:martin.richard.kong@gmail.com" target="_blank">martin.richard.kong@gmail.com</a>>; 'Matt Martineau'<br>
>>> <<a href="mailto:m.martineau@bristol.ac.uk" target="_blank">m.martineau@bristol.ac.uk</a>>; Menard, Lorri<br>
>>> <<a href="mailto:lorri.menard@intel.com" target="_blank">lorri.menard@intel.com</a>>; Monteleone, Robert<br>
>>> <<a href="mailto:robert.monteleone@intel.com" target="_blank">robert.monteleone@intel.com</a>>; <a href="mailto:oscar@ornl.gov" target="_blank">oscar@ornl.gov</a>; Rao, Premanand M<br>
>>> <<a href="mailto:premanand.m.rao@intel.com" target="_blank">premanand.m.rao@intel.com</a>>; Rice, Michael P<br>
>>> <<a href="mailto:michael.p.rice@intel.com" target="_blank">michael.p.rice@intel.com</a>>; Robichaux, Joseph<br>
>>> <<a href="mailto:joseph.robichaux@intel.com" target="_blank">joseph.robichaux@intel.com</a>>; <a href="mailto:gregory.rodgers@amd.com" target="_blank">gregory.rodgers@amd.com</a>; Rokos,<br>
>>> Georgios <<a href="mailto:georgios.rokos@intel.com" target="_blank">georgios.rokos@intel.com</a>>; '<a href="mailto:samuel.antao@ibm.com" target="_blank">samuel.antao@ibm.com</a>'<br>
>>> <<a href="mailto:samuel.antao@ibm.com" target="_blank">samuel.antao@ibm.com</a>>; 'Sarah McNamara' <<a href="mailto:mcnamara@ca.ibm.com" target="_blank">mcnamara@ca.ibm.com</a>>;<br>
>>> '<a href="mailto:sergey.y.ostanevich@gmail.com" target="_blank">sergey.y.ostanevich@gmail.com</a>' <<a href="mailto:sergey.y.ostanevich@gmail.com" target="_blank">sergey.y.ostanevich@gmail.com</a>>;<br>
>>> 'Sergio Pino Gallardo' <<a href="mailto:sergiop@udel.edu" target="_blank">sergiop@udel.edu</a>>; 'Sharif, Hashim'<br>
>>> <<a href="mailto:hsharif3@illinois.edu" target="_blank">hsharif3@illinois.edu</a>>; Sjodin, Jan <<a href="mailto:Jan.Sjodin@amd.com" target="_blank">Jan.Sjodin@amd.com</a>>; Sunil<br>
>>> Shrestha <<a href="mailto:sshrestha@cray.com" target="_blank">sshrestha@cray.com</a>>; 'Sunita Chandrasekaran'<br>
>>> <<a href="mailto:schandra@udel.edu" target="_blank">schandra@udel.edu</a>>; Tian, Xinmin <<a href="mailto:xinmin.tian@intel.com" target="_blank">xinmin.tian@intel.com</a>>; Tianyi<br>
>>> Zhang <<a href="mailto:tzhan18@lsu.edu" target="_blank">tzhan18@lsu.edu</a>>; '<a href="mailto:vadve@illinois.edu" target="_blank">vadve@illinois.edu</a>'<br>
>>> <<a href="mailto:vadve@illinois.edu" target="_blank">vadve@illinois.edu</a>>; 'Wael Yehia' <<a href="mailto:wyehia@ca.ibm.com" target="_blank">wyehia@ca.ibm.com</a>>; 'Wang<br>
>>> Chen' <<a href="mailto:wdchen@ca.ibm.com" target="_blank">wdchen@ca.ibm.com</a>>; Wilmarth, Terry L<br>
>>> <<a href="mailto:terry.l.wilmarth@intel.com" target="_blank">terry.l.wilmarth@intel.com</a>><br>
>>> Subject: Comparison of 2 schemes to implement OpenMP 5.0 declare<br>
>>> mapper codegen<br>
>>> <br>
>>> Hi,<br>
>>> <br>
>>> Alexey and I would like to have your attention on an ongoing<br>
>>> discussion of 2 schemes to implement the declare mapper in OpenMP<br>
>>> 5.0. The detailed discussion can be found at<br>
>>> <a href="https://reviews.llvm.org/D59474" rel="noreferrer" target="_blank">https://reviews.llvm.org/D59474</a> [1]<br>
>>> <br>
>>> Scheme 1 (the one has been implemented by me in<br>
>>> <a href="https://reviews.llvm.org/D59474" rel="noreferrer" target="_blank">https://reviews.llvm.org/D59474</a> [1]):<br>
>>> The detailed design can be found at<br>
>>> <br>
>> <br>
> <a href="https://github.com/lingda-li/public-sharing/blob/master/mapper_runtime_design.pptx" rel="noreferrer" target="_blank">https://github.com/lingda-li/public-sharing/blob/master/mapper_runtime_design.pptx</a><br>
>>> [2]<br>
>>> For each mapper function, the compiler generates a function like<br>
>>> this:<br>
>>> <br>
>>> ```<br>
>>> void <type>.mapper(void *base, void *begin, size_t size, int64_t<br>
>>> type) {<br>
>>> // Allocate space for an array section first.<br>
>>> if (size > 1 && !maptype.IsDelete)<br>
>>> <push>(base, begin, size*sizeof(Ty), clearToFrom(type));<br>
>>> <br>
>>> // Map members.<br>
>>> for (unsigned i = 0; i < size; i++) {<br>
>>> // For each component specified by this mapper:<br>
>>> for (auto c : components) {<br>
>>> ...; // code to generate c.arg_base, c.arg_begin, c.arg_size,<br>
>>> c.arg_type<br>
>>> if (c.hasMapper())<br>
>>> (*c.Mapper())(c.arg_base, c.arg_begin, c.arg_size, c.arg_type);<br>
>>> else<br>
>>> <push>(c.arg_base, c.arg_begin, c.arg_size, c.arg_type);<br>
>>> }<br>
>>> }<br>
>>> // Delete the array section.<br>
>>> if (size > 1 && maptype.IsDelete)<br>
>>> <push>(base, begin, size*sizeof(Ty), clearToFrom(type));<br>
>>> }<br>
>>> ```<br>
>>> This function is passed to the OpenMP runtime, and the runtime<br>
>>> will call this function to finish the data mapping.<br>
>>> <br>
>>> Scheme 2 (which Alexey proposes):<br>
>>> Alexey proposes to move parts of the mapper function above into<br>
>>> the OpenMP runtime, so the compiler will generate code below:<br>
>>> ```<br>
>>> void <type>.mapper(void *base, void *begin, size_t size, int64_t<br>
>>> type) {<br>
>>> ...; // code to generate arg_base, arg_begin, arg_size, arg_type,<br>
>>> arg_mapper.<br>
>>> auto sub_components[] = {...}; // fill in generated begin, base,<br>
>>> ...<br>
>>> __tgt_mapper(base, begin, size, type, sub_components);<br>
>>> }<br>
>>> ```<br>
>>> <br>
>>> `__tgt_mapper` is a runtime function as below:<br>
>>> ```<br>
>>> void __tgt_mapper(void *base, void *begin, size_t size, int64_t<br>
>>> type, auto components[]) {<br>
>>> // Allocate space for an array section first.<br>
>>> if (size > 1 && !maptype.IsDelete)<br>
>>> <push>(base, begin, size*sizeof(Ty), clearToFrom(type));<br>
>>> <br>
>>> // Map members.<br>
>>> for (unsigned i = 0; i < size; i++) {<br>
>>> // For each component specified by this mapper:<br>
>>> for (auto c : components) {<br>
>>> if (c.hasMapper())<br>
>>> (*c.Mapper())(c.arg_base, c.arg_begin, c.arg_size, c.arg_type);<br>
>>> else<br>
>>> <push>(c.arg_base, c.arg_begin, c.arg_size, c.arg_type);<br>
>>> }<br>
>>> }<br>
>>> // Delete the array section.<br>
>>> if (size > 1 && maptype.IsDelete)<br>
>>> <push>(base, begin, size*sizeof(Ty), clearToFrom(type));<br>
>>> }<br>
>>> ```<br>
>>> <br>
>>> Comparison:<br>
>>> Why to choose 1 over 2:<br>
>>> 1. In scheme 2, the compiler needs to generate all map types and<br>
>>> pass them to __tgt_mapper through sub_components. But in this<br>
>>> case, the compiler won't be able to generate the correct MEMBER_OF<br>
>>> field in map type. As a result, the runtime has to fix it using<br>
>>> the mechanism we already have here: __tgt_mapper_num_components.<br>
>>> This not only increases complexity, but also, it means the runtime<br>
>>> needs further manipulation of the map type, which creates locality<br>
>>> issues. While in the current scheme, the map type is generated by<br>
>>> compiler once, so the data locality will be very good in this<br>
>>> case.<br>
>>> 2. In scheme 2, sub_components includes all components that should<br>
>>> be mapped. If we are mapping an array, this means we need to map<br>
>>> many components, which will need to allocate memory for<br>
>>> sub_components in the heap. This creates further memory management<br>
>>> burden and is not an efficient way to use memory.<br>
>>> 3. In scheme 1, we are able to inline nested mapper functions. As<br>
>>> a result, the compiler can do further optimizations to optimize<br>
>>> the mapper function, e.g., eliminate redundant computation, loop<br>
>>> unrolling, and thus achieve potentially better performance. We<br>
>>> cannot achieve these optimizations in scheme 2.<br>
>>> <br>
>>> Why to choose 2 over 1:<br>
>>> 1. Less code in the mapper function codegen (I doubt this because<br>
>>> the codegen function of scheme 1 uses less than 200 loc)<br>
>>> Alexey may have other reasons.<br>
>>> <br>
>>> We will appreciate if you can share your thoughts.<br>
>>> <br>
>>> Thanks,<br>
>>> Lingda Li<br>
>>> <br>
>>> -------------------------<br>
>>> FROM: Narayanaswamy, Ravi <<a href="mailto:ravi.narayanaswamy@intel.com" target="_blank">ravi.narayanaswamy@intel.com</a>><br>
>>> Sent: Wednesday, June 19, 2019 3:09 PM<br>
>>> To: 'Alexandre Eichenberger'; 'Alexey Bataev';<br>
>>> '<a href="mailto:barbara.chapman@stonybrook.edu" target="_blank">barbara.chapman@stonybrook.edu</a>'; Bobrovsky, Konstantin S; 'Carlo<br>
>>> Bertolli'; 'Chan, SiuChi'; Cownie, James H; David Oehmke; Deepak<br>
>>> Eachempati; 'Denny, Joel E.'; Dmitriev, Serguei N; Doerfert,<br>
>>> Johannes Rudolf ; '<a href="mailto:estotzer@ti.com" target="_blank">estotzer@ti.com</a>'; 'Ettore Tiotto';<br>
>>> '<a href="mailto:fraggamuffin@gmail.com" target="_blank">fraggamuffin@gmail.com</a>'; 'Gheorghe-Teod Bercea';<br>
>>> '<a href="mailto:hfinkel@anl.gov" target="_blank">hfinkel@anl.gov</a>'; '<a href="mailto:jbeyer@nvidia.com" target="_blank">jbeyer@nvidia.com</a>'; 'Jeeva Paudel'; 'Jeff<br>
>>> Heath'; Jeffrey Sandoval; Jones, Jeff C; '<a href="mailto:josem@udel.edu" target="_blank">josem@udel.edu</a>'; 'Kelvin<br>
>>> Li'; 'Kevin K O'Brien'; '<a href="mailto:khaldi.dounia@gmail.com" target="_blank">khaldi.dounia@gmail.com</a>'; 'Kotsifakou,<br>
>>> Maria'; Krishnaiyer, Rakesh; Lieberman, Ron ; '<a href="mailto:lli@bnl.gov" target="_blank">lli@bnl.gov</a>';<br>
>>> 'Lopez, Matthew Graham'; '<a href="mailto:lopezmg@ornl.org" target="_blank">lopezmg@ornl.org</a>'; 'Martin Kong'; 'Matt<br>
>>> Martineau'; Menard, Lorri; Monteleone, Robert; Narayanaswamy,<br>
>>> Ravi; 'Oscar R. Hernandez'; Rao, Premanand M; Rice, Michael P;<br>
>>> Robichaux, Joseph; Rodgers, Gregory; Rokos, Georgios;<br>
>>> '<a href="mailto:samuel.antao@ibm.com" target="_blank">samuel.antao@ibm.com</a>'; 'Sarah McNamara';<br>
>>> '<a href="mailto:sergey.y.ostanevich@gmail.com" target="_blank">sergey.y.ostanevich@gmail.com</a>'; 'Sergio Pino Gallardo'; 'Sharif,<br>
>>> Hashim'; Sjodin, Jan ; Sunil Shrestha (<a href="mailto:sshrestha@cray.com" target="_blank">sshrestha@cray.com</a>);<br>
>>> 'Sunita Chandrasekaran'; Tian, Xinmin; Tianyi Zhang;<br>
>>> '<a href="mailto:vadve@illinois.edu" target="_blank">vadve@illinois.edu</a>'; 'Wael Yehia'; 'Wang Chen'; Wilmarth, Terry L<br>
>>> Subject: OpenMP / HPC in Clang / LLVM Multi-company Telecom<br>
>>> Meeting Minutes June 19th 2019<br>
>>> <br>
>>> NEXT MEETING : JULY 10TH (MOVED FROM JULY 3RD)<br>
>>> <br>
>>> OPENS :<br>
>>> - DOCUMENTATION<br>
>>> - Greg : Can we have documents for libopenmp and Libomptarget.<br>
>>> - Alexey suggested having 3 documents: libopenmp, Libomptarget and<br>
>>> device plugin<br>
>>> - Hal will convert the existing libomptarget document. Once done<br>
>>> others can update document to capture the existing implementation<br>
>>> Future addition to libomptarget will also require update to<br>
>>> document.<br>
>>> - Next libopenmp document will be created if it does not exist or<br>
>>> updated if one exists.<br>
>>> <br>
>>> LTO FOR FAT BINARY LINKING<br>
>>> - Serguei (Intel) has an implementation which enables LTO and<br>
>>> doing away with linker scripts.<br>
>>> Everybody agreed this is a good idea, especially some linkers<br>
>>> don’t have support for linker scripts.<br>
>>> AMD is interested in enabling enabling LTO and will like to see<br>
>>> the code<br>
>>> Serguei to post the code to get feedback from all<br>
>>> - Hal to present in next meeting his proposal to support static<br>
>>> fat archives using LTO.<br>
>>> <br>
>>> OPENMP 5.0 FEATURES<br>
>>> - No update on setting up the public website. Johannes was out<br>
>>> attending ISC.<br>
>>> - New features added since last release (courtesy of Kelvin)<br>
>>> - allocate clause/allocate directive - parsing+sema, codegen<br>
>>> - mutexinout dependence-type for task<br>
>>> - user-defined mapper (declare mapper) - parsing+sema.<br>
>>> - omp_get_device_num() API routine<br>
>>> <br>
>>> DEVELOPMENT ACTIVITY<br>
>>> - ASYNC API<br>
>>> Support in Clang and libopenmp including lit test had been checked<br>
>>> in by Doru<br>
>>> <br>
>>> - MAPPER SUPPORT<br>
>>> Initial support for Mapper has been posted for review Lingda. Once<br>
>>> approved, the rest of the support will be done<br>
>>> Lingda : Should the old API being replaced by the new similar API<br>
>>> with extra mapper argument be obsoleted<br>
>>> Suggestion was for clang to not generated but keep the API in<br>
>>> libomptarget for backward compatible. In the future it can be<br>
>>> obsoleted<br>
>>> <br>
>>> - REQUIRED DIRECTIVES<br>
>>> Support for required directives has been checked in by Doru.<br>
>>> There was one issue with checking for requires directive and<br>
>>> confirming it the Declare type is TO or LINK.<br>
>>> Doru removed the check and added note to make sure if things<br>
>>> change in future need to modify this code.<br>
>>> <br>
>>> ROLL CALL :<br>
>>> <br>
>>> COMPANY<br>
>>> ATTENDEES<br>
>>> <br>
>>> 19-JUN<br>
>>> <br>
>>> AMD<br>
>>> <br>
>>> Greg Rodgers<br>
>>> <br>
>>> x<br>
>>> <br>
>>> Ashwin Aji<br>
>>> <br>
>>> Jan Sjodin<br>
>>> <br>
>>> x<br>
>>> <br>
>>> Ron Lieberman<br>
>>> <br>
>>> x<br>
>>> <br>
>>> sameer Sahasrabuddhe<br>
>>> <br>
>>> Andrey Kasaurov<br>
>>> <br>
>>> ANL<br>
>>> Hal Finkel<br>
>>> <br>
>>> x<br>
>>> <br>
>>> Johannes Doerfert<br>
>>> <br>
>>> IBM<br>
>>> Alexandre Eichenberger<br>
>>> <br>
>>> Carlo Bertolli<br>
>>> <br>
>>> Kelvin Li<br>
>>> <br>
>>> Doru<br>
>>> <br>
>>> x<br>
>>> <br>
>>> Alexey Bataev<br>
>>> <br>
>>> x<br>
>>> <br>
>>> INTEL<br>
>>> Andrey Churbanov<br>
>>> <br>
>>> Ravi Narayanaswamy<br>
>>> <br>
>>> x<br>
>>> <br>
>>> Serguei Dmitriev<br>
>>> <br>
>>> x<br>
>>> <br>
>>> Rajiv Deodhar<br>
>>> <br>
>>> Lorri Menard<br>
>>> <br>
>>> Terry Wilmarth<br>
>>> <br>
>>> Rao, Prem<br>
>>> <br>
>>> Hansang Bae<br>
>>> <br>
>>> George Rokos<br>
>>> <br>
>>> x<br>
>>> <br>
>>> CRAY<br>
>>> Deepak Eachempati<br>
>>> <br>
>>> x<br>
>>> <br>
>>> MICRON<br>
>>> John Leidel<br>
>>> <br>
>>> NVIDIA<br>
>>> James Beyer<br>
>>> <br>
>>> x<br>
>>> <br>
>>> ORNL<br>
>>> Graham Lopez<br>
>>> <br>
>>> Joel Denny<br>
>>> <br>
>>> Geoffroy Vallee<br>
>>> <br>
>>> Oscar Hernandez<br>
>>> <br>
>>> SBU/BNL<br>
>>> Lingda Li<br>
>>> <br>
>>> x<br>
>>> <br>
>>> Jose Monlsave<br>
>>> <br>
>>> Martin Kong<br>
>>> <br>
>>> TI<br>
>>> Eric Stotzer<br>
>>> <br>
>>> U OF BRISTOL<br>
>>> Mat Martineau<br>
>>> <br>
>>> U OF DELAWARE<br>
>>> Sunita Chandrasekaran<br>
>>> <br>
>>> U OF ILLINOIS<br>
>>> Hashim Sharif<br>
>>> <br>
>>> RICE<br>
>>> John Mellor-Crummey<br>
>>> <br>
>>> LSU<br>
>>> Tianyi Zhang<br>
>>> <br>
>>> <br>
>> <br>
> .........................................................................................................................................<br>
>>> àJoin Skype Meeting [3]<br>
>>> <br>
>>> Trouble Joining? Try Skype Web App [4]<br>
>>> <br>
>>> Join by phone<br>
>>> +1(916)356-2663 (or your local bridge access #) Choose bridge 5.<br>
>>> [5] (Global) English (United States)<br>
>>> Find a local number [6]<br>
>>> <br>
>>> Conference ID: 7607896966<br>
>>> Forgot your dial-in PIN? [6] |Help [7]<br>
>>> <br>
>>> [!OC([1033])!]<br>
>>> <br>
>> <br>
> .........................................................................................................................................<br>
> <br>
> <br>
> Links:<br>
> ------<br>
> [1]<br>
> <a href="https://urldefense.proofpoint.com/v2/url?u=https-3A__reviews.llvm.org_D59474&d=DwMFaQ&c=aTOVZmpUfPKZuaG9NO7J7Mh6imZbfhL47t9CpZ-pCOw&r=RLUU7gQynM_GwGu2QR7zHw&m=0c8CuLZZzM3R7PecCmFPYLuPYEOtCJHYTIGjSgIPaWU&s=EVaPRpEtSzi0Y56zmjD5fXRzN87UZDOaYp5PY3TXiVQ&e=" rel="noreferrer" target="_blank">https://urldefense.proofpoint.com/v2/url?u=https-3A__reviews.llvm.org_D59474&d=DwMFaQ&c=aTOVZmpUfPKZuaG9NO7J7Mh6imZbfhL47t9CpZ-pCOw&r=RLUU7gQynM_GwGu2QR7zHw&m=0c8CuLZZzM3R7PecCmFPYLuPYEOtCJHYTIGjSgIPaWU&s=EVaPRpEtSzi0Y56zmjD5fXRzN87UZDOaYp5PY3TXiVQ&e=</a><br>
> [2]<br>
> <a href="https://github.com/lingda-li/public-sharing/blob/master/mapper_runtime_design.pptx" rel="noreferrer" target="_blank">https://github.com/lingda-li/public-sharing/blob/master/mapper_runtime_design.pptx</a><br>
> [3]<br>
> <a href="https://urldefense.proofpoint.com/v2/url?u=https-3A__meet.intel.com_ravi.narayanaswamy_DK7943NR&d=DwMFaQ&c=aTOVZmpUfPKZuaG9NO7J7Mh6imZbfhL47t9CpZ-pCOw&r=RLUU7gQynM_GwGu2QR7zHw&m=0c8CuLZZzM3R7PecCmFPYLuPYEOtCJHYTIGjSgIPaWU&s=K4msFCmDvK4n0MdVQd7UTXRRvRkaNwLzMaP8fnX0iOg&e=" rel="noreferrer" target="_blank">https://urldefense.proofpoint.com/v2/url?u=https-3A__meet.intel.com_ravi.narayanaswamy_DK7943NR&d=DwMFaQ&c=aTOVZmpUfPKZuaG9NO7J7Mh6imZbfhL47t9CpZ-pCOw&r=RLUU7gQynM_GwGu2QR7zHw&m=0c8CuLZZzM3R7PecCmFPYLuPYEOtCJHYTIGjSgIPaWU&s=K4msFCmDvK4n0MdVQd7UTXRRvRkaNwLzMaP8fnX0iOg&e=</a><br>
> [4]<br>
> <a href="https://urldefense.proofpoint.com/v2/url?u=https-3A__meet.intel.com_ravi.narayanaswamy_DK7943NR-3Fsl-3D1&d=DwMFaQ&c=aTOVZmpUfPKZuaG9NO7J7Mh6imZbfhL47t9CpZ-pCOw&r=RLUU7gQynM_GwGu2QR7zHw&m=0c8CuLZZzM3R7PecCmFPYLuPYEOtCJHYTIGjSgIPaWU&s=krI3wEp2z8GhcZt6feFq3WgaBjcEoTDRk-GvI1BIdO8&e=" rel="noreferrer" target="_blank">https://urldefense.proofpoint.com/v2/url?u=https-3A__meet.intel.com_ravi.narayanaswamy_DK7943NR-3Fsl-3D1&d=DwMFaQ&c=aTOVZmpUfPKZuaG9NO7J7Mh6imZbfhL47t9CpZ-pCOw&r=RLUU7gQynM_GwGu2QR7zHw&m=0c8CuLZZzM3R7PecCmFPYLuPYEOtCJHYTIGjSgIPaWU&s=krI3wEp2z8GhcZt6feFq3WgaBjcEoTDRk-GvI1BIdO8&e=</a><br>
> [5]<br>
> tel:+1(916)356-2663%20(or%20your%20local%20bridge%20access%20#)%20Choose%20bridge%205.<br>
> [6]<br>
> <a href="https://urldefense.proofpoint.com/v2/url?u=https-3A__dial.intel.com&d=DwMFaQ&c=aTOVZmpUfPKZuaG9NO7J7Mh6imZbfhL47t9CpZ-pCOw&r=RLUU7gQynM_GwGu2QR7zHw&m=0c8CuLZZzM3R7PecCmFPYLuPYEOtCJHYTIGjSgIPaWU&s=g2dQtoTqaRXyBMaIUpfyoPFDRTtrQbgbWbb9b90tgBg&e=" rel="noreferrer" target="_blank">https://urldefense.proofpoint.com/v2/url?u=https-3A__dial.intel.com&d=DwMFaQ&c=aTOVZmpUfPKZuaG9NO7J7Mh6imZbfhL47t9CpZ-pCOw&r=RLUU7gQynM_GwGu2QR7zHw&m=0c8CuLZZzM3R7PecCmFPYLuPYEOtCJHYTIGjSgIPaWU&s=g2dQtoTqaRXyBMaIUpfyoPFDRTtrQbgbWbb9b90tgBg&e=</a><br>
> [7]<br>
> <a href="https://urldefense.proofpoint.com/v2/url?u=https-3A__o15.officeredir.microsoft.com_r_rlidLync15-3Fclid-3D1033-26p1-3D5-26p2-3D2009&d=DwMFaQ&c=aTOVZmpUfPKZuaG9NO7J7Mh6imZbfhL47t9CpZ-pCOw&r=RLUU7gQynM_GwGu2QR7zHw&m=0c8CuLZZzM3R7PecCmFPYLuPYEOtCJHYTIGjSgIPaWU&s=6OCBXxzOIJfra2Pewq_p-l2pY3MyKnuG-TLr7M1xq-s&e=" rel="noreferrer" target="_blank">https://urldefense.proofpoint.com/v2/url?u=https-3A__o15.officeredir.microsoft.com_r_rlidLync15-3Fclid-3D1033-26p1-3D5-26p2-3D2009&d=DwMFaQ&c=aTOVZmpUfPKZuaG9NO7J7Mh6imZbfhL47t9CpZ-pCOw&r=RLUU7gQynM_GwGu2QR7zHw&m=0c8CuLZZzM3R7PecCmFPYLuPYEOtCJHYTIGjSgIPaWU&s=6OCBXxzOIJfra2Pewq_p-l2pY3MyKnuG-TLr7M1xq-s&e=</a><br>
> _______________________________________________<br>
> cfe-dev mailing list<br>
> <a href="mailto:cfe-dev@lists.llvm.org" target="_blank">cfe-dev@lists.llvm.org</a><br>
> <a href="https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-dev" rel="noreferrer" target="_blank">https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-dev</a><br>
</blockquote></div>