diff --git a/README.textile b/README.textile
index 20367ce..4e293bc 100644
--- a/README.textile
+++ b/README.textile
@@ -6,7 +6,7 @@ The OpenCL API is not known for its terseness; consider this "'Hello World' exam
 	
 (def source
   "__kernel void square (
-       __global const float *a
+       __global const float *a,
        __global float *b) {
     int gid = get_global_id(0);
     b[gid] = a[gid] * a[gid];
@@ -15,7 +15,7 @@ The OpenCL API is not known for its terseness; consider this "'Hello World' exam
 (with-cl
   (with-program (compile-program source)
     (let [a (wrap [1 2 3] :float)
-          b (create-buffer 3 :float)]
+          b (mimic a)]
       (enqueue-kernel :square 3 a b)
       (enqueue-read b))))
 </code></pre>
diff --git a/src/calx/core.clj b/src/calx/core.clj
index fe49bd0..9935c6d 100644
--- a/src/calx/core.clj
+++ b/src/calx/core.clj
@@ -282,8 +282,9 @@
     (sequential? x) (into-array x)))
 
 (defn enqueue-kernel
-  ([kernel global-size block-size & args]
+  ([kernel global-size & args]
      (let [kernel ^CLKernel ((program) kernel)]
        (doseq [[idx arg] (indexed (map get-cl-object args))]
+	 (println idx arg)
 	 (.setArg kernel idx arg))
        (.enqueueNDRange kernel (queue) (to-dim-array global-size) *workgroup-size* (make-array CLEvent 0)))))
diff --git a/src/calx/data.clj b/src/calx/data.clj
index 9fb52c4..9f8f9ca 100644
--- a/src/calx/data.clj
+++ b/src/calx/data.clj
@@ -80,8 +80,8 @@
 ;;;
 
 (defprotocol Data
-  (mimic [d] [d dim]
-    "Create a different buffer of the same type. 'dim' defaults to the buffer's dimensions.")
+  (mimic [d] [d dim usage]
+    "Create a different buffer of the same type. 'dim' and 'usage' default to those of the original buffer.")
   (enqueue-read [d] [d range]
     "Asynchronously copies a subset of the buffer into local memory. 'range' defaults to the full buffer.
      Returns an object that, when dereferenced, halts execution until the copy is complete, then returns a seq.")
@@ -97,8 +97,13 @@
     ^{:cl-object cl-buf}
     (reify
       Data
-      (mimic [this] (mimic this dim))
-      (mimic [_ dim] (.createByteBuffer (context) (usage-types usage) (* dim element-bytes) false))
+      (mimic [this] (mimic this dim usage))
+      (mimic [_ dim usage]
+	(create-buffer-
+	  (.createByteBuffer (context) (usage-types usage) (* dim element-bytes))
+	  sig
+	  dim
+	  usage))
       (signature [_] sig)
       (dim [_] dim)
       (enqueue-read [this] (enqueue-read this (interval 0 dim)))