[cfe-dev] EXTERNAL: Re: OpenCL compile object file symbol tables

[Dorrington, Albert]

James, Thanks. I have been trading email with Tom regarding updates to the Clover code.

I have been trying to learn more about the ELF format and just this morning came across some documentation regarding the ELF format that AMD uses with their GCN/OpenCL compiler environment (which appears to use a custom LLVM implementation.)

>From what I read, it sounds like AMD chose to create a nested ELF format, which contains both the LLVM IR and the target specific binaries.

>From http://openwall.info/wiki/john/development/GCN-ISA

The generated ELF contains the following sections:
.shstrtab
.strtab
.symtab
.llvmir - LLVM IR?
.comment - unrecognized, binary data
.rodata - contains OpenCL information, flags, SDK version, etc.
.text - contains an inner ELF

The .text section contains another ELF file. This is where the microcode (GCN bytecode) is actually stored. These are the sections of the inner ELF:
.shstrtab
.text - contains the microcode
.data - this was completely empty in a sample binary (we should check out more binaries)
.symtab
.strtab

It sounds like the r600 backend for LLVM may need to be updated to either do something similar, or provide additional information.

It would seem I still have a lot to learn/understand about this environment. :-)

Thanks
-Al

From: mankeyrabbit at gmail.com [mailto:mankeyrabbit at gmail.com] On Behalf Of James Molloy
Sent: Tuesday, January 21, 2014 8:55 AM
To: Dorrington, Albert; Tom Stellard
Cc: cfe-dev at cs.uiuc.edu
Subject: Re: EXTERNAL: Re: [cfe-dev] OpenCL compile object file symbol tables

Hi Al,

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).

Cheers,

James

On 21 January 2014 12:40, Dorrington, Albert <albert.dorrington at lmco.com<mailto:albert.dorrington at lmco.com>> wrote:
Hi James,
 Thanks for your response and explanation. It sounds as if this may be LLVM, instead of Clang.
 I have a very simple OpenCL Kernel:

__kernel void vecAdd(__global float* a) {
  int gid = get_global_id(0);
  a[gid] += a[gid];
}

I am using the Clang/LLVM tools to reproduce the GPU specific binary being generated by Mesa Clover with the following commands:

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 kernel.cl<http://kernel.cl> -o kernel.bc
llvm-link kernel.bc /usr/local/lib/clc/turks-r600--.bc -o kernel-linked.bc
opt -O2 -internalize-public-api-list=vecAdd -internalize -inline -inline-threshold=1000000000 kernel-linked.bc -o kernel-linked-opt.bc
llc -march=r600 -mcpu=turks kernel-linked-opt.bc -filetype=obj -o kernel.o

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.

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.

It seems that, using the build steps I described above, that the information I am looking for is lost with the llc command.

Thanks
-Al


From: mankeyrabbit at gmail.com<mailto:mankeyrabbit at gmail.com> [mailto:mankeyrabbit at gmail.com<mailto:mankeyrabbit at gmail.com>] On Behalf Of James Molloy
Sent: Tuesday, January 21, 2014 4:16 AM
To: Dorrington, Albert
Cc: cfe-dev at cs.uiuc.edu<mailto:cfe-dev at cs.uiuc.edu>
Subject: EXTERNAL: Re: [cfe-dev] OpenCL compile object file symbol tables

Hi Albert,

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.

Clang is often used as *part of* a CL compiler - for example in POCL (http://pocl.sourceforge.net/). 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.

Have you looked at POCL as an example of using Clang in a CL stack?

Cheers,

James

On 20 January 2014 20:15, Dorrington, Albert <albert.dorrington at lmco.com<mailto:albert.dorrington at lmco.com>> wrote:
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.

Is this by design?

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.
For instance, after loading a pre-compiled binary with clCreateProgramWithBinary() and then validating that the kernel function contained therein is the expected kernel.

Al Dorrington
Software Engineer Sr
Lockheed Martin, Mission Systems and Training


_______________________________________________
cfe-dev mailing list
cfe-dev at cs.uiuc.edu<mailto:cfe-dev at cs.uiuc.edu>
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-dev


-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/cfe-dev/attachments/20140121/9d5ae19e/attachment.html>