Hi.
First.
#include <Im-doing-something-somewhat-odd.h>
I'm trying to use a current clang/llvm (current as in git checkout from just the other day) to build an opencl kernel and then link that with some code which has been compiled with gcc/g++.
When the clang .o is linked to the gcc/gcc+, I'm getting /home/tgall/opencl/SNU/tmp2/cl_temp_1.tkl uses VFP register arguments, /home/tgall/opencl/SNU/tmp2/cl_temp_1.o does not
the cl_temp_1.o was produced with clang. the cl_temp_1.tkl via gcc/g++.
Let's dive into details.
This is following in the footsteps of an open source framework called SNU which implements OpenCL. Within SNU they had a fairly old version of clang+llvm which wouldn't even build on ARM so step one has been to figure out what SNU was doing with clang and replicate this using latest clang.
So given the following minimal test kernel placed into cl_temp_1.cl
/* Header to make Clang compatible with OpenCL */ #define __global __attribute__((address_space(1)))
int get_global_id(int index);
/* Test kernel */ __kernel void test(__global float *in, __global float *out) { int index = get_global_id(0); out[index] = 3.14159f * in[index] + in[index]; }
then we following the following steps: clang -mfloat-abi=hard -mfpu=neon -S -emit-llvm -x cl -I/home/tgall/opencl/SNU/src/compiler/tools/clang/lib/Headers -I/home/tgall/opencl/SNU/inc -include /home/tgall/opencl/SNU/inc/comp/cl_kernel.h /home/tgall/opencl/SNU/tmp2/cl_temp_1.cl -o /home/tgall/opencl/SNU/tmp2/cl_temp_1.ll
with the resulting cl_temp_1.ll we:
llc /home/tgall/opencl/SNU/tmp2/cl_temp_1.ll
which results in cl_temp_1.s. Then:
clang -c -mfloat-abi=hard -mfpu=neon -o /home/tgall/opencl/SNU/tmp2/cl_temp_1.o /home/tgall/opencl/SNU/tmp2/cl_temp_1.s
so now in theory we should have a perflectly good cl_temp_1.o ready for linking.
But first let's get the bits ready that will be built by the traditional gnu toolschain. We have:
gcc -shared -fPIC -O3 -o /home/tgall/opencl/SNU/tmp2/cl_temp_1_info.so /home/tgall/opencl/SNU/tmp2/cl_temp_1_info.c
and
gcc -shared -fPIC -march=armv7-a -mtune=cortex-a9 -mfloat-abi=hard -mfpu=neon -fsigned-char -DDEF_INCLUDE_ARM -I. -I /home/tgall/opencl/SNU/src/runtime/hal/device/cpu -I /home/tgall/opencl/SNU/src/runtime/include -I /home/tgall/opencl/SNU/src/runtime/core -I /home/tgall/opencl/SNU/src/runtime/core/device -I /home/tgall/opencl/SNU/src/runtime/hal -I /home/tgall/opencl/SNU/src/runtime/hal/device -DTARGET_MACH_CPU -O3 -c /home/tgall/opencl/SNU/src/runtime/hal/device/cpu/hal.c -o /home/tgall/opencl/SNU/tmp2/hal.o
And here we try to link it all together.
g++ -shared -fPIC -march=armv7-a -mtune=cortex-a9 -mfloat-abi=hard -mfpu=neon -fsigned-char -DDEF_INCLUDE_ARM -O3 -o /home/tgall/opencl/SNU/tmp2/cl_temp_1.tkl /home/tgall/opencl/SNU/tmp2/hal.o /home/tgall/opencl/SNU/tmp2/cl_temp_1.o -L/home/tgall/opencl/SNU/lib/lnx_arm -lsnusamsung_opencl_builtin_lnx_arm -lpthread -lm
and bang we're back to the error I first mentioned: /usr/bin/ld: error: /home/tgall/opencl/SNU/tmp2/cl_temp_1.tkl uses VFP register arguments, /home/tgall/opencl/SNU/tmp2/cl_temp_1.o does not
so first obvious question is -mfloat-abi=hard -mfpu=neon correct for clang?
tgall@miranda:~/opencl/SNU/tmp2$ clang --version clang version 3.3 Target: armv7l-unknown-linux-gnueabihf Thread model: posix
Thanks for any suggestions!
-- Regards, Tom
"Where's the kaboom!? There was supposed to be an earth-shattering kaboom!" Marvin Martian Tech Lead, Graphics Working Group | Linaro.org │ Open source software for ARM SoCs w) tom.gall att linaro.org h) tom_gall att mac.com