From d917334beda8e1a7b73ffe09f9252028747f5833 Mon Sep 17 00:00:00 2001
From: Olivier Chafik <olivier.chafik@gmail.com>
Date: Sun, 6 Oct 2013 20:28:28 +0200
Subject: [PATCH] Clear buffers upon creation (not needed on Apple CPU
 implementation, needed on GPU implementations, issue #3)

---
 .../scala/scalacl/impl/ScheduledBuffer.scala  | 31 ++++++++++++++-----
 1 file changed, 24 insertions(+), 7 deletions(-)

diff --git a/src/main/scala/scalacl/impl/ScheduledBuffer.scala b/src/main/scala/scalacl/impl/ScheduledBuffer.scala
index 3441cd3..d45403a 100644
--- a/src/main/scala/scalacl/impl/ScheduledBuffer.scala
+++ b/src/main/scala/scalacl/impl/ScheduledBuffer.scala
@@ -45,6 +45,18 @@ ScheduledBuffer:
 
  */
 
+object ScheduledBuffer {
+  private val clearBytesKernel = new Kernel(
+    """
+    kernel void f(global char* buffer, long length) {
+      size_t i = get_global_id(0);
+      if (i >= length) return;
+      buffer[i] = 0;
+    }
+    """
+  )
+}
+
 private[scalacl] class ScheduledBuffer[T](initialBuffer: CLBuffer[T], clearBuffer: Boolean = true)(implicit context: Context) extends DefaultScheduledData {
 
   private var buffer_ = initialBuffer
@@ -60,13 +72,18 @@ private[scalacl] class ScheduledBuffer[T](initialBuffer: CLBuffer[T], clearBuffe
   def buffer = buffer_
 
   def clear() = {
-    val byteLength = buffer_.getByteCount
-    kernel {
-      // TODO: optimize clear
-      for (i <- 0L until byteLength) {
-        this(i) = 0.asInstanceOf[T]
-      }
-    }
+    val b = buffer
+    val byteCount = b.getByteCount
+    ScheduledData.schedule(
+      Array(),
+      Array(this),
+      ScheduledBuffer.clearBytesKernel.enqueue(
+        context,
+        KernelExecutionParameters(globalSizes = Array(byteCount)),
+        args = Array(b, byteCount.asInstanceOf[AnyRef]),
+        _
+      )
+    )
   }
 
   def apply(index: Long): T = ???