<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:"\@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]-->
</head>
<body lang="EN-US" link="#0563C1" vlink="#954F72">
<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!<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">Change Made
</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">We proposed new module level attribute to represent target device information.<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"><i><span style="font-family:""",serif;color:#0070C0;border:none windowtext 1.0pt;padding:0in">/// Get the target device information, which is a comma-separated string</span></i><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"><i><span style="font-family:""",serif;color:#0070C0;border:none windowtext 1.0pt;padding:0in">/// describing one or more devices.</span></i><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"><i><span style="font-family:""",serif;color:#0070C0;border:none windowtext 1.0pt;padding:0in">const std::string &getTargetDevices() const { return TargetDevices; }</span></i><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" style="text-indent:.5in"><i><span style="font-family:""",serif;color:#0070C0;border:none windowtext 1.0pt;padding:0in">/// set the target device information.</span></i><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"><i><span style="font-family:""",serif;color:#0070C0;border:none windowtext 1.0pt;padding:0in">void setTargetDevices(StringRef T) { TargetDevices = T; }</span></i><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" style="text-indent:.5in"><span style="font-family:""",serif;color:#0070C0;border:none windowtext 1.0pt;padding:0in">IR Dump (the extension indicated in red font)<o:p></o:p></span></p>
<p class="MsoNormal" style="text-indent:.5in"><i><span style="font-family:""",serif;color:#0070C0;border:none windowtext 1.0pt;padding:0in">target triple = "x86_64-unknown-linux-gnu"</span></i><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"><i><span style="font-family:""",serif;color:red;border:none windowtext 1.0pt;padding:0in">target device_triples = "x86_64-mic,i386-pc-linux-gnu"</span></i><span style="font-family:""",serif;color:red;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">========</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">Summary:</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">========<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-family:""",serif;color:#0070C0;border:none windowtext 1.0pt;padding:0in">We propose a new module-level attribute to represent the target device information embedded in the LLVM IR. The change is simple and straightforward.
The patch is uploaded with this RFC for code review.</span><span style="font-size:12.0pt;line-height:106%"><o:p></o:p></span></p>
<p class="MsoNormal"><a href="https://reviews.llvm.org/D46071"><span style="font-size:10.0pt;line-height:106%;font-family:""",serif">https://reviews.llvm.org/D46071</span></a><span style="font-size:10.0pt;line-height:106%;font-family:""",serif;color:#0070C0"><o:p></o:p></span></p>
<p class="MsoNormal"><a href="https://reviews.llvm.org/D46074"><span style="font-size:10.0pt;line-height:106%;font-family:""",serif">https://reviews.llvm.org/D46074</span></a><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-size:10.0pt;line-height:106%;font-family:""",serif;color:#0070C0"><o:p> </o:p></span></p>
<p class="MsoNormal"><o:p> </o:p></p>
</div>
</body>
</html>