ScalaCL: added customCode function to allow for execution of hand-written OpenCL kernels

Author: ochafik

Core/src/main/velocity/com/nativelibs4java/opencl/CLContext.java

@@ -206,7 +206,7 @@ public String toString() {
 	public CLQueue createDefaultOutOfOrderQueueIfPossible() {
     		try {
     			return createDefaultOutOfOrderQueue();
-    		} catch (CLException.InvalidQueueProperties ex) {
+    		} catch (Throwable th) {//CLException.InvalidQueueProperties ex) {
     			return createDefaultQueue();
     		}
     }

ScalaCL/src/main/scala/scalacl/ScalaCL.scala

@@ -12,7 +12,7 @@ package scalacl {
     override def canEqual(that: Any) = getClass.isInstance(that.asInstanceOf[AnyRef])
   }
   class Context(val context: CLContext, val queue: CLQueue) extends AbstractProduct {
-    def this(context: CLContext) = this(context, context.createDefaultOutOfOrderQueueIfPossible())
+    def this(context: CLContext) = this(context, context.createDefaultQueue())//createDefaultOutOfOrderQueueIfPossible())
     def this() = this(JavaCL.createBestContext(CLPlatform.DeviceFeature.OutOfOrderQueueSupport, CLPlatform.DeviceFeature.MaxComputeUnits))
 
     def release = {
@@ -74,6 +74,18 @@ package object scalacl {
   val OutOfOrderQueueSupport = CLPlatform.DeviceFeature.OutOfOrderQueueSupport
   val MostImageFormats = CLPlatform.DeviceFeature.MostImageFormats
 
+  def customCode(
+    source: String,
+    compilerArguments: Array[String] = Array(),
+    macros: Map[String, String] = Map()
+  ): CLCode = {
+    new CLSimpleCode(
+      Array(source),
+      compilerArguments,
+      macros
+    )
+  }
+  
   private[scalacl] def reuse[T](value: Any, create: => T): T =
     if (value != null && value.isInstanceOf[T])
       value.asInstanceOf[T]

ScalaCL/src/main/scala/scalacl/impl/CLCode.scala

@@ -66,4 +66,54 @@ trait CLCode {
 
   override def hashCode = hc
   override def equals(o: Any) = o.isInstanceOf[CLCode] && strs.equals(o.asInstanceOf[CLCode].strs)
+  
+  /**
+   * Execute the code using the provided arguments (may be CLArray instances, CLGuardedBuffer instances, CLBuffer instances or primitive values), using the provided global sizes (and optional local group sizes) and respecting the execution order of arguments declared as "read from" and "written to".<br>
+   * Example :
+   * <code>
+     import scalacl._
+     implicit val context = Context.best(CPU)
+     val n = 100
+     val f = 0.5f
+     val sinCosOutputs: CLArray[Float] = new CLArray[Float](2 * n)
+     val sinCosCode = customCode("""
+       __kernel void sinCos(__global float2* outputs, float f) {
+        int i = get_global_id(0);
+        float c, s = sincos(i * f, &c);
+        outputs[i] = (float2)(s, c);
+       }
+     """)
+     sinCosCode.execute(
+       args = Array(sinCosOutputs, f),
+       writes = Array(sinCosOutputs),
+       globalSizes = Array(n)
+     )
+     val resCL = sinCosOutputs.toArray
+   * </code>
+   */
+  def execute(
+    args: Array[Any],
+    globalSizes: Array[Int],
+    localSizes: Array[Int] = null,
+    reads: Array[CLEventBoundContainer] = Array(),
+    writes: Array[CLEventBoundContainer] = Array(),
+    kernelName: String = null
+  )(implicit context: Context): Unit = {
+    val flatArgs: Array[Object] = args.flatMap(_ match {
+      case b: CLArray[_] =>
+        b.buffers.map(_.buffer.asInstanceOf[Object])
+      case b: CLGuardedBuffer[_] =>
+        Array[Object](b.buffer)
+      case a =>
+        Array[Object](a.asInstanceOf[Object])
+    })
+    val kernel = getKernel(context, kernelName)
+    kernel.synchronized {
+      //println("flatArgs = " + flatArgs.map(a => a + ": " + a.getClass.getSimpleName).mkString(", "))
+      kernel.setArgs(flatArgs:_*)
+      CLEventBound.syncBlock(CLEventBound.flatten(reads), CLEventBound.flatten(writes), evts => {
+        kernel.enqueueNDRange(context.queue, globalSizes, localSizes, evts:_*)
+      })
+    }
+  }
 }

ScalaCL/src/main/scala/scalacl/impl/CLEventBound.scala

@@ -60,6 +60,9 @@ trait CLEventBound extends CLEventBoundContainer {
 }
 
 object CLEventBound {
+  def flatten(containers: Array[CLEventBoundContainer]) =
+    containers.flatMap(_.eventBoundComponents)
+    
   def syncBlock(reads: Array[CLEventBound], writes: Array[CLEventBound], action: Array[CLEvent] => CLEvent): CLEvent = {
 
     def recursiveSync(ebs: List[(CLEventBound, Boolean)], evts: ArrayBuilder[CLEvent]): CLEvent = {
@@ -94,6 +97,10 @@ object CLEventBound {
       lb += ((eb, false))
     for (eb <- writes)
       lb += ((eb, true))
-    recursiveSync(lb.result, Array.newBuilder[CLEvent])
+      
+    if (lb.isEmpty)
+      action(Array())
+    else
+      recursiveSync(lb.result, Array.newBuilder[CLEvent])
   }
 }

ScalaCL/src/main/scala/scalacl/impl/CLFunction.scala

@@ -17,8 +17,8 @@ trait CLRunnable {
     if (dims.sum > 0) {
       lazy val defaultContainers = args collect { case c: CLEventBoundContainer => c }
       CLEventBound.syncBlock(
-        Option(reads).getOrElse(defaultContainers).flatMap(_.eventBoundComponents), 
-        Option(writes).getOrElse(defaultContainers).flatMap(_.eventBoundComponents), 
+        CLEventBound.flatten(Option(reads).getOrElse(defaultContainers)), 
+        CLEventBound.flatten(Option(writes).getOrElse(defaultContainers)), 
         evts => {
           run(dims = dims, args = args, eventsToWaitFor = evts)
         }

ScalaCL/src/test/scala/scalacl/CLCustomCodeTest.scala

@@ -0,0 +1,44 @@
+package scalacl
+
+import impl._
+
+import com.nativelibs4java.opencl._
+import org.bridj.Pointer
+import org.bridj.Pointer._
+
+import org.junit._
+import Assert._
+
+import scala.math._
+
+class CLCustomCodeTest {
+  
+  @Test
+  def testSinCos = {
+    import scalacl._
+    implicit val context = Context.best(CPU)
+    val n = 100
+    val f = 0.5f
+    val sinCosOutputs: CLArray[Float] = new CLArray[Float](2 * n)
+    val sinCosCode = customCode("""
+      __kernel void sinCos(__global float2* outputs, float f) {
+       int i = get_global_id(0);
+       float c, s = sincos(i * f, &c);
+       outputs[i] = (float2)(s, c);
+      }
+    """)
+    sinCosCode.execute(
+      args = Array(sinCosOutputs, f),
+      writes = Array(sinCosOutputs),
+      globalSizes = Array(n)
+    )
+    val resCL = sinCosOutputs.toArray
+    
+    val resJava = (0 until n).flatMap(i => {
+      val x = i * f
+      Array(sin(x).toFloat, cos(x).toFloat)
+    }).toArray
+    
+    assertArrayEquals(resJava, resCL, 0.00001f)
+  }
+}