<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=utf-8">
<meta name="Generator" content="Microsoft Word 15 (filtered medium)">
<style><!--
/* Font Definitions */
@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;}
/* Style Definitions */
p.MsoNormal, li.MsoNormal, div.MsoNormal
{margin:0in;
margin-bottom:.0001pt;
font-size:11.0pt;
font-family:"Calibri",sans-serif;}
a:link, span.MsoHyperlink
{mso-style-priority:99;
color:blue;
text-decoration:underline;}
a:visited, span.MsoHyperlinkFollowed
{mso-style-priority:99;
color:purple;
text-decoration:underline;}
p.msonormal0, li.msonormal0, div.msonormal0
{mso-style-name:msonormal;
mso-margin-top-alt:auto;
margin-right:0in;
mso-margin-bottom-alt:auto;
margin-left:0in;
font-size:11.0pt;
font-family:"Calibri",sans-serif;}
span.EmailStyle19
{mso-style-type:personal-reply;
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="blue" vlink="purple">
<div class="WordSection1">
<p class="MsoNormal">The table is used for looking up the corresponding device function or variable for host function or variable
<o:p></o:p></p>
<p class="MsoNormal">When you call a kernel, you are actually using the host address of the entry point. The runtime uses the table to look up the corresponding address in the device image to make the call.<o:p></o:p></p>
<p class="MsoNormal">The lookup is done positional entry in the table, so the table on the host must match the table on the device.<o:p></o:p></p>
<p class="MsoNormal">Ravi<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal"><b>From:</b> Openmp-dev <openmp-dev-bounces@lists.llvm.org> <b>
On Behalf Of </b>Yonghong Yan via Openmp-dev<br>
<b>Sent:</b> Wednesday, July 15, 2020 11:03 AM<br>
<b>To:</b> openmp-dev <openmp-dev@lists.llvm.org><br>
<b>Subject:</b> [Openmp-dev] The purpose of __tgt_offload_entry and how compiler generates the content for the offloading entry table and omp_offloading_entries section<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<div>
<div>
<p class="MsoNormal">Hi Friends, hope you all are cool and safe. <o:p></o:p></p>
</div>
<div>
<p class="MsoNormal"><o:p> </o:p></p>
</div>
<p class="MsoNormal">We are experimenting with the libomptarget library to directly use with source code. I read <a href="https://reviews.llvm.org/D64943">https://reviews.llvm.org/D64943</a> about how to create a fat binary using clang-offload-wrapper utility.
Given a CUDA program, I am able to create a source file for what the clang-offload-wrapper can produce. I however am not sure what to initialize the <span style="font-size:9.0pt;font-family:Consolas;color:#6A737D">__start_omp_offloading_entries</span> and <span style="font-size:9.0pt;font-family:Consolas;color:#6A737D">__stop_omp_offloading_entries
</span>symbols. They are implicitly created by the linker when linking and creating the omp_offloading_entries section. What I am not sure are what is the purpose of the offload entry table and what could be put in the entry table by the compiler/linker (the
content of the omp_offloading_entries section)? From what I read of the libomptarget design document, it seems they are used to store global variables and functions that could be mapped to devices and they are used at the runtime to help the mapping. How far
of my understanding is to what it really is?<o:p></o:p></p>
<div>
<p class="MsoNormal"><o:p> </o:p></p>
</div>
<div>
<p class="MsoNormal">For example, for the following axpy omp offloading, what could be put in the offloading entry table? I built the program and did an objdump to see the content of the section, but not much clue what they are. <o:p></o:p></p>
</div>
<div>
<p class="MsoNormal"><o:p> </o:p></p>
</div>
<div>
<p class="MsoNormal" style="margin-bottom:12.0pt">void axpy(int N, float *Y, float *X, float a) {<br>
int i,j;<br>
#pragma omp target map(to:X[0:N], N) map(tofrom:Y[0:N])<br>
#pragma omp parallel for<br>
for (i = 0; i < N; ++i) {<br>
Y[i] += a * X[i];<br>
}<br>
}<o:p></o:p></p>
</div>
<div>
<p class="MsoNormal">Thank you<o:p></o:p></p>
</div>
<div>
<p class="MsoNormal">Yonghong<o:p></o:p></p>
</div>
</div>
</div>
</body>
</html>