1 ! Copyright (C) 2010 Erik Charlebois.
2 ! See http://factorcode.org/license.txt for BSD license.
3 USING: multiline locals io.encodings.ascii io.encodings.string sequences
4 math specialized-arrays alien.c-types math.order alien opencl tools.test
5 accessors arrays destructors kernel namespaces alien.data ;
6 FROM: alien.c-types => float ;
7 SPECIALIZED-ARRAY: float
12 __global float* input,
13 __global float* output,
14 const unsigned int count)
16 int i = get_global_id(0);
18 output[i] = input[i] * input[i];
22 :: opencl-square ( in -- out )
24 in byte-length :> num-bytes
25 in length :> num-floats
26 cl-platforms first devices>> first :> device
27 device 1array <cl-context> &dispose :> context
28 context device f f <cl-queue> &dispose :> queue
30 context device queue [
31 "" kernel-source 1array <cl-program> &dispose "square" <cl-kernel> &dispose :> kernel
32 cl-read-access num-bytes in <cl-buffer> &dispose :> in-buffer
33 cl-write-access num-bytes f <cl-buffer> &dispose :> out-buffer
35 kernel in-buffer out-buffer num-floats uint <ref> 3array
36 { num-floats } [ ] cl-queue-kernel &dispose drop
40 out-buffer 0 num-bytes <cl-buffer-range>
41 cl-read-buffer num-floats \ float <c-direct-array>
45 { float-array{ 1.0 4.0 9.0 16.0 100.0 } }
46 [ float-array{ 1.0 2.0 3.0 4.0 10.0 } opencl-square ] unit-test