<html xmlns:v="urn:schemas-microsoft-com:vml" xmlns:o="urn:schemas-microsoft-com:office:office" xmlns:w="urn:schemas-microsoft-com:office:word" xmlns:m="http://schemas.microsoft.com/office/2004/12/omml" xmlns="http://www.w3.org/TR/REC-html40">
<head>
<meta http-equiv="Content-Type" content="text/html; charset=us-ascii">
<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:Consolas;
panose-1:2 11 6 9 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:105%;
font-size:11.0pt;
font-family:"Calibri",sans-serif;
color:black;}
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;}
pre
{mso-style-priority:99;
mso-style-link:"HTML Preformatted Char";
margin:0in;
margin-bottom:.0001pt;
font-size:10.0pt;
font-family:"Courier New";
color:black;}
span.EmailStyle17
{mso-style-type:personal;
font-family:"Calibri",sans-serif;
color:windowtext;}
span.HTMLPreformattedChar
{mso-style-name:"HTML Preformatted Char";
mso-style-priority:99;
mso-style-link:"HTML Preformatted";
font-family:Consolas;
color:black;}
span.EmailStyle20
{mso-style-type:personal-reply;
font-family:"Calibri",sans-serif;
color:#1F497D;}
.MsoChpDefault
{mso-style-type:export-only;
font-size:10.0pt;}
@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]-->
</head>
<body bgcolor="white" lang="EN-US" link="#0563C1" vlink="#954F72">
<div class="WordSection1">
<p class="MsoNormal"><span style="color:#0070C0">Given a global variable @gg, the compiler has to generate code on the host to specify whether it is passed by value or passed by reference. In the following example, if the compiler generates the code for passing
by value, the outlined function on the target </span><span style="font-family:""",serif;color:#0070C0;border:none windowtext 1.0pt;padding:0in">i386-pc-linux-gnu</span><span style="color:#0070C0"> cannot get the correct value since it assumes the variable
@gg is passed by reference. </span><span style="color:#0070C0"><o:p></o:p></span></p>
<p class="MsoNormal"><span style="color:#0070C0"><o:p> </o:p></span></p>
<p class="MsoNormal"><span style="color:#0070C0">Here is the corresponding IR on the host side.<o:p></o:p></span></p>
<p class="MsoNormal"><span style="color:red"> %0 = load double, double* @gg, align 8, !tbaa !3<o:p></o:p></span></p>
<p class="MsoNormal"><span style="color:red"> %1 = bitcast double %0 to i64<o:p></o:p></span></p>
<p class="MsoNormal"><span style="color:#1F497D"> …<o:p></o:p></span></p>
<p class="MsoNormal"><span style="color:#0070C0"> %12 = getelementptr inbounds [4 x i8*], [4 x i8*]* %.offload_baseptrs, i32 0, i32 2<o:p></o:p></span></p>
<p class="MsoNormal"><span style="color:#0070C0"> %13 = bitcast i8** %12 to i64*<o:p></o:p></span></p>
<p class="MsoNormal"><span style="color:red"> store i64 %1, i64* %13, align 8<o:p></o:p></span></p>
<p class="MsoNormal"><span style="color:#0070C0"> …<o:p></o:p></span></p>
<p class="MsoNormal"><span style="color:#0070C0"> %20 = call i32 @__tgt_target(i64 -1, i8* @.omp_offload.region_id, i32 4, i8** %4, i8** %6, i64* getelementptr inbounds ([4 x i64], [4 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([4
x i64], [4 x i64]* @.offload_maptypes, i32 0, i32 0)) <o:p></o:p></span></p>
<p class="MsoNormal"><span style="color:#0070C0"><o:p> </o:p></span></p>
<p class="MsoNormal"><span style="color:#0070C0">Thanks,<o:p></o:p></span></p>
<p class="MsoNormal"><span style="color:#0070C0"><o:p> </o:p></span></p>
<p class="MsoNormal"><span style="color:#0070C0">Jin<o:p></o:p></span></p>
<p class="MsoNormal"><span style="color:#1F497D"><o:p> </o:p></span></p>
<div>
<div style="border:none;border-top:solid #E1E1E1 1.0pt;padding:3.0pt 0in 0in 0in">
<p class="MsoNormal" style="margin-bottom:0in;margin-bottom:.0001pt;line-height:normal">
<a name="_____replyseparator"></a><b><span style="color:windowtext">From:</span></b><span style="color:windowtext"> Friedman, Eli [mailto:efriedma@codeaurora.org]
<br>
<b>Sent:</b> Wednesday, April 25, 2018 3:18 PM<br>
<b>To:</b> Lin, Jin <jin.lin@intel.com>; 'llvm-dev@lists.llvm.org' <llvm-dev@lists.llvm.org><br>
<b>Subject:</b> Re: [llvm-dev] [LLVM][RFC] Representing the target device information in the LLVM IR<o:p></o:p></span></p>
</div>
</div>
<p class="MsoNormal"><o:p> </o:p></p>
<div>
<p class="MsoNormal">On 4/25/2018 3:05 PM, Lin, Jin via llvm-dev wrote:<span style="font-size:12.0pt;line-height:105%"><o:p></o:p></span></p>
</div>
<blockquote style="margin-top:5.0pt;margin-bottom:5.0pt">
<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><o:p></o:p></p>
<p class="MsoNormal"><span style="font-family:""",serif;color:#0070C0;border:none windowtext 1.0pt;padding:0in">===========================================================================</span><o:p></o:p></p>
<p class="MsoNormal"><span style="font-family:""",serif;color:#0070C0;border:none windowtext 1.0pt;padding:0in"> </span><o:p></o:p></p>
<p class="MsoNormal"><span style="font-family:""",serif;color:#0070C0;border:none windowtext 1.0pt;padding:0in">Why this RFC change?</span><o:p></o:p></p>
<p class="MsoNormal"><span style="font-family:""",serif;color:#0070C0;border:none windowtext 1.0pt;padding:0in">=================</span><o:p></o:p></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.
</span><o:p></o:p></p>
<p class="MsoNormal"><span style="font-family:""",serif;color:#0070C0;border:none windowtext 1.0pt;padding:0in"> </span><o:p></o:p></p>
<p class="MsoNormal"><span style="font-family:""",serif;color:#0070C0;border:none windowtext 1.0pt;padding:0in">Existing code: 64-bit firstprivate variable</span><o:p></o:p></p>
<p class="MsoNormal"><span style="font-family:""",serif;color:#0070C0;border:none windowtext 1.0pt;padding:0in"> </span><o:p></o:p></p>
<p class="MsoNormal" style="text-indent:.5in"><span style="font-family:""",serif;color:#0070C0;border:none windowtext 1.0pt;padding:0in">void foo() {</span><o:p></o:p></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;</span><o:p></o:p></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)</span><o:p></o:p></p>
<p class="MsoNormal" style="text-indent:.5in"><span style="font-family:""",serif;color:#0070C0;border:none windowtext 1.0pt;padding:0in"> {}</span><o:p></o:p></p>
<p class="MsoNormal" style="text-indent:.5in"><span style="font-family:""",serif;color:#0070C0;border:none windowtext 1.0pt;padding:0in">}</span><o:p></o:p></p>
<p class="MsoNormal"><span style="font-family:""",serif;color:#0070C0;border:none windowtext 1.0pt;padding:0in"> </span><o:p></o:p></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 …</span><o:p></o:p></p>
<p class="MsoNormal"><span style="font-family:""",serif;color:#0070C0;border:none windowtext 1.0pt;padding:0in"> </span><o:p></o:p></p>
<p class="MsoNormal"><span style="font-family:""",serif;color:#0070C0;border:none windowtext 1.0pt;padding:0in">x86_64-mic</span><o:p></o:p></p>
<p class="MsoNormal"><span style="font-family:""",serif;color:#0070C0;border:none windowtext 1.0pt;padding:0in"> </span><o:p></o:p></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 {</span><o:p></o:p></p>
<p class="MsoNormal" style="text-indent:.5in"><span style="font-family:""",serif;color:#0070C0;border:none windowtext 1.0pt;padding:0in">entry:</span><o:p></o:p></p>
<p class="MsoNormal" style="text-indent:.5in"><span style="font-family:""",serif;color:#0070C0;border:none windowtext 1.0pt;padding:0in">…</span><o:p></o:p></p>
<p class="MsoNormal" style="text-indent:.5in"><span style="font-family:""",serif;color:#0070C0;border:none windowtext 1.0pt;padding:0in">}</span><o:p></o:p></p>
<p class="MsoNormal"><span style="font-family:""",serif;color:#0070C0;border:none windowtext 1.0pt;padding:0in"> </span><o:p></o:p></p>
<p class="MsoNormal"><span style="font-family:""",serif;color:#0070C0;border:none windowtext 1.0pt;padding:0in">i386-pc-linux-gnu</span><o:p></o:p></p>
<p class="MsoNormal"><span style="font-family:""",serif;color:#0070C0;border:none windowtext 1.0pt;padding:0in"> </span><o:p></o:p></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 {</span><o:p></o:p></p>
<p class="MsoNormal" style="text-indent:.5in"><span style="font-family:""",serif;color:#0070C0;border:none windowtext 1.0pt;padding:0in">entry:</span><o:p></o:p></p>
<p class="MsoNormal" style="text-indent:.5in"><span style="font-family:""",serif;color:#0070C0;border:none windowtext 1.0pt;padding:0in"> …</span><o:p></o:p></p>
<p class="MsoNormal" style="text-indent:.5in"><span style="font-family:""",serif;color:#0070C0;border:none windowtext 1.0pt;padding:0in">}</span><o:p></o:p></p>
<p class="MsoNormal"><span style="font-family:""",serif;color:#0070C0;border:none windowtext 1.0pt;padding:0in"> </span><o:p></o:p></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><o:p></o:p></p>
</blockquote>
<p class="MsoNormal" style="margin-bottom:0in;margin-bottom:.0001pt;line-height:normal">
<span style="font-size:12.0pt;font-family:"Times New Roman",serif"><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>
<br>
<o:p></o:p></span></p>
<pre>-- <o:p></o:p></pre>
<pre>Employee of Qualcomm Innovation Center, Inc.<o:p></o:p></pre>
<pre>Qualcomm Innovation Center, Inc. is a member of Code Aurora Forum, a Linux Foundation Collaborative Project<o:p></o:p></pre>
</div>
</body>
</html>