extra/opencl/opencl-tests.factor

   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 ;
   6 FROM: alien.c-types => float ;
   7 SPECIALIZED-ARRAY: float
   8 IN: opencl.tests
   9
  10 STRING: kernel-source
  11 __kernel void square(
  12     __global float* input,
  13     __global float* output,
  14     const unsigned int count)
  15 {
  16     int i = get_global_id(0);
  17     if (i < count)
  18         output[i] = input[i] * input[i];
  19 }
  20 ;
  21
  22 :: opencl-square ( in -- out )
  23     [
  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
  29
  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
  34
  35             kernel in-buffer out-buffer num-floats <uint> 3array
  36             { num-floats } [ ] cl-queue-kernel &dispose drop
  37
  38             cl-finish
  39             out-buffer 0 num-bytes <cl-buffer-range> cl-read-buffer num-floats <direct-float-array>
  40         ] with-cl-state
  41     ] with-destructors ;
  42
  43 [ float-array{ 1.0 4.0 9.0 16.0 100.0 } ]
  44 [ float-array{ 1.0 2.0 3.0 4.0 10.0 } opencl-square ] unit-test