OpenCL under FreeBSD-10?

[Eduardo Morras]

On Sat, 05 Jul 2014 15:11:51 +0930
Shane Ambler <FreeBSD at ShaneWare.Biz> wrote:

> On 05/07/2014 08:31, Sam Fourman Jr. wrote:
> > I am also interested in this..
> > 
> > 
> > On Fri, Jul 4, 2014 at 5:49 PM, Eduardo Morras <emorrasg at yahoo.es>
> > wrote:
> > 
> >> On Thu, 03 Jul 2014 18:41:00 -0700
> >> Dennis Glatting <freebsd at pki2.com> wrote:
> >>
> >>> Is there a working version of OpenCL somewhere for FreeBSD10?
> >>> freeocl is broken, badly.
> >>
> >> Googling a bit shows llvm libclc project that outputs PTX (Nvidia)
> >> code. Clang+llvm is part of the system, don't know if this piece
> >> of llvm project has been ported too.
> >>
> >> http://llvm.org
> >>
> >> http://libclc.llvm.org/
> 
> According to http://llvm.org/docs/doxygen/html/Triple_8h_source.html
> llvm supports spir and spir64 for "standard portable IR for OpenCL"
> 
> Unfortunately we don't appear to have this enabled on freebsd.
> From a ports installed llvm - `llc33 --version`
> (base installed llc is a shorter list of only supported host cpu's)
> 
> LLVM (http://llvm.org/):
>   LLVM version 3.3
>   Optimized build.
>   Built Mar 21 2014 (05:38:50).
>   Default target: amd64-portbld-freebsd9.2
>   Host CPU: corei7-avx
> 
>   Registered Targets:

>     nvptx    - NVIDIA PTX 32-bit
>     nvptx64  - NVIDIA PTX 64-bit

Then the target for nvidia assembler is enabled. My output for llc33 and llc34 --version shows the same results. This permit us to convert from llvm IR to nvidia PTX.

Digging a bit more I found this page: 
http://steckdenis.be/post-2011-05-02-using-clang-to-compile-opencl-kernels.html

Following it's examples show me:
camibar% clang -v
FreeBSD clang version 3.3 (tags/RELEASE_33/final 183502) 20130610
Target: x86_64-unknown-freebsd9.2
Thread model: posix
%mkdir opencl
%cd opencl

create file test.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];
}

%clang -S -emit-llvm -o test.ll -x cl test.cl
%ls
test.cl test.ll
%more test.ll
; ModuleID = 'test.cl'
target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v64:64:64-v128:128:128-a0:0:64-s0:64:64-f80:128:128-n8:16:32:64-S128"
target triple = "x86_64-unknown-freebsd9.2"

; Function Attrs: nounwind uwtable
define void @test(float addrspace(1)* nocapture %in, float addrspace(1)* nocapture %out) #0 {
entry:
  %call = tail call i32 @get_global_id(i32 0) #3
  %idxprom = sext i32 %call to i64
  %arrayidx = getelementptr inbounds float addrspace(1)* %in, i64 %idxprom
  %0 = load float addrspace(1)* %arrayidx, align 4, !tbaa !1
...

So, clang3.3 in FreeBSD can compile opencl code to llvm assm (llvm IR). 

Will dig more when I'll have more time. 

Feel free to expand this mail thread.


---   ---
Eduardo Morras <emorrasg at yahoo.es>