--- OpenCL/OpenCL.xs	2011/11/15 20:38:07	1.3
+++ OpenCL/OpenCL.xs	2011/11/15 21:13:42	1.4
@@ -105,6 +105,22 @@
    croak ("%s: %s is not of type %s", func, svname, pkg);
 }
 
+static void *
+tmpbuf (size_t size)
+{
+  static void *buf;
+  static size_t len;
+
+  if (len < size)
+    {
+      free (buf);
+      len = ((size + 31) & ~4095) + 4096 - 32;
+      buf = malloc (len);
+    }
+
+  return buf;
+}
+
 MODULE = OpenCL		PACKAGE = OpenCL
 
 PROTOTYPES: ENABLE
@@ -144,14 +160,12 @@
         int i;
 
 	NEED_SUCCESS (GetPlatformIDs, (0, 0, &count));
-        Newx (list, count, cl_platform_id);
+        list = tmpbuf (sizeof (*list) * count);
 	NEED_SUCCESS (GetPlatformIDs, (count, list, 0));
 
         EXTEND (SP, count);
         for (i = 0; i < count; ++i)
           PUSHs (NEW_MORTAL_OBJ ("OpenCL::Platform", list [i]));
-
-        Safefree (list);
 }
 
 void
@@ -193,14 +207,12 @@
         int i;
 
 	NEED_SUCCESS (GetDeviceIDs, (this, type, 0, 0, &count));
-        Newx (list, count, cl_device_id);
+        list = tmpbuf (sizeof (*list) * count);
 	NEED_SUCCESS (GetDeviceIDs, (this, type, count, list, 0));
 
         EXTEND (SP, count);
         for (i = 0; i < count; ++i)
           PUSHs (sv_setref_pv (sv_newmortal (), "OpenCL::Device", list [i]));
-
-        Safefree (list);
 }
 
 void
@@ -560,7 +572,59 @@
           XPUSH_NEW_OBJ ("OpenCL::Event", ev);
 }
 
- /*TODO http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueNDRangeKernel.html */
+void
+enqueue_nd_range_kernel (OpenCL::Queue this, OpenCL::Kernel kernel, SV *global_work_offset, SV *global_work_size, SV *local_work_size = &PL_sv_undef, ...)
+	PPCODE:
+{
+	cl_event ev = 0;
+        size_t *gwo = 0, *gws, *lws = 0;
+        int gws_len;
+        size_t *lists;
+        int i;
+        EVENT_LIST (2, items - 2);
+
+        if (!SvROK (global_work_size) || SvTYPE (SvRV (global_work_size)) != SVt_PVAV)
+          croak ("clEnqueueNDRangeKernel: global_work_size must be an array reference");
+
+        gws_len = AvFILLp (SvRV (global_work_size)) + 1;
+
+        lists = tmpbuf (sizeof (size_t) * 3 * gws_len);
+
+        gws = lists + gws_len * 0;
+        for (i = 0; i < gws_len; ++i)
+          gws [i] = SvIV (AvARRAY (SvRV (global_work_size))[i]);
+
+        if (SvOK (global_work_offset))
+          {
+            if (!SvROK (global_work_offset) || SvTYPE (SvRV (global_work_offset)) != SVt_PVAV)
+              croak ("clEnqueueNDRangeKernel: global_work_offset must be undef or an array reference");
+
+            if (AvFILLp (SvRV (global_work_size)) + 1 != gws_len)
+              croak ("clEnqueueNDRangeKernel: global_work_offset must be undef or an array of same size as global_work_size");
+
+            gwo = lists + gws_len * 1;
+            for (i = 0; i < gws_len; ++i)
+              gwo [i] = SvIV (AvARRAY (SvRV (global_work_offset))[i]);
+          }
+
+        if (SvOK (local_work_size))
+          {
+            if (SvOK (local_work_size) && !SvROK (local_work_size) || SvTYPE (SvRV (local_work_size)) != SVt_PVAV)
+              croak ("clEnqueueNDRangeKernel: global_work_size must be undef or an array reference");
+
+            if (AvFILLp (SvRV (local_work_size)) + 1 != gws_len)
+              croak ("clEnqueueNDRangeKernel: local_work_local must be undef or an array of same size as global_work_size");
+
+            lws = lists + gws_len * 2;
+            for (i = 0; i < gws_len; ++i)
+              lws [i] = SvIV (AvARRAY (SvRV (local_work_size))[i]);
+          }
+
+        NEED_SUCCESS (EnqueueNDRangeKernel, (this, kernel, gws_len, gwo, gws, lws, event_list_count, event_list_ptr, GIMME_V != G_VOID ? &ev : 0));
+
+        if (ev)
+          XPUSH_NEW_OBJ ("OpenCL::Event", ev);
+}
 
 void
 enqueue_marker (OpenCL::Queue this)