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?
The other question is how can I use POCL with our LLVM backend. I do not
know how to start.
Thanks in advance.
Ramin.
------------------------------------------------------------------------------
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
