<div dir="ltr">Hi Al,<div><br></div><div>So this is something to do with the r600 backend and how it exports to ELF. I've reproduced your commands and see an ELF file with all anonymous symbols, as you said. I suggest that this is a question for the r600 maintainer, Tom Stellard (CC'd).</div>
<div><br></div><div>Cheers,</div><div><br></div><div>James</div></div><div class="gmail_extra"><br><br><div class="gmail_quote">On 21 January 2014 12:40, Dorrington, Albert <span dir="ltr"><<a href="mailto:albert.dorrington@lmco.com" target="_blank">albert.dorrington@lmco.com</a>></span> wrote:<br>
<blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
<div lang="EN-US" link="blue" vlink="purple">
<div>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1f497d">Hi James,<u></u><u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1f497d"><u></u> <u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1f497d">Thanks for your response and explanation. It sounds as if this may be LLVM, instead of Clang.<u></u><u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1f497d"><u></u> <u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1f497d">I have a very simple OpenCL Kernel:<u></u><u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1f497d"><u></u> <u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1f497d">__kernel void vecAdd(__global float* a) {<u></u><u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1f497d"> int gid = get_global_id(0);<u></u><u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1f497d"> a[gid] += a[gid];<u></u><u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1f497d">}<u></u><u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1f497d"><u></u> <u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1f497d">I am using the Clang/LLVM tools to reproduce the GPU specific binary being generated by Mesa Clover with the following commands:<u></u><u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1f497d"><u></u> <u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1f497d">clang –O0 –emit-llvm –include /usr/local/include/clc.clc.h –I /usr/local/include –Dcl_clang_storage_class_specifies –target r600 –mcpu=turks –c <a href="http://kernel.cl" target="_blank">kernel.cl</a> –o
kernel.bc<u></u><u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1f497d">llvm-link kernel.bc /usr/local/lib/clc/turks-r600--.bc –o kernel-linked.bc<u></u><u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1f497d">opt –O2 –internalize-public-api-list=vecAdd –internalize –inline –inline-threshold=1000000000 kernel-linked.bc –o kernel-linked-opt.bc<u></u><u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1f497d">llc –march=r600 –mcpu=turks kernel-linked-opt.bc –filetype=obj –o kernel.o<u></u><u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1f497d"><u></u> <u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1f497d">Currently the clCreateProgramWithBinary() accepts the LLVM IR, not the ELF binary objects. I am looking to change that, for an embedded environment where the
kernels would be pre-compiled using the LLVM/Clang tools.<u></u><u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1f497d"><u></u> <u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1f497d">The problem that I see is that the kernel.o ELF file does not appear to list the function names in the symbol table. So, I’m not seeing how I could implement
the clCreateKernel() call to lookup the function within the ELF object. Or for that matter, if the ELF contained more than one kernel function, how I would retrieve the names within a clCreateKernelsInProgram() call.<u></u><u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1f497d"><u></u> <u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1f497d">It seems that, using the build steps I described above, that the information I am looking for is lost with the llc command.<u></u><u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1f497d"><u></u> <u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1f497d">Thanks<u></u><u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1f497d">-Al<u></u><u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1f497d"><u></u> <u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1f497d"><u></u> <u></u></span></p>
<div style="border:none;border-left:solid blue 1.5pt;padding:0in 0in 0in 4.0pt">
<div>
<div style="border:none;border-top:solid #b5c4df 1.0pt;padding:3.0pt 0in 0in 0in">
<p class="MsoNormal"><b><span style="font-size:10.0pt;font-family:"Tahoma","sans-serif"">From:</span></b><span style="font-size:10.0pt;font-family:"Tahoma","sans-serif""> <a href="mailto:mankeyrabbit@gmail.com" target="_blank">mankeyrabbit@gmail.com</a> [mailto:<a href="mailto:mankeyrabbit@gmail.com" target="_blank">mankeyrabbit@gmail.com</a>]
<b>On Behalf Of </b>James Molloy<br>
<b>Sent:</b> Tuesday, January 21, 2014 4:16 AM<br>
<b>To:</b> Dorrington, Albert<br>
<b>Cc:</b> <a href="mailto:cfe-dev@cs.uiuc.edu" target="_blank">cfe-dev@cs.uiuc.edu</a><br>
<b>Subject:</b> EXTERNAL: Re: [cfe-dev] OpenCL compile object file symbol tables<u></u><u></u></span></p>
</div>
</div>
<p class="MsoNormal"><u></u> <u></u></p>
<div>
<p class="MsoNormal">Hi Albert,<u></u><u></u></p>
<div>
<p class="MsoNormal"><u></u> <u></u></p>
</div>
<div>
<p class="MsoNormal">Clang is not in and of itself an OpenCL compiler. It has a frontend for OpenCL-C, and can produce LLVM-IR from that (which can then be pushed through LLVM to produce some machine code). But those generated functions would be useless.<u></u><u></u></p>
</div>
<div>
<p class="MsoNormal"><u></u> <u></u></p>
</div>
<div>
<p class="MsoNormal">Clang is often used as *part of* a CL compiler - for example in POCL (<a href="http://pocl.sourceforge.net/" target="_blank">http://pocl.sourceforge.net/</a>). The IR generated from Clang for CL-C code doesn't contain any details of how it is going to
be executed. For example, is it going on a GPU or a CPU? In the latter case, loops will need to be inserted and calls to get_local_id() will need to reference the loop induction variables. There'll need to be some way inserted of being able to pass the group
ID and other payload-global data too.<u></u><u></u></p>
</div>
<div>
<p class="MsoNormal"><u></u> <u></u></p>
</div>
<div>
<p class="MsoNormal">Have you looked at POCL as an example of using Clang in a CL stack?<u></u><u></u></p>
</div>
<div>
<p class="MsoNormal"><u></u> <u></u></p>
</div>
<div>
<p class="MsoNormal">Cheers,<u></u><u></u></p>
</div>
<div>
<p class="MsoNormal"><u></u> <u></u></p>
</div>
<div>
<p class="MsoNormal">James<u></u><u></u></p>
</div>
</div>
<div>
<p class="MsoNormal" style="margin-bottom:12.0pt"><u></u> <u></u></p>
<div>
<p class="MsoNormal">On 20 January 2014 20:15, Dorrington, Albert <<a href="mailto:albert.dorrington@lmco.com" target="_blank">albert.dorrington@lmco.com</a>> wrote:<u></u><u></u></p>
<div>
<div>
<p class="MsoNormal">When I compile an OpenCL kernel to a binary file, using llvm/clang, I don’t see the kernel functions defined within the symbol table.<u></u><u></u></p>
<p class="MsoNormal"> <u></u><u></u></p>
<p class="MsoNormal">Is this by design?<u></u><u></u></p>
<p class="MsoNormal"> <u></u><u></u></p>
<p class="MsoNormal">Without the function names in the symbol table, I’m not sure how I would look up a function within a pre-compiled object file.<u></u><u></u></p>
<p class="MsoNormal">For instance, after loading a pre-compiled binary with clCreateProgramWithBinary() and then validating that the kernel function contained therein is the expected kernel.<u></u><u></u></p>
<p class="MsoNormal"> <u></u><u></u></p>
<p class="MsoNormal"><b><span style="font-size:14.0pt">Al Dorrington</span></b><u></u><u></u></p>
<p class="MsoNormal"><i><span style="font-size:10.0pt">Software Engineer Sr</span></i><u></u><u></u></p>
<p class="MsoNormal"><i><span style="font-size:10.0pt">Lockheed Martin, Mission Systems and Training</span></i><u></u><u></u></p>
<p class="MsoNormal"> <u></u><u></u></p>
</div>
</div>
<p class="MsoNormal" style="margin-bottom:12.0pt"><br>
_______________________________________________<br>
cfe-dev mailing list<br>
<a href="mailto:cfe-dev@cs.uiuc.edu" target="_blank">cfe-dev@cs.uiuc.edu</a><br>
<a href="http://lists.cs.uiuc.edu/mailman/listinfo/cfe-dev" target="_blank">http://lists.cs.uiuc.edu/mailman/listinfo/cfe-dev</a><u></u><u></u></p>
</div>
<p class="MsoNormal"><u></u> <u></u></p>
</div>
</div>
</div>
</div>
</blockquote></div><br></div>