<html>
<head>
<meta http-equiv="Content-Type" content="text/html;
charset=windows-1252">
</head>
<body text="#000000" bgcolor="#FFFFFF">
<div class="moz-cite-prefix">On 4/25/2018 3:05 PM, Lin, Jin via
llvm-dev wrote:<br>
</div>
<blockquote type="cite"
cite="mid:262352419881AB4EA6E7EB413CFCA81CC4DB8030@ORSMSX108.amr.corp.intel.com">
<meta http-equiv="Content-Type" content="text/html;
charset=windows-1252">
<meta name="Generator" content="Microsoft Word 15 (filtered
medium)">
<style><!--
/* Font Definitions */
@font-face
{font-family:SimSun;
panose-1:2 1 6 0 3 1 1 1 1 1;}
@font-face
{font-family:"Cambria Math";
panose-1:2 4 5 3 5 4 6 3 2 4;}
@font-face
{font-family:Calibri;
panose-1:2 15 5 2 2 2 4 3 2 4;}
@font-face
{font-family:"\@SimSun";
panose-1:2 1 6 0 3 1 1 1 1 1;}
@font-face
{font-family:"
panose-1:0 0 0 0 0 0 0 0 0 0;}
/* Style Definitions */
p.MsoNormal, li.MsoNormal, div.MsoNormal
{margin-top:0in;
margin-right:0in;
margin-bottom:8.0pt;
margin-left:0in;
line-height:106%;
font-size:11.0pt;
font-family:"Calibri",sans-serif;}
a:link, span.MsoHyperlink
{mso-style-priority:99;
color:#0563C1;
text-decoration:underline;}
a:visited, span.MsoHyperlinkFollowed
{mso-style-priority:99;
color:#954F72;
text-decoration:underline;}
span.EmailStyle17
{mso-style-type:personal-compose;
font-family:"Calibri",sans-serif;
color:windowtext;}
.MsoChpDefault
{mso-style-type:export-only;
font-family:"Calibri",sans-serif;}
@page WordSection1
{size:8.5in 11.0in;
margin:1.0in 1.0in 1.0in 1.0in;}
div.WordSection1
{page:WordSection1;}
--></style><!--[if gte mso 9]><xml>
<o:shapedefaults v:ext="edit" spidmax="1026" />
</xml><![endif]--><!--[if gte mso 9]><xml>
<o:shapelayout v:ext="edit">
<o:idmap v:ext="edit" data="1" />
</o:shapelayout></xml><![endif]-->
<div class="WordSection1">
<p class="MsoNormal"><span
style="font-family:""",serif;color:#0070C0;border:none
windowtext 1.0pt;padding:0in">RFC: Representing the target
device information in the LLVM IR</span><span
style="font-size:10.0pt;line-height:106%;font-family:""",serif;color:#0070C0"><o:p></o:p></span></p>
<p class="MsoNormal"><span
style="font-family:""",serif;color:#0070C0;border:none
windowtext 1.0pt;padding:0in">===========================================================================</span><span
style="font-size:10.0pt;line-height:106%;font-family:""",serif;color:#0070C0"><o:p></o:p></span></p>
<p class="MsoNormal"><span
style="font-family:""",serif;color:#0070C0;border:none
windowtext 1.0pt;padding:0in"> </span><span
style="font-size:10.0pt;line-height:106%;font-family:""",serif;color:#0070C0"><o:p></o:p></span></p>
<p class="MsoNormal"><span
style="font-family:""",serif;color:#0070C0;border:none
windowtext 1.0pt;padding:0in">Why this RFC change?</span><span
style="font-size:10.0pt;line-height:106%;font-family:""",serif;color:#0070C0"><o:p></o:p></span></p>
<p class="MsoNormal"><span
style="font-family:""",serif;color:#0070C0;border:none
windowtext 1.0pt;padding:0in">=================</span><span
style="font-size:10.0pt;line-height:106%;font-family:""",serif;color:#0070C0"><o:p></o:p></span></p>
<p class="MsoNormal"><span
style="font-family:""",serif;color:#0070C0;border:none
windowtext 1.0pt;padding:0in">The target device information
needs to be passed to the LLVM backend when OpenMP backend
outlining is enabled. For example, for multiple target
devices, the target compilation has to generate a single
host to support all the targets. In order to make sure all
the target outlined functions have the same interface, the
information of all the target architectures is needed during
host and target compilation. In the following example, the
firstprivate variable ‘d’ is passed by value under
x86_64-mic and passed by reference under i386-pc-linux-gnu.
In order to avoid this inconsistency, the compiler needs all
the target architecture information so it can find a common
interface. In this example, it will change the x86_64-mic
interface for ‘d’ to pass by reference.
<o:p></o:p></span></p>
<p class="MsoNormal"><span
style="font-family:""",serif;color:#0070C0;border:none
windowtext 1.0pt;padding:0in"><o:p> </o:p></span></p>
<p class="MsoNormal"><span
style="font-family:""",serif;color:#0070C0;border:none
windowtext 1.0pt;padding:0in">Existing code: 64-bit
firstprivate variable<o:p></o:p></span></p>
<p class="MsoNormal"><span
style="font-family:""",serif;color:#0070C0;border:none
windowtext 1.0pt;padding:0in"><o:p> </o:p></span></p>
<p class="MsoNormal" style="text-indent:.5in"><span
style="font-family:""",serif;color:#0070C0;border:none
windowtext 1.0pt;padding:0in">void foo() {<o:p></o:p></span></p>
<p class="MsoNormal" style="text-indent:.5in"><span
style="font-family:""",serif;color:#0070C0;border:none
windowtext 1.0pt;padding:0in"> double d = 1.0;<o:p></o:p></span></p>
<p class="MsoNormal" style="text-indent:.5in"><span
style="font-family:""",serif;color:#0070C0;border:none
windowtext 1.0pt;padding:0in"> #pragma omp target
firstprivate(d)<o:p></o:p></span></p>
<p class="MsoNormal" style="text-indent:.5in"><span
style="font-family:""",serif;color:#0070C0;border:none
windowtext 1.0pt;padding:0in"> {}<o:p></o:p></span></p>
<p class="MsoNormal" style="text-indent:.5in"><span
style="font-family:""",serif;color:#0070C0;border:none
windowtext 1.0pt;padding:0in">}<o:p></o:p></span></p>
<p class="MsoNormal"><span
style="font-family:""",serif;color:#0070C0;border:none
windowtext 1.0pt;padding:0in"><o:p> </o:p></span></p>
<p class="MsoNormal"><span
style="font-family:""",serif;color:#0070C0;border:none
windowtext 1.0pt;padding:0in">$clang –fopenmp-backend
-fopenmp-targets=x86_64-mic, i386-pc-linux-gnu …<o:p></o:p></span></p>
<p class="MsoNormal"><span
style="font-family:""",serif;color:#0070C0;border:none
windowtext 1.0pt;padding:0in"><o:p> </o:p></span></p>
<p class="MsoNormal"><span
style="font-family:""",serif;color:#0070C0;border:none
windowtext 1.0pt;padding:0in">x86_64-mic<o:p></o:p></span></p>
<p class="MsoNormal"><span
style="font-family:""",serif;color:#0070C0;border:none
windowtext 1.0pt;padding:0in"><o:p> </o:p></span></p>
<p class="MsoNormal" style="text-indent:.5in"><span
style="font-family:""",serif;color:#0070C0;border:none
windowtext 1.0pt;padding:0in">define void
@__omp_offloading…(i64 %d) #0 {<o:p></o:p></span></p>
<p class="MsoNormal" style="text-indent:.5in"><span
style="font-family:""",serif;color:#0070C0;border:none
windowtext 1.0pt;padding:0in">entry:<o:p></o:p></span></p>
<p class="MsoNormal" style="text-indent:.5in"><span
style="font-family:""",serif;color:#0070C0;border:none
windowtext 1.0pt;padding:0in">…<o:p></o:p></span></p>
<p class="MsoNormal" style="text-indent:.5in"><span
style="font-family:""",serif;color:#0070C0;border:none
windowtext 1.0pt;padding:0in">}<o:p></o:p></span></p>
<p class="MsoNormal"><span
style="font-family:""",serif;color:#0070C0;border:none
windowtext 1.0pt;padding:0in"><o:p> </o:p></span></p>
<p class="MsoNormal"><span
style="font-family:""",serif;color:#0070C0;border:none
windowtext 1.0pt;padding:0in">i386-pc-linux-gnu<o:p></o:p></span></p>
<p class="MsoNormal"><span
style="font-family:""",serif;color:#0070C0;border:none
windowtext 1.0pt;padding:0in"><o:p> </o:p></span></p>
<p class="MsoNormal" style="text-indent:.5in"><span
style="font-family:""",serif;color:#0070C0;border:none
windowtext 1.0pt;padding:0in">define void
@__omp_offloading…(double* dereferenceable(8) %d) #0 {<o:p></o:p></span></p>
<p class="MsoNormal" style="text-indent:.5in"><span
style="font-family:""",serif;color:#0070C0;border:none
windowtext 1.0pt;padding:0in">entry:<o:p></o:p></span></p>
<p class="MsoNormal" style="text-indent:.5in"><span
style="font-family:""",serif;color:#0070C0;border:none
windowtext 1.0pt;padding:0in"> …<o:p></o:p></span></p>
<p class="MsoNormal" style="text-indent:.5in"><span
style="font-family:""",serif;color:#0070C0;border:none
windowtext 1.0pt;padding:0in">}<o:p></o:p></span></p>
<p class="MsoNormal"><span
style="font-family:""",serif;color:#0070C0;border:none
windowtext 1.0pt;padding:0in"><o:p> </o:p></span></p>
<p class="MsoNormal"><span
style="font-family:""",serif;color:#0070C0;border:none
windowtext 1.0pt;padding:0in">There is an inconsistency
between host and target part(s) of the program!</span></p>
</div>
</blockquote>
<br>
I don't see how this inconsistency is a problem... at least, not on
its own. The host code doesn't call either of these functions
directly; it calls the OpenMP runtime, which should invoke the
offloaded function correctly. (If it doesn't, that's a bug in the
OpenMP lowering, not the LLVM backend.)<br>
<br>
-Eli<br>
<pre class="moz-signature" cols="72">--
Employee of Qualcomm Innovation Center, Inc.
Qualcomm Innovation Center, Inc. is a member of Code Aurora Forum, a Linux Foundation Collaborative Project</pre>
</body>
</html>