On Jul 7, 2014, at 3:15 , Ramin Guner <[email protected]> wrote: > Hi all, > > We are developing a device that supports OpenCL but not all of the features. > For example, our device supports only i32 and f32 types. > > We almost finished to write an LLVM backend for it. The last thing is > functions. To use sin, cos and other functions I defined custom intrinsics. > Now, I am able to compile OpenCL kernels using the steps below. > > clang --target=spir-unknown-unknown -S -emit-llvm -o cl.ll -x cl o.cl > ./replace > opt -always-inline cl.ll -S -o cl.ll > llc -march=tosun -relocation-model=pic -filetype=asm cl.ll -o out.asm > > `replace` script replaces the undefined functions as below. > OpenCL Kernel: > float fsin(float); > __kernel void test(__global float* A){ > A[0] = fsin(A[0]); > } > > LLVM IR: > ; ModuleID = 'o.cl' > target datalayout = > "e-p:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v24:32:32-v32:32:32-v48:64:64-v64:64:64-v96:128:128-v128:128:128-v192:256:256-v256:256:256-v512:512:512-v1024:1024:1024" > target triple = "spir-unknown-unknown" > > ; Function Attrs: nounwind > define void @test(float addrspace(1)* nocapture %A) #0 { > entry: > %0 = load float addrspace(1)* %A, align 4, !tbaa !2 > %call = tail call float @fsin(float %0) #2 > store float %call, float addrspace(1)* %A, align 4, !tbaa !2 > ret void > } > > declare float @fsin(float) #1 > > After `./replace`: > ... > define float @fsin(float %in) alwaysinline { > %val = call float @llvm.sin.f32(float %in) > ret float %val > } > declare float @llvm.sin.f32(float) > > My first question is if POCL makes this replacement automatically. Or, will I > do this with or without POCL?
Ramin The math functions (and other library functions) that OpenCL C code uses are declared in a header file _kernel.h, and are all explicitly provided in a kernel library. In addition, Pocl provides optimized implementations of many math functions for many architectures. If you want to provide your own functions for sin, cos etc., then you have two choices: (1) Provide intrinsics such as the ones you describe. Pocl should be using them automatically, if these intrinsics are available via <math.h>. It may also be necessary to disable Pocl's optimized implementations (i.e., disable "Vecmathlib"). (2) Implement your functions in C or C++, or provide a C or C++ interface to them, and tell Pocl to use these; this does not require defining any intrinsics. > The other question is how can I use POCL with our LLVM backend. I do not know > how to start. I assume that you mean "target" when you say "backend". The targets that pocl supports (e.g. x86-64, armv7, ...) are handled in configure.ac. You'll find several switch statements there. This translates to the equivalent of specifying a --target=... option when calling clang. > Thanks in advance. Thank you for your interest. If you have further questions, then it may be easier to schedule a Google Hangout or a Skype call than to discuss this via email. Let us know. -erik -- Erik Schnetter <[email protected]> http://www.perimeterinstitute.ca/personal/eschnetter/ My email is as private as my paper mail. I therefore support encrypting and signing email messages. Get my PGP key from http://pgp.mit.edu/.
signature.asc
Description: Message signed with OpenPGP using GPGMail
------------------------------------------------------------------------------ Open source business process management suite built on Java and Eclipse Turn processes into business applications with Bonita BPM Community Edition Quickly connect people, data, and systems into organized workflows Winner of BOSSIE, CODIE, OW2 and Gartner awards http://p.sf.net/sfu/Bonitasoft
_______________________________________________ pocl-devel mailing list [email protected] https://lists.sourceforge.net/lists/listinfo/pocl-devel
