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

[Dorrington, Albert]

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 -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] On Behalf Of James Molloy
Sent: Tuesday, January 21, 2014 4:16 AM
To: Dorrington, Albert
Cc: 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/6bd3be69/attachment.html>