Merge lp:~schnetter/pocl/main into lp:~pocl/pocl/trunk

[Erik Schnetter]

I made many improvements to the OpenCL run-time library, mostly to clGet* and buffer read/write/copy routines. The current state of the library is now sufficient to run my application code (although it is still much too slow).

=== modified file 'clconfig.h.in'
--- clconfig.h.in	2011-12-15 04:18:23 +0000
+++ clconfig.h.in	2011-12-18 04:51:24 +0000
@@ -1,13 +1,20 @@
+/* Define to 1 if __fp16 can be converted to float. */
+#undef CONVERT___FP16_FLOAT
+
+/* The size of `double', as computed by sizeof. */
+#undef SIZEOF_DOUBLE
+
 /* The size of `long', as computed by sizeof. */
 #undef SIZEOF_LONG
 
-/* The size of `half', as computed by sizeof. */
+/* The size of `void *', as computed by sizeof. */
+#undef SIZEOF_VOID_P
+
+/* The size of `__fp16', as computed by sizeof. */
 #undef SIZEOF___FP16
+
+
+
 /* The OpenCL type `half' is called `__fp16' in C */
 #define SIZEOF_HALF SIZEOF___FP16
-
-/* The size of `double', as computed by sizeof. */
-#undef SIZEOF_DOUBLE
-
-/* The size of `void *', as computed by sizeof. */
-#undef SIZEOF_VOID_P
+#define CONVERT_HALF_FLOAT CONVERT___FP16_FLOAT

=== modified file 'configure.ac'
--- configure.ac	2011-12-15 04:18:23 +0000
+++ configure.ac	2011-12-18 04:51:24 +0000
@@ -223,9 +223,22 @@
 # Checks for typedefs, structures, and compiler characteristics.
 AC_CHECK_SIZEOF([long])
 AC_CHECK_SIZEOF([__fp16])
+AC_MSG_CHECKING([whether __fp16 can be converted to float])
+AC_COMPILE_IFELSE([AC_LANG_PROGRAM([[__fp16 x; float y;]],
+                                   [[y=x]])],
+                  [
+                    AC_MSG_RESULT([yes])
+                    AC_DEFINE_UNQUOTED([CONVERT___FP16_FLOAT], 1, [Define to 1 if __fp16 can be converted to float.])
+                  ],
+                  [
+                    AC_MSG_RESULT([no])
+                    AC_DEFINE_UNQUOTED([CONVERT___FP16_FLOAT], 0, [Define to 1 if __fp16 can be converted to float.])
+                  ]
+                 )
 AC_CHECK_SIZEOF([double])
 AC_CHECK_SIZEOF([void *])
 AC_CHECK_ALIGNOF([float16], [typedef float float16  __attribute__((__ext_vector_type__(16)));])
+AC_CHECK_ALIGNOF([double16], [typedef double double16  __attribute__((__ext_vector_type__(16)));])
 
 CC="$old_CC"
 

=== modified file 'examples/kernel/test_bitselect.cl'
--- examples/kernel/test_bitselect.cl	2011-12-14 15:15:44 +0000
+++ examples/kernel/test_bitselect.cl	2011-12-18 04:51:24 +0000
@@ -1170,7 +1170,7 @@
  })
  )
 
-void test_bitselect()
+kernel void test_bitselect()
 {
   CALL_FUNC_G(test_bitselect)
 }

=== modified file 'examples/kernel/test_fabs.cl'
--- examples/kernel/test_fabs.cl	2011-12-14 15:15:44 +0000
+++ examples/kernel/test_fabs.cl	2011-12-18 04:51:24 +0000
@@ -172,7 +172,7 @@
  })
  )
 
-void test_fabs()
+kernel void test_fabs()
 {
   CALL_FUNC_V(test_fabs)
 }

=== modified file 'examples/kernel/test_rotate.cl'
--- examples/kernel/test_rotate.cl	2011-12-14 15:15:44 +0000
+++ examples/kernel/test_rotate.cl	2011-12-18 04:51:24 +0000
@@ -198,7 +198,7 @@
  })
  )
 
-void test_rotate()
+kernel void test_rotate()
 {
   CALL_FUNC_G(test_rotate)
 }

=== modified file 'examples/scalarwave/scalarwave.c'
--- examples/scalarwave/scalarwave.c	2011-12-14 15:45:14 +0000
+++ examples/scalarwave/scalarwave.c	2011-12-18 04:51:24 +0000
@@ -11,7 +11,7 @@
 
 
 
-#define GRID_GRANULARITY 1 // TODO 2
+#define GRID_GRANULARITY 2
 
 typedef struct grid_t {
   cl_double dt;                 // time step
@@ -96,7 +96,7 @@
   if (ierr) return -1;
   
   size_t const global_work_size[3] =
-    {grid->ni, grid->nj, grid->nk};
+    {grid->ai, grid->aj, grid->ak};
   size_t const local_work_size[3] =
     {GRID_GRANULARITY, GRID_GRANULARITY, GRID_GRANULARITY};
   
@@ -207,9 +207,8 @@
       phi = tmp;
     }
     
-    // TODO: We create the program and allocate the buffers each time,
-    // which is slow. But then, we only want to test correctness, not
-    // performance. (Yet?)
+    // TODO: We allocate the buffers each time, which is slow. But
+    // then, we only want to test correctness, not performance. (Yet?)
     int const ierr =
       exec_scalarwave_kernel (source, phi, phi_p, phi_p_p, &grid);
     assert(!ierr);

=== modified file 'examples/scalarwave/scalarwave.cl'
--- examples/scalarwave/scalarwave.cl	2011-12-14 15:45:14 +0000
+++ examples/scalarwave/scalarwave.cl	2011-12-18 04:51:24 +0000
@@ -21,7 +21,7 @@
   int ni, nj, nk;               // used size
 } grid_t;
 
-void
+kernel void
 scalarwave(global double       *restrict const phi,
            global double const *restrict const phi_p,
            global double const *restrict const phi_p_p,
@@ -72,11 +72,13 @@
   size_t const k = get_global_id(2);
   
   // If outside the domain, do nothing
-  if (i>=ni || j>=nj || k>=nk) return;
+  if (__builtin_expect(i>=ni || j>=nj || k>=nk, false)) return;
   
   size_t const ind3d = di*i + dj*j + dk*k;
   
-  if (i==0 || i==ni-1 || j==0 || j==nj-1 || k==0 || k==nk-1) {
+  if (__builtin_expect(i==0 || j==0 || k==0 || i==ni-1 || j==nj-1 || k==nk-1,
+                       false))
+  {
     // Boundary condition
     
     phi[ind3d] = 0.0;

=== modified file 'include/CL/cl.h'
--- include/CL/cl.h	2011-02-08 22:30:08 +0000
+++ include/CL/cl.h	2011-12-18 04:51:24 +0000
@@ -814,7 +814,7 @@
                          size_t              /* buffer_row_pitch */,
                          size_t              /* buffer_slice_pitch */,
                          size_t              /* host_row_pitch */,
-                         size_t              /* host_slice_pitch */,                        
+                         size_t              /* host_slice_pitch */,
                          const void *        /* ptr */,
                          cl_uint             /* num_events_in_wait_list */,
                          const cl_event *    /* event_wait_list */,

=== modified file 'include/_kernel.h'
--- include/_kernel.h	2011-12-15 19:03:41 +0000
+++ include/_kernel.h	2011-12-18 04:51:24 +0000
@@ -32,6 +32,9 @@
    only. Seems the pragma does not add the macro, so we have the target
    define the macro and the pragma is conditionally enabled.
 */
+#ifdef cl_khr_fp16
+#  pragma OPENCL EXTENSION cl_khr_fp16: enable
+#endif
 #ifdef cl_khr_fp64
 #  pragma OPENCL EXTENSION cl_khr_fp64: enable
 #endif

=== modified file 'include/arm/types.h'
--- include/arm/types.h	2011-12-14 23:01:01 +0000
+++ include/arm/types.h	2011-12-18 04:51:24 +0000
@@ -4,7 +4,7 @@
 
 #define __EMBEDDED_PROFILE__ 1
 #undef cles_khr_int64
-#define cl_khr_fp16             /* ES: is this correct? */
+#define cl_khr_fp16
 #undef cl_khr_fp64
 
 typedef uint size_t;

=== modified file 'include/tce/types.h'
--- include/tce/types.h	2011-12-14 23:01:01 +0000
+++ include/tce/types.h	2011-12-18 04:51:24 +0000
@@ -4,7 +4,7 @@
 
 #define __EMBEDDED_PROFILE__ 1
 #undef cles_khr_int64
-#define cl_khr_fp16             /* ES: is this correct? */
+#define cl_khr_fp16
 #undef cl_khr_fp64
 
 typedef uint size_t;

=== modified file 'include/types.h'
--- include/types.h	2011-12-14 23:01:01 +0000
+++ include/types.h	2011-12-18 04:51:24 +0000
@@ -18,7 +18,7 @@
 #  undef cles_khr_int64
 #endif
 
-#if SIZEOF_HALF == 2
+#if SIZEOF_HALF == 2 /* && CONVERT_HALF_FLOAT */
 #  define cl_khr_fp16
 #else
 #  undef cl_khr_fp16

=== modified file 'lib/CL/Makefile.am'
--- lib/CL/Makefile.am	2011-12-07 15:45:30 +0000
+++ lib/CL/Makefile.am	2011-12-18 04:51:24 +0000
@@ -31,6 +31,7 @@
                    clReleaseCommandQueue.c	\
                    clCreateBuffer.c		\
 	           clEnqueueReadBuffer.c	\
+	           clEnqueueReadBufferRect.c	\
                    clReleaseMemObject.c		\
                    clCreateProgramWithSource.c	\
                    clReleaseProgram.c		\
@@ -46,8 +47,14 @@
                    clCreateContext.c		\
                    clGetProgramBuildInfo.c	\
                    clEnqueueWriteBuffer.c	\
+                   clEnqueueWriteBufferRect.c	\
+                   clEnqueueCopyBuffer.c	\
+                   clEnqueueCopyBufferRect.c	\
                    clFinish.c			\
+                   clEnqueueBarrier.c		\
                    clGetKernelWorkGroupInfo.c	\
+                   clGetProgramInfo.c		\
+                   clGetEventProfilingInfo.c	\
                    pocl_cl.h			\
                    devices/devices.h		\
                    devices/devices.c		\

=== modified file 'lib/CL/clCreateBuffer.c'
--- lib/CL/clCreateBuffer.c	2011-12-08 02:45:48 +0000
+++ lib/CL/clCreateBuffer.c	2011-12-18 04:51:24 +0000
@@ -69,6 +69,8 @@
   mem->mem_host_ptr = host_ptr;
   mem->reference_count = 1;
   mem->context = context;
-      
+  
+  if (errcode_ret != NULL)
+    *errcode_ret = CL_SUCCESS;
   return mem;
 }

=== modified file 'lib/CL/clCreateCommandQueue.c'
--- lib/CL/clCreateCommandQueue.c	2011-10-14 10:31:27 +0000
+++ lib/CL/clCreateCommandQueue.c	2011-12-18 04:51:24 +0000
@@ -38,5 +38,7 @@
   command_queue->reference_count = 1;
   command_queue->properties = properties;
 
+  if (errcode_ret != NULL)
+    *errcode_ret = CL_SUCCESS;
   return command_queue;
 }

=== modified file 'lib/CL/clCreateContext.c'
--- lib/CL/clCreateContext.c	2011-12-07 15:00:07 +0000
+++ lib/CL/clCreateContext.c	2011-12-18 04:51:24 +0000
@@ -72,8 +72,7 @@
   context->properties = properties;
   context->reference_count = 1;
 
-  if (errcode_ret != NULL)
+  if (errcode_ret)
     *errcode_ret = CL_SUCCESS;
-
   return context;
 }

=== modified file 'lib/CL/clCreateContextFromType.c'
--- lib/CL/clCreateContextFromType.c	2011-12-02 17:27:49 +0000
+++ lib/CL/clCreateContextFromType.c	2011-12-18 04:51:24 +0000
@@ -68,5 +68,7 @@
   context->properties = properties;
   context->reference_count = 1;
 
+  if (errcode_ret != NULL)
+    *errcode_ret = CL_SUCCESS;
   return context;
 }

=== modified file 'lib/CL/clCreateKernel.c'
--- lib/CL/clCreateKernel.c	2011-12-14 18:52:39 +0000
+++ lib/CL/clCreateKernel.c	2011-12-18 04:51:24 +0000
@@ -145,5 +145,7 @@
   program->kernels = kernel;
   kernel->next = k;
 
+  if (errcode_ret != NULL)
+    *errcode_ret = CL_SUCCESS;
   return kernel;
 }

=== modified file 'lib/CL/clCreateProgramWithSource.c'
--- lib/CL/clCreateProgramWithSource.c	2011-10-14 10:31:27 +0000
+++ lib/CL/clCreateProgramWithSource.c	2011-12-18 04:51:24 +0000
@@ -92,5 +92,7 @@
   program->binary = NULL;
   program->kernels = NULL;
   
+  if (errcode_ret != NULL)
+    *errcode_ret = CL_SUCCESS;
   return program;
 }

=== added file 'lib/CL/clEnqueueBarrier.c'
--- lib/CL/clEnqueueBarrier.c	1970-01-01 00:00:00 +0000
+++ lib/CL/clEnqueueBarrier.c	2011-12-18 04:51:24 +0000
@@ -0,0 +1,31 @@
+/* OpenCL runtime library: clEnqueueBarrier()
+
+   Copyright (c) 2011 Erik Schnetter
+   
+   Permission is hereby granted, free of charge, to any person obtaining a copy
+   of this software and associated documentation files (the "Software"), to deal
+   in the Software without restriction, including without limitation the rights
+   to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+   copies of the Software, and to permit persons to whom the Software is
+   furnished to do so, subject to the following conditions:
+   
+   The above copyright notice and this permission notice shall be included in
+   all copies or substantial portions of the Software.
+   
+   THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+   IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+   FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+   AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+   LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+   OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+   THE SOFTWARE.
+*/
+
+#include "pocl_cl.h"
+
+CL_API_ENTRY cl_int CL_API_CALL
+clEnqueueBarrier(cl_command_queue command_queue) CL_API_SUFFIX__VERSION_1_0
+{
+  /* All operations are serialised, so we never have to wait */
+  return CL_SUCCESS;
+}

=== added file 'lib/CL/clEnqueueCopyBuffer.c'
--- lib/CL/clEnqueueCopyBuffer.c	1970-01-01 00:00:00 +0000
+++ lib/CL/clEnqueueCopyBuffer.c	2011-12-18 04:51:24 +0000
@@ -0,0 +1,67 @@
+/* OpenCL runtime library: clEnqueueCopyBuffer()
+
+   Copyright (c) 2011 Universidad Rey Juan Carlos
+   
+   Permission is hereby granted, free of charge, to any person obtaining a copy
+   of this software and associated documentation files (the "Software"), to deal
+   in the Software without restriction, including without limitation the rights
+   to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+   copies of the Software, and to permit persons to whom the Software is
+   furnished to do so, subject to the following conditions:
+   
+   The above copyright notice and this permission notice shall be included in
+   all copies or substantial portions of the Software.
+   
+   THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+   IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+   FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+   AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+   LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+   OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+   THE SOFTWARE.
+*/
+
+#include "pocl_cl.h"
+#include <assert.h>
+
+CL_API_ENTRY cl_int CL_API_CALL
+clEnqueueCopyBuffer(cl_command_queue command_queue,
+                    cl_mem src_buffer,
+                    cl_mem dst_buffer,
+                    size_t src_offset,
+                    size_t dst_offset,
+                    size_t cb, 
+                    cl_uint num_events_in_wait_list,
+                    const cl_event *event_wait_list,
+                    cl_event *event) CL_API_SUFFIX__VERSION_1_0
+{
+  cl_device_id device_id;
+  unsigned i;
+
+  if (command_queue == NULL)
+    return CL_INVALID_COMMAND_QUEUE;
+
+  if ((src_buffer == NULL) || (dst_buffer == NULL))
+    return CL_INVALID_MEM_OBJECT;
+
+  if ((command_queue->context != src_buffer->context) ||
+      (command_queue->context != dst_buffer->context))
+    return CL_INVALID_CONTEXT;
+
+  if ((src_offset + cb > src_buffer->size) ||
+      (dst_offset + cb > dst_buffer->size))
+    return CL_INVALID_VALUE;
+
+  device_id = command_queue->device;
+  for (i = 0; i < command_queue->context->num_devices; ++i)
+    {
+      if (command_queue->context->devices[i] == device_id)
+        break;
+    }
+
+  assert(i < command_queue->context->num_devices);
+
+  device_id->copy(device_id->data, src_buffer->device_ptrs[i], dst_buffer->device_ptrs[i], cb);
+
+  return CL_SUCCESS;
+}

=== added file 'lib/CL/clEnqueueCopyBufferRect.c'
--- lib/CL/clEnqueueCopyBufferRect.c	1970-01-01 00:00:00 +0000
+++ lib/CL/clEnqueueCopyBufferRect.c	2011-12-18 04:51:24 +0000
@@ -0,0 +1,87 @@
+/* OpenCL runtime library: clEnqueueCopyBufferRect()
+
+   Copyright (c) 2011 Universidad Rey Juan Carlos
+   
+   Permission is hereby granted, free of charge, to any person obtaining a copy
+   of this software and associated documentation files (the "Software"), to deal
+   in the Software without restriction, including without limitation the rights
+   to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+   copies of the Software, and to permit persons to whom the Software is
+   furnished to do so, subject to the following conditions:
+   
+   The above copyright notice and this permission notice shall be included in
+   all copies or substantial portions of the Software.
+   
+   THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+   IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+   FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+   AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+   LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+   OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+   THE SOFTWARE.
+*/
+
+#include "pocl_cl.h"
+#include <assert.h>
+
+CL_API_ENTRY cl_int CL_API_CALL
+clEnqueueCopyBufferRect(cl_command_queue command_queue,
+                        cl_mem src_buffer,
+                        cl_mem dst_buffer,
+                        const size_t *src_origin,
+                        const size_t *dst_origin, 
+                        const size_t *region,
+                        size_t src_row_pitch,
+                        size_t src_slice_pitch,
+                        size_t dst_row_pitch,
+                        size_t dst_slice_pitch,
+                        cl_uint num_events_in_wait_list,
+                        const cl_event *event_wait_list,
+                        cl_event *event) CL_API_SUFFIX__VERSION_1_1
+{
+  cl_device_id device_id;
+  unsigned i;
+
+  if (command_queue == NULL)
+    return CL_INVALID_COMMAND_QUEUE;
+
+  if ((src_buffer == NULL) || (dst_buffer == NULL))
+    return CL_INVALID_MEM_OBJECT;
+
+  if ((command_queue->context != src_buffer->context) ||
+      (command_queue->context != dst_buffer->context))
+    return CL_INVALID_CONTEXT;
+
+  if ((src_origin == NULL) ||
+      (dst_origin == NULL) ||
+      (region == NULL))
+    return CL_INVALID_VALUE;
+
+  if ((region[0]*region[1]*region[2] > 0) &&
+      (src_origin[0] + region[0]-1 +
+       src_row_pitch * (src_origin[1] + region[1]-1) +
+       src_slice_pitch * (src_origin[2] + region[2]-1) >= src_buffer->size))
+    return CL_INVALID_VALUE;
+  if ((region[0]*region[1]*region[2] > 0) &&
+      (dst_origin[0] + region[0]-1 +
+       dst_row_pitch * (dst_origin[1] + region[1]-1) +
+       dst_slice_pitch * (dst_origin[2] + region[2]-1) >= dst_buffer->size))
+    return CL_INVALID_VALUE;
+
+  device_id = command_queue->device;
+  for (i = 0; i < command_queue->context->num_devices; ++i)
+    {
+      if (command_queue->context->devices[i] == device_id)
+        break;
+    }
+
+  assert(i < command_queue->context->num_devices);
+
+  device_id->copy_rect(device_id->data,
+                       src_buffer->device_ptrs[i], dst_buffer->device_ptrs[i],
+                       src_origin, dst_origin, region,
+                       src_row_pitch, src_slice_pitch,
+                       dst_row_pitch, dst_slice_pitch);
+
+  return CL_SUCCESS;
+}

=== modified file 'lib/CL/clEnqueueReadBuffer.c'
--- lib/CL/clEnqueueReadBuffer.c	2011-11-30 16:00:41 +0000
+++ lib/CL/clEnqueueReadBuffer.c	2011-12-18 04:51:24 +0000
@@ -1,4 +1,4 @@
-/* OpenCL runtime library: clEnqueueWriteBuffer()
+/* OpenCL runtime library: clEnqueueReadBuffer()
 
    Copyright (c) 2011 Universidad Rey Juan Carlos
    

=== added file 'lib/CL/clEnqueueReadBufferRect.c'
--- lib/CL/clEnqueueReadBufferRect.c	1970-01-01 00:00:00 +0000
+++ lib/CL/clEnqueueReadBufferRect.c	2011-12-18 04:51:24 +0000
@@ -0,0 +1,82 @@
+/* OpenCL runtime library: clEnqueueReadBufferRect()
+
+   Copyright (c) 2011 Universidad Rey Juan Carlos
+   
+   Permission is hereby granted, free of charge, to any person obtaining a copy
+   of this software and associated documentation files (the "Software"), to deal
+   in the Software without restriction, including without limitation the rights
+   to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+   copies of the Software, and to permit persons to whom the Software is
+   furnished to do so, subject to the following conditions:
+   
+   The above copyright notice and this permission notice shall be included in
+   all copies or substantial portions of the Software.
+   
+   THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+   IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+   FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+   AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+   LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+   OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+   THE SOFTWARE.
+*/
+
+#include "pocl_cl.h"
+#include <assert.h>
+
+CL_API_ENTRY cl_int CL_API_CALL
+clEnqueueReadBufferRect(cl_command_queue command_queue,
+                        cl_mem buffer,
+                        cl_bool blocking_read,
+                        const size_t *buffer_origin,
+                        const size_t *host_origin,
+                        const size_t *region,
+                        size_t buffer_row_pitch,
+                        size_t buffer_slice_pitch,
+                        size_t host_row_pitch,
+                        size_t host_slice_pitch,
+                        void *ptr,
+                        cl_uint num_events_in_wait_list,
+                        const cl_event *event_wait_list,
+                        cl_event *event) CL_API_SUFFIX__VERSION_1_1
+{
+  cl_device_id device_id;
+  unsigned i;
+
+  if (command_queue == NULL)
+    return CL_INVALID_COMMAND_QUEUE;
+
+  if (buffer == NULL)
+    return CL_INVALID_MEM_OBJECT;
+
+  if (command_queue->context != buffer->context)
+    return CL_INVALID_CONTEXT;
+
+  if ((ptr == NULL) ||
+      (buffer_origin == NULL) ||
+      (host_origin == NULL) ||
+      (region == NULL))
+    return CL_INVALID_VALUE;
+  
+  if ((region[0]*region[1]*region[2] > 0) &&
+      (buffer_origin[0] + region[0]-1 +
+       buffer_row_pitch * (buffer_origin[1] + region[1]-1) +
+       buffer_slice_pitch * (buffer_origin[2] + region[2]-1) >= buffer->size))
+    return CL_INVALID_VALUE;
+
+  device_id = command_queue->device;
+  for (i = 0; i < command_queue->context->num_devices; ++i)
+    {
+      if (command_queue->context->devices[i] == device_id)
+        break;
+    }
+
+  assert(i < command_queue->context->num_devices);
+
+  device_id->read_rect(device_id->data, ptr, buffer->device_ptrs[i],
+                       buffer_origin, host_origin, region,
+                       buffer_row_pitch, buffer_slice_pitch,
+                       host_row_pitch, host_slice_pitch);
+
+  return CL_SUCCESS;
+}

=== modified file 'lib/CL/clEnqueueWriteBuffer.c'
--- lib/CL/clEnqueueWriteBuffer.c	2011-11-30 16:00:41 +0000
+++ lib/CL/clEnqueueWriteBuffer.c	2011-12-18 04:51:24 +0000
@@ -1,4 +1,4 @@
-/* OpenCL runtime library: clEnqueueReadBuffer()
+/* OpenCL runtime library: clEnqueueWriteBuffer()
 
    Copyright (c) 2011 Universidad Rey Juan Carlos
    

=== added file 'lib/CL/clEnqueueWriteBufferRect.c'
--- lib/CL/clEnqueueWriteBufferRect.c	1970-01-01 00:00:00 +0000
+++ lib/CL/clEnqueueWriteBufferRect.c	2011-12-18 04:51:24 +0000
@@ -0,0 +1,97 @@
+/* OpenCL runtime library: clEnqueueWriteBufferRect()
+
+   Copyright (c) 2011 Universidad Rey Juan Carlos
+   
+   Permission is hereby granted, free of charge, to any person obtaining a copy
+   of this software and associated documentation files (the "Software"), to deal
+   in the Software without restriction, including without limitation the rights
+   to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+   copies of the Software, and to permit persons to whom the Software is
+   furnished to do so, subject to the following conditions:
+   
+   The above copyright notice and this permission notice shall be included in
+   all copies or substantial portions of the Software.
+   
+   THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+   IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+   FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+   AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+   LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+   OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+   THE SOFTWARE.
+*/
+
+#include "pocl_cl.h"
+#include <assert.h>
+
+CL_API_ENTRY cl_int CL_API_CALL
+clEnqueueWriteBufferRect(cl_command_queue command_queue,
+                         cl_mem buffer,
+                         cl_bool blocking_write,
+                         const size_t *buffer_origin,
+                         const size_t *host_origin, 
+                         const size_t *region,
+                         size_t buffer_row_pitch,
+                         size_t buffer_slice_pitch,
+                         size_t host_row_pitch,
+                         size_t host_slice_pitch,
+                         const void *ptr,
+                         cl_uint num_events_in_wait_list,
+                         const cl_event *event_wait_list,
+                         cl_event *event) CL_API_SUFFIX__VERSION_1_1
+{
+  cl_device_id device_id;
+  unsigned i;
+
+  if (command_queue == NULL)
+    return CL_INVALID_COMMAND_QUEUE;
+
+  if (buffer == NULL)
+    return CL_INVALID_MEM_OBJECT;
+
+  if (command_queue->context != buffer->context)
+    return CL_INVALID_CONTEXT;
+
+  if ((ptr == NULL) ||
+      (buffer_origin == NULL) ||
+      (host_origin == NULL) ||
+      (region == NULL))
+    return CL_INVALID_VALUE;
+  
+  if ((region[0]*region[1]*region[2] > 0) &&
+      (buffer_origin[0] + region[0]-1 +
+       buffer_row_pitch * (buffer_origin[1] + region[1]-1) +
+       buffer_slice_pitch * (buffer_origin[2] + region[2]-1) >= buffer->size))
+  {
+#warning "TODO"
+    printf("bo=[%d,%d,%d]\n"
+           "ho=[%d,%d,%d]\n"
+           "re=[%d,%d,%d]\n"
+           "bp=[,%d,%d]\n"
+           "hp=[,%d,%d]\n"
+           "bs=[%d]\n",
+           (int)buffer_origin[0], (int)buffer_origin[1], (int)buffer_origin[2],
+           (int)host_origin[0], (int)host_origin[1], (int)host_origin[2],
+           (int)region[0], (int)region[1], (int)region[2],
+           (int)buffer_row_pitch, (int)buffer_slice_pitch,
+           (int)host_row_pitch, (int)host_slice_pitch,
+           (int)buffer->size);
+    return CL_INVALID_VALUE;
+  }
+
+  device_id = command_queue->device;
+  for (i = 0; i < command_queue->context->num_devices; ++i)
+    {
+      if (command_queue->context->devices[i] == device_id)
+	break;
+    }
+
+  assert(i < command_queue->context->num_devices);
+
+  device_id->write_rect(device_id->data, ptr, buffer->device_ptrs[i],
+                        buffer_origin, host_origin, region,
+                        buffer_row_pitch, buffer_slice_pitch,
+                        host_row_pitch, host_slice_pitch);
+
+  return CL_SUCCESS;
+}

=== modified file 'lib/CL/clFinish.c'
--- lib/CL/clFinish.c	2011-12-05 22:08:57 +0000
+++ lib/CL/clFinish.c	2011-12-18 04:51:24 +0000
@@ -24,7 +24,7 @@
 #include "pocl_cl.h"
 
 CL_API_ENTRY cl_int CL_API_CALL
-clFinish(cl_command_queue command_queue)
+clFinish(cl_command_queue command_queue) CL_API_SUFFIX__VERSION_1_0
 {
   /* All operations are serialised, so we never have to wait */
   return CL_SUCCESS;

=== modified file 'lib/CL/clGetContextInfo.c'
--- lib/CL/clGetContextInfo.c	2011-10-14 10:31:27 +0000
+++ lib/CL/clGetContextInfo.c	2011-12-18 04:51:24 +0000
@@ -33,28 +33,20 @@
                  size_t *param_value_size_ret) CL_API_SUFFIX__VERSION_1_0
 {
   size_t value_size;
-
+  
   switch (param_name) {
-  default:
-    assert(0);
+    
   case CL_CONTEXT_DEVICES:
-    value_size = context->num_devices * sizeof(cl_device_id);
-  }
-
-  if (param_value_size_ret != NULL)
-    *param_value_size_ret = value_size;
-
-  if (param_value != NULL) {
-    if (value_size > param_value_size)
-      return CL_INVALID_VALUE;
-
-    switch (param_name) {
-    default:
-      assert(0);
-    case CL_CONTEXT_DEVICES:
-      memcpy(param_value, context->devices, value_size);
+    {
+      value_size = context->num_devices * sizeof(cl_device_id);
+      if (param_value != NULL) {
+        if (param_value_size < value_size)
+          return CL_INVALID_VALUE;
+        memcpy(param_value, context->devices, value_size);
+      }
+      if (param_value_size_ret != NULL)
+        *param_value_size_ret = value_size;
+      return CL_SUCCESS;
     }
   }
-  
-  return CL_SUCCESS;
 }

=== modified file 'lib/CL/clGetDeviceIDs.c'
--- lib/CL/clGetDeviceIDs.c	2011-12-06 03:00:30 +0000
+++ lib/CL/clGetDeviceIDs.c	2011-12-18 04:51:24 +0000
@@ -23,6 +23,7 @@
 
 #include "pocl_cl.h"
 #include "devices/devices.h"
+#include <string.h>
 
 /* Note: this is a kludge. This will require a thorough re-write when pocl
  * supports multiple devices
@@ -34,41 +35,35 @@
                cl_device_id *   devices, 
                cl_uint *        num_devices) CL_API_SUFFIX__VERSION_1_0
 {
-  int num = 0;
+  int num;
+  int i;
 
-  // TODO: OpenCL API specification allows implementation dependant behaviour
-  // if platform == NULL. Should we just allow for it?	
+  /* TODO: OpenCL API specification allows implementation dependent
+     behaviour if platform == NULL. Should we just allow it? */
   if (platform == NULL || ( platform->magic != 42 ))
     return CL_INVALID_PLATFORM;
-	
-  // Currently - POCL supports only the host device - i.e. a CPU
+  
+  /* Currently POCL supports only the host device, i.e. a CPU */
   if ((device_type & CL_DEVICE_TYPE_CPU) ||
       (device_type & CL_DEVICE_TYPE_DEFAULT))
     num = 1;
-  else if ((device_type | CL_DEVICE_TYPE_GPU) ||
-           (device_type | CL_DEVICE_TYPE_ACCELERATOR))
+  else if ((device_type & CL_DEVICE_TYPE_GPU) ||
+           (device_type & CL_DEVICE_TYPE_ACCELERATOR))
     num = 0;
   else
     return CL_INVALID_DEVICE_TYPE;
-
-	// no room for any response
-  if (devices == NULL && num_devices == NULL)
-    return CL_INVALID_VALUE;
-
-  // user forgot to allocate space for response
-  if (num_entries > 0 && devices == NULL )
-    return CL_INVALID_VALUE;
-
-	
+  
+  if (devices != NULL) {
+    if (num < num_entries)
+      return CL_INVALID_VALUE;
+    
+    for (i=0; i<num; ++i)
+      devices[i] = &pocl_devices[i];
+  }
+  
   if (num_devices != NULL)
     *num_devices = num;
-	
-  if (num_entries > 0 && devices!= NULL)
-	{
-      if (num)
-        devices[0] = &pocl_devices[0];
-	}
-
+  
   if (num > 0)
     return CL_SUCCESS;
   else

=== modified file 'lib/CL/clGetDeviceInfo.c'
--- lib/CL/clGetDeviceInfo.c	2011-12-02 13:09:02 +0000
+++ lib/CL/clGetDeviceInfo.c	2011-12-18 04:51:24 +0000
@@ -24,13 +24,175 @@
 #include "pocl_cl.h"
 
 CL_API_ENTRY cl_int CL_API_CALL
-clGetDeviceInfo(cl_device_id     device,
-                cl_device_info   param_name, 
-                size_t           param_value_size, 
-                void *           param_value,
-                size_t *         param_value_size_ret ) CL_API_SUFFIX__VERSION_1_0
+clGetDeviceInfo(cl_device_id   device,
+                cl_device_info param_name, 
+                size_t         param_value_size, 
+                void *         param_value,
+                size_t *       param_value_size_ret) CL_API_SUFFIX__VERSION_1_0
 {
-	// TODO: dig up the info
-	return CL_INVALID_VALUE;
+  switch (param_name)
+  {
+  case CL_DEVICE_TYPE:
+    {
+      size_t const value_size = sizeof(cl_device_type);
+      if (param_value)
+      {
+        if (param_value_size < value_size) return CL_INVALID_VALUE;
+        *(cl_device_type*)param_value = device->type;
+      }
+      if (param_value_size_ret)
+        *param_value_size_ret = value_size;
+      return CL_SUCCESS;
+    }
+    
+  case CL_DEVICE_VENDOR_ID                         : break;
+  case CL_DEVICE_MAX_COMPUTE_UNITS                 : break;
+  case CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS          : break;
+  case CL_DEVICE_MAX_WORK_GROUP_SIZE               : break;
+  case CL_DEVICE_MAX_WORK_ITEM_SIZES               : break;
+    
+  case CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR:
+    {
+      size_t const value_size = sizeof(cl_uint);
+      if (param_value)
+      {
+        if (param_value_size < value_size) return CL_INVALID_VALUE;
+        *(cl_uint*)param_value = device->preferred_vector_width_char;
+      }
+      if (param_value_size_ret)
+        *param_value_size_ret = value_size;
+      return CL_SUCCESS;
+    }
+    
+  case CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT:
+    {
+      size_t const value_size = sizeof(cl_uint);
+      if (param_value)
+      {
+        if (param_value_size < value_size) return CL_INVALID_VALUE;
+        *(cl_uint*)param_value = device->preferred_vector_width_short;
+      }
+      if (param_value_size_ret)
+        *param_value_size_ret = value_size;
+      return CL_SUCCESS;
+    }
+    
+  case CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT:
+    {
+      size_t const value_size = sizeof(cl_uint);
+      if (param_value)
+      {
+        if (param_value_size < value_size) return CL_INVALID_VALUE;
+        *(cl_uint*)param_value = device->preferred_vector_width_int;
+      }
+      if (param_value_size_ret)
+        *param_value_size_ret = value_size;
+      return CL_SUCCESS;
+    }
+    
+  case CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG:
+    {
+      size_t const value_size = sizeof(cl_uint);
+      if (param_value)
+      {
+        if (param_value_size < value_size) return CL_INVALID_VALUE;
+        *(cl_uint*)param_value = device->preferred_vector_width_long;
+      }
+      if (param_value_size_ret)
+        *param_value_size_ret = value_size;
+      return CL_SUCCESS;
+    }
+    
+  case CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT:
+    {
+      size_t const value_size = sizeof(cl_uint);
+      if (param_value)
+      {
+        if (param_value_size < value_size) return CL_INVALID_VALUE;
+        *(cl_uint*)param_value = device->preferred_vector_width_float;
+      }
+      if (param_value_size_ret)
+        *param_value_size_ret = value_size;
+      return CL_SUCCESS;
+    }
+    
+  case CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE:
+    {
+      size_t const value_size = sizeof(cl_uint);
+      if (param_value)
+      {
+        if (param_value_size < value_size) return CL_INVALID_VALUE;
+        *(cl_uint*)param_value = device->preferred_vector_width_double;
+      }
+      if (param_value_size_ret)
+        *param_value_size_ret = value_size;
+      return CL_SUCCESS;
+    }
+
+  case CL_DEVICE_MAX_CLOCK_FREQUENCY               : break;
+  case CL_DEVICE_ADDRESS_BITS                      : break;
+  case CL_DEVICE_MAX_READ_IMAGE_ARGS               : break;
+  case CL_DEVICE_MAX_WRITE_IMAGE_ARGS              : break;
+  case CL_DEVICE_MAX_MEM_ALLOC_SIZE                : break;
+  case CL_DEVICE_IMAGE2D_MAX_WIDTH                 : break;
+  case CL_DEVICE_IMAGE2D_MAX_HEIGHT                : break;
+  case CL_DEVICE_IMAGE3D_MAX_WIDTH                 : break;
+  case CL_DEVICE_IMAGE3D_MAX_HEIGHT                : break;
+  case CL_DEVICE_IMAGE3D_MAX_DEPTH                 : break;
+  case CL_DEVICE_IMAGE_SUPPORT                     : break;
+  case CL_DEVICE_MAX_PARAMETER_SIZE                : break;
+  case CL_DEVICE_MAX_SAMPLERS                      : break;
+  case CL_DEVICE_MEM_BASE_ADDR_ALIGN               : break;
+  case CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE          : break;
+  case CL_DEVICE_SINGLE_FP_CONFIG                  : break;
+  case CL_DEVICE_GLOBAL_MEM_CACHE_TYPE             : break;
+  case CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE         : break;
+  case CL_DEVICE_GLOBAL_MEM_CACHE_SIZE             : break;
+  case CL_DEVICE_GLOBAL_MEM_SIZE                   : break;
+  case CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE          : break;
+  case CL_DEVICE_MAX_CONSTANT_ARGS                 : break;
+  case CL_DEVICE_LOCAL_MEM_TYPE                    : break;
+  case CL_DEVICE_LOCAL_MEM_SIZE                    : break;
+  case CL_DEVICE_ERROR_CORRECTION_SUPPORT          : break;
+  case CL_DEVICE_PROFILING_TIMER_RESOLUTION        : break;
+  case CL_DEVICE_ENDIAN_LITTLE                     : break;
+  case CL_DEVICE_AVAILABLE                         : break;
+  case CL_DEVICE_COMPILER_AVAILABLE                : break;
+  case CL_DEVICE_EXECUTION_CAPABILITIES            : break;
+  case CL_DEVICE_QUEUE_PROPERTIES                  : break;
+    
+  case CL_DEVICE_NAME:
+    {
+      size_t const value_size = strlen(device->name) + 1;
+      if (param_value)
+      {
+        if (param_value_size < value_size) return CL_INVALID_VALUE;
+        memcpy(param_value, device->name, value_size);
+      }
+      if (param_value_size_ret)
+        *param_value_size_ret = value_size;
+      return CL_SUCCESS;
+    }
+    
+  case CL_DEVICE_VENDOR                            : break;
+  case CL_DRIVER_VERSION                           : break;
+  case CL_DEVICE_PROFILE                           : break;
+  case CL_DEVICE_VERSION                           : break;
+  case CL_DEVICE_EXTENSIONS                        : break;
+  case CL_DEVICE_PLATFORM                          : break;
+  case CL_DEVICE_DOUBLE_FP_CONFIG                  : break;
+  case CL_DEVICE_HALF_FP_CONFIG                    : break;
+  case CL_DEVICE_PREFERRED_VECTOR_WIDTH_HALF       : break;
+  case CL_DEVICE_HOST_UNIFIED_MEMORY               : break;
+  case CL_DEVICE_NATIVE_VECTOR_WIDTH_CHAR          : break;
+  case CL_DEVICE_NATIVE_VECTOR_WIDTH_SHORT         : break;
+  case CL_DEVICE_NATIVE_VECTOR_WIDTH_INT           : break;
+  case CL_DEVICE_NATIVE_VECTOR_WIDTH_LONG          : break;
+  case CL_DEVICE_NATIVE_VECTOR_WIDTH_FLOAT         : break;
+  case CL_DEVICE_NATIVE_VECTOR_WIDTH_DOUBLE        : break;
+  case CL_DEVICE_NATIVE_VECTOR_WIDTH_HALF          : break;
+  case CL_DEVICE_OPENCL_C_VERSION                  : break;
+  }
+  
+  return CL_INVALID_VALUE;
 }
-

=== added file 'lib/CL/clGetEventProfilingInfo.c'
--- lib/CL/clGetEventProfilingInfo.c	1970-01-01 00:00:00 +0000
+++ lib/CL/clGetEventProfilingInfo.c	2011-12-18 04:51:24 +0000
@@ -0,0 +1,62 @@
+/* OpenCL runtime library: clGetEventProfilingInfo()
+
+   Copyright (c) 2011 Erik Schnetter
+   
+   Permission is hereby granted, free of charge, to any person obtaining a copy
+   of this software and associated documentation files (the "Software"), to deal
+   in the Software without restriction, including without limitation the rights
+   to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+   copies of the Software, and to permit persons to whom the Software is
+   furnished to do so, subject to the following conditions:
+   
+   The above copyright notice and this permission notice shall be included in
+   all copies or substantial portions of the Software.
+   
+   THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+   IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+   FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+   AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+   LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+   OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+   THE SOFTWARE.
+*/
+
+#include "pocl_cl.h"
+#include <string.h>
+
+CL_API_ENTRY cl_int CL_API_CALL
+clGetEventProfilingInfo(cl_event event,
+                        cl_profiling_info param_name,
+                        size_t param_value_size,
+                        void *param_value,
+                        size_t *param_value_size_ret) CL_API_SUFFIX__VERSION_1_0
+{
+  size_t const value_size = sizeof(cl_ulong);
+  if (param_value)
+  {
+    if (param_value_size < value_size) return CL_INVALID_VALUE;
+    
+    switch (param_name)
+    {
+    case CL_PROFILING_COMMAND_QUEUED:
+      *(cl_ulong*)param_value = 0; /* TODO: return correct value */
+      break;
+    case CL_PROFILING_COMMAND_SUBMIT:
+      *(cl_ulong*)param_value = 0; /* TODO: return correct value */
+      break;
+    case CL_PROFILING_COMMAND_START:
+      *(cl_ulong*)param_value = 0; /* TODO: return correct value */
+      break;
+    case CL_PROFILING_COMMAND_END:
+      *(cl_ulong*)param_value = 0; /* TODO: return correct value */
+      break;
+    default:
+      return CL_INVALID_VALUE;
+    }
+  }
+  
+  if (param_value_size_ret)
+    *param_value_size_ret = value_size;
+  
+  return CL_SUCCESS;
+}

=== modified file 'lib/CL/clGetPlatformIDs.c'
--- lib/CL/clGetPlatformIDs.c	2011-12-02 13:09:02 +0000
+++ lib/CL/clGetPlatformIDs.c	2011-12-18 04:51:24 +0000
@@ -36,27 +36,19 @@
                  cl_platform_id *  platforms,
                  cl_uint *         num_platforms) CL_API_SUFFIX__VERSION_1_0
 {	
-
-	// user requests only number of platforms - not types
-	if (num_entries == 0 && num_platforms != NULL)
-      {
-		*num_platforms = 1;
-		return CL_SUCCESS;
-      }
-
-	// Bad request - no place to store response
-	if (num_entries > 0 && platforms == NULL)
-      return CL_INVALID_VALUE;
-	
-	// Check required by spec
-	if (platforms == NULL && num_platforms == NULL)
-      return CL_INVALID_VALUE;
-	
-	// platform is not used now - just mark this platform as 'valid'
-	platforms[0] = &(_platforms[0]);
-	
-	if (num_platforms != NULL)
-      *num_platforms = 1;
-
-	return CL_SUCCESS;
+  int const num = 1;
+  int i;
+  
+  if (platforms != NULL) {
+    if (num_entries < num)
+      return CL_INVALID_VALUE;
+    
+    for (i=0; i<num; ++i)
+      platforms[i] = &_platforms[i];
+  }
+  
+  if (num_platforms != NULL)
+    *num_platforms = num;
+  
+  return CL_SUCCESS;
 }

=== modified file 'lib/CL/clGetPlatformInfo.c'
--- lib/CL/clGetPlatformInfo.c	2011-12-02 13:09:02 +0000
+++ lib/CL/clGetPlatformInfo.c	2011-12-18 04:51:24 +0000
@@ -31,14 +31,14 @@
                   void *           param_value,
                   size_t *         param_value_size_ret) CL_API_SUFFIX__VERSION_1_0
 {
-  char *ret=0;
-  int retlen;	
+  const char *ret;
+  int retlen;
 
   if (platform == NULL || (platform->magic != 42))
     return CL_INVALID_PLATFORM;
 	
   switch (param_name)
-	{
+  {
     case CL_PLATFORM_PROFILE:
       // TODO: figure this out depending on the native execution host.
       // assume FULL_PROFILE for now.
@@ -59,23 +59,21 @@
       break;
     default: 
       return CL_INVALID_VALUE;
-	}
+  }
 
-  // Specs say (section 4.1) to "ignore param_value" should it be NULL
-  if (param_value == NULL)
-    return CL_SUCCESS;	
-	
   // the OpenCL API docs *seem* to count the trailing NULL
   retlen = strlen(ret) + 1;
 	
-  if (param_value_size < retlen)
-    return CL_INVALID_VALUE;
-
-  strncpy(param_value, ret, retlen); 
+  if (param_value != NULL)
+  {
+    if (param_value_size < retlen)
+      return CL_INVALID_VALUE;
+    
+    memcpy(param_value, ret, retlen); 
+  }
 	
   if (param_value_size_ret != NULL)
-    *param_value_size_ret=retlen;
+    *param_value_size_ret = retlen;
 	
   return CL_SUCCESS;
-
 }

=== modified file 'lib/CL/clGetProgramBuildInfo.c'
--- lib/CL/clGetProgramBuildInfo.c	2011-12-02 13:09:02 +0000
+++ lib/CL/clGetProgramBuildInfo.c	2011-12-18 04:51:24 +0000
@@ -32,31 +32,48 @@
                       void *                param_value,
                       size_t *              param_value_size_ret) CL_API_SUFFIX__VERSION_1_0
 {
-  char *retval = "";
-  int retlen;
-
-  if (program == NULL)
-    return CL_INVALID_PROGRAM;
-
-  // currently - just stub-implmement this.
-  // there doesn't seem to exist an "CL_INTERNAL_ERROR" return code :(
-  if (param_name != CL_PROGRAM_BUILD_LOG)
-    return CL_INVALID_OPERATION;
-
-  retlen = strlen(retval) + 1;
-	
-  if (param_value == NULL)
-    return CL_SUCCESS; 	
-
-  if (param_value_size < retlen)
-    return CL_INVALID_VALUE;
-	
-  strncpy(param_value, retval, retlen);
-	
-  if (param_value_size_ret != NULL)
-    *param_value_size_ret = retlen;
-
-  return CL_SUCCESS;
+  const char *retval = "";      /* dummy return value */
+  
+  switch (param_name) {
+  case CL_PROGRAM_BUILD_STATUS:
+    {
+      size_t const value_size = strlen(retval) + 1;
+      if (param_value)
+      {
+        if (param_value_size < value_size) return CL_INVALID_VALUE;
+        memcpy(param_value, retval, value_size);
+      }
+      if (param_value_size_ret)
+        *param_value_size_ret = value_size;
+      return CL_SUCCESS;
+    }
+    
+  case CL_PROGRAM_BUILD_OPTIONS:
+    {
+      size_t const value_size = strlen(retval) + 1;
+      if (param_value)
+      {
+        if (param_value_size < value_size) return CL_INVALID_VALUE;
+        memcpy(param_value, retval, value_size);
+      }
+      if (param_value_size_ret)
+        *param_value_size_ret = value_size;
+      return CL_SUCCESS;
+    }
+    
+  case CL_PROGRAM_BUILD_LOG:
+    {
+      size_t const value_size = strlen(retval) + 1;
+      if (param_value)
+      {
+        if (param_value_size < value_size) return CL_INVALID_VALUE;
+        memcpy(param_value, retval, value_size);
+      }
+      if (param_value_size_ret)
+        *param_value_size_ret = value_size;
+      return CL_SUCCESS;
+    }
+  }
+  
+  return CL_INVALID_VALUE;
 }
-
-

=== added file 'lib/CL/clGetProgramInfo.c'
--- lib/CL/clGetProgramInfo.c	1970-01-01 00:00:00 +0000
+++ lib/CL/clGetProgramInfo.c	2011-12-18 04:51:24 +0000
@@ -0,0 +1,89 @@
+/* OpenCL runtime library: clGetProgramInfo()
+
+   Copyright (c) 2011 Erik Schnetter
+   
+   Permission is hereby granted, free of charge, to any person obtaining a copy
+   of this software and associated documentation files (the "Software"), to deal
+   in the Software without restriction, including without limitation the rights
+   to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+   copies of the Software, and to permit persons to whom the Software is
+   furnished to do so, subject to the following conditions:
+   
+   The above copyright notice and this permission notice shall be included in
+   all copies or substantial portions of the Software.
+   
+   THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+   IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+   FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+   AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+   LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+   OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+   THE SOFTWARE.
+*/
+
+#include "pocl_cl.h"
+#include <string.h>
+
+CL_API_ENTRY cl_int CL_API_CALL
+clGetProgramInfo(cl_program program,
+                 cl_program_info param_name,
+                 size_t param_value_size,
+                 void *param_value,
+                 size_t *param_value_size_ret) CL_API_SUFFIX__VERSION_1_0
+{
+  switch (param_name)
+  {
+  case CL_PROGRAM_REFERENCE_COUNT:
+    return CL_INVALID_VALUE;    /* not yet implemented */
+    
+  case CL_PROGRAM_CONTEXT:
+    return CL_INVALID_VALUE;    /* not yet implemented */
+    
+  case CL_PROGRAM_NUM_DEVICES:
+    return CL_INVALID_VALUE;    /* not yet implemented */
+    
+  case CL_PROGRAM_DEVICES:
+    return CL_INVALID_VALUE;    /* not yet implemented */
+    
+  case CL_PROGRAM_SOURCE:
+    {
+      size_t const value_size = strlen(program->source) + 1;
+      if (param_value)
+      {
+        if (param_value_size < value_size) return CL_INVALID_VALUE;
+        memcpy(param_value, program->source, value_size);
+      }
+      if (param_value_size_ret)
+        *param_value_size_ret = value_size;
+      return CL_SUCCESS;
+    }
+    
+  case CL_PROGRAM_BINARY_SIZES:
+    {
+      size_t const value_size = sizeof(size_t);
+      if (param_value)
+      {
+        if (param_value_size < value_size) return CL_INVALID_VALUE;
+        *(size_t*)param_value = program->binary_size;
+      }
+      if (param_value_size_ret)
+        *param_value_size_ret = value_size;
+      return CL_SUCCESS;
+    }
+    
+  case CL_PROGRAM_BINARIES:
+    {
+      size_t const value_size = sizeof(unsigned char *);
+      if (param_value)
+      {
+        if (param_value_size < value_size) return CL_INVALID_VALUE;
+        *(unsigned char **)param_value = program->binary;
+      }
+      if (param_value_size_ret)
+        *param_value_size_ret = value_size;
+      return CL_SUCCESS;
+    }
+  }
+  
+  return CL_INVALID_VALUE;
+}

=== modified file 'lib/CL/devices/native/native.c'
--- lib/CL/devices/native/native.c	2011-12-08 02:45:48 +0000
+++ lib/CL/devices/native/native.c	2011-12-18 04:51:24 +0000
@@ -27,9 +27,13 @@
 #include <stdlib.h>
 #include <unistd.h>
 
+#define max(a,b) (((a) > (b)) ? (a) : (b))
+
 #define COMMAND_LENGTH 256
 #define WORKGROUP_STRING_LENGTH 128
 
+#define ALIGNMENT (max(ALIGNOF_FLOAT16, ALIGNOF_DOUBLE16))
+
 struct data {
   /* Currently loaded kernel. */
   cl_kernel current_kernel;
@@ -60,7 +64,7 @@
 
   if (flags & CL_MEM_COPY_HOST_PTR)
     {
-      if (posix_memalign (&b, ALIGNOF_FLOAT16, size) == 0)
+      if (posix_memalign (&b, ALIGNMENT, size) == 0)
 	{
 	  memcpy (b, host_ptr, size);
 	  return b;
@@ -74,7 +78,7 @@
       return host_ptr;
     }
 
-  if (posix_memalign (&b, ALIGNOF_FLOAT16, size) == 0)
+  if (posix_memalign (&b, ALIGNMENT, size) == 0)
     return b;
 
   return NULL;

=== modified file 'lib/CL/devices/native/native.h'
--- lib/CL/devices/native/native.h	2011-12-08 02:45:48 +0000
+++ lib/CL/devices/native/native.h	2011-12-18 04:51:24 +0000
@@ -94,7 +94,11 @@
   pocl_native_malloc, /* malloc */					\
   pocl_native_free, /* free */						\
   pocl_native_read, /* read */						\
+  NULL, /* read_rect */                                                 \
   pocl_native_write, /* write */                                        \
+  NULL, /* write_rect */                                                \
+  NULL, /* copy */                                                      \
+  NULL, /* copy_rect */                                                 \
   pocl_native_run, /* run */						\
   NULL /* data */                                                       \
 }

=== modified file 'lib/CL/devices/pthread/pthread.c'
--- lib/CL/devices/pthread/pthread.c	2011-12-08 02:45:48 +0000
+++ lib/CL/devices/pthread/pthread.c	2011-12-18 04:51:24 +0000
@@ -29,10 +29,13 @@
 #include <unistd.h>
 
 #define min(a,b) (((a) < (b)) ? (a) : (b))
+#define max(a,b) (((a) > (b)) ? (a) : (b))
 
 #define COMMAND_LENGTH 256
 #define WORKGROUP_STRING_LENGTH 128
 
+#define ALIGNMENT (max(ALIGNOF_FLOAT16, ALIGNOF_DOUBLE16))
+
 /* The name of the environment variable used to force a certain max thread count
    for the thread execution. */
 #define THREAD_COUNT_ENV "POCL_MAX_PTHREAD_COUNT"
@@ -80,7 +83,7 @@
 
   if (flags & CL_MEM_COPY_HOST_PTR)
     {
-      if (posix_memalign (&b, ALIGNOF_FLOAT16, size) == 0)
+      if (posix_memalign (&b, ALIGNMENT, size) == 0)
 	{
 	  memcpy (b, host_ptr, size);
 	  return b;
@@ -94,7 +97,7 @@
       return host_ptr;
     }
 
-  if (posix_memalign (&b, ALIGNOF_FLOAT16, size) == 0)
+  if (posix_memalign (&b, ALIGNMENT, size) == 0)
     return b;
   
   return NULL;
@@ -118,7 +121,38 @@
   memcpy (host_ptr, device_ptr, cb);
 }
 
-void pocl_pthread_write (void *data, const void *host_ptr, void *device_ptr, size_t cb)
+void
+pocl_pthread_read_rect (void *data,
+                        void *__restrict__ const host_ptr,
+                        void *__restrict__ const device_ptr,
+                        const size_t *__restrict__ const buffer_origin,
+                        const size_t *__restrict__ const host_origin, 
+                        const size_t *__restrict__ const region,
+                        size_t const buffer_row_pitch,
+                        size_t const buffer_slice_pitch,
+                        size_t const host_row_pitch,
+                        size_t const host_slice_pitch)
+{
+  char const *__restrict const adjusted_device_ptr = 
+    (char const*)device_ptr +
+    buffer_origin[0] + buffer_row_pitch * (buffer_origin[1] + buffer_slice_pitch * buffer_origin[2]);
+  char *__restrict__ const adjusted_host_ptr = 
+    (char*)host_ptr +
+    host_origin[0] + host_row_pitch * (host_origin[1] + host_slice_pitch * host_origin[2]);
+  
+  size_t j, k;
+  
+  /* TODO: handle overlaping regions */
+  
+  for (k = 0; k < region[2]; ++k)
+    for (j = 0; j < region[1]; ++j)
+      memcpy (adjusted_host_ptr + host_row_pitch * j + host_slice_pitch * k,
+              adjusted_device_ptr + buffer_row_pitch * j + buffer_slice_pitch * k,
+              region[0]);
+}
+
+void
+pocl_pthread_write (void *data, const void *host_ptr, void *device_ptr, size_t cb)
 {
   if (host_ptr == device_ptr)
     return;
@@ -126,6 +160,75 @@
   memcpy (device_ptr, host_ptr, cb);
 }
 
+void
+pocl_pthread_write_rect (void *data,
+                         const void *__restrict__ const host_ptr,
+                         void *__restrict__ const device_ptr,
+                         const size_t *__restrict__ const buffer_origin,
+                         const size_t *__restrict__ const host_origin, 
+                         const size_t *__restrict__ const region,
+                         size_t const buffer_row_pitch,
+                         size_t const buffer_slice_pitch,
+                         size_t const host_row_pitch,
+                         size_t const host_slice_pitch)
+{
+  char *__restrict const adjusted_device_ptr = 
+    (char*)device_ptr +
+    buffer_origin[0] + buffer_row_pitch * (buffer_origin[1] + buffer_slice_pitch * buffer_origin[2]);
+  char const *__restrict__ const adjusted_host_ptr = 
+    (char const*)host_ptr +
+    host_origin[0] + host_row_pitch * (host_origin[1] + host_slice_pitch * host_origin[2]);
+  
+  size_t j, k;
+
+  /* TODO: handle overlaping regions */
+  
+  for (k = 0; k < region[2]; ++k)
+    for (j = 0; j < region[1]; ++j)
+      memcpy (adjusted_device_ptr + buffer_row_pitch * j + buffer_slice_pitch * k,
+              adjusted_host_ptr + host_row_pitch * j + host_slice_pitch * k,
+              region[0]);
+}
+
+void
+pocl_pthread_copy (void *data, const void *src_ptr, const void *dst_ptr, size_t cb)
+{
+  if (src_ptr == dst_ptr)
+    return;
+  
+  memcpy (dst_ptr, src_ptr, cb);
+}
+
+void
+pocl_pthread_copy_rect (void *data,
+                        const void *__restrict const src_ptr,
+                        void *__restrict__ const dst_ptr,
+                        const size_t *__restrict__ const src_origin,
+                        const size_t *__restrict__ const dst_origin, 
+                        const size_t *__restrict__ const region,
+                        size_t const src_row_pitch,
+                        size_t const src_slice_pitch,
+                        size_t const dst_row_pitch,
+                        size_t const dst_slice_pitch)
+{
+  char const *__restrict const adjusted_src_ptr = 
+    (char const*)src_ptr +
+    src_origin[0] + src_row_pitch * (src_origin[1] + src_slice_pitch * src_origin[2]);
+  char *__restrict__ const adjusted_dst_ptr = 
+    (char*)dst_ptr +
+    dst_origin[0] + dst_row_pitch * (dst_origin[1] + dst_slice_pitch * dst_origin[2]);
+  
+  size_t j, k;
+
+  /* TODO: handle overlaping regions */
+  
+  for (k = 0; k < region[2]; ++k)
+    for (j = 0; j < region[1]; ++j)
+      memcpy (adjusted_dst_ptr + dst_row_pitch * j + dst_slice_pitch * k,
+              adjusted_src_ptr + src_row_pitch * j + src_slice_pitch * k,
+              region[0]);
+}
+
 //#define DEBUG_MT
 //#define DEBUG_MAX_THREAD_COUNT
 /**

=== modified file 'lib/CL/devices/pthread/pthread.h'
--- lib/CL/devices/pthread/pthread.h	2011-12-08 02:45:48 +0000
+++ lib/CL/devices/pthread/pthread.h	2011-12-18 04:51:24 +0000
@@ -31,13 +31,82 @@
 			  size_t size, void *host_ptr);
 void pocl_pthread_free (void *data, cl_mem_flags flags, void *ptr);
 void pocl_pthread_read (void *data, void *host_ptr, const void *device_ptr, size_t cb);
+void pocl_pthread_read_rect (void *data, void *host_ptr, void *device_ptr,
+                             const size_t *buffer_origin,
+                             const size_t *host_origin, 
+                             const size_t *region,
+                             size_t buffer_row_pitch,
+                             size_t buffer_slice_pitch,
+                             size_t host_row_pitch,
+                             size_t host_slice_pitch);
 void pocl_pthread_write (void *data, const void *host_ptr, void *device_ptr, size_t cb);
+void pocl_pthread_write_rect (void *data, const void *host_ptr, void *device_ptr,
+                              const size_t *buffer_origin,
+                              const size_t *host_origin, 
+                              const size_t *region,
+                              size_t buffer_row_pitch,
+                              size_t buffer_slice_pitch,
+                              size_t host_row_pitch,
+                              size_t host_slice_pitch);
+void pocl_pthread_copy (void *data, const void *src_ptr, const void *dst_ptr, size_t cb);
+void pocl_pthread_copy_rect (void *data, const void *src_ptr, void *dst_ptr,
+                             const size_t *src_origin,
+                             const size_t *dst_origin, 
+                             const size_t *region,
+                             size_t src_row_pitch,
+                             size_t src_slice_pitch,
+                             size_t dst_row_pitch,
+                             size_t dst_slice_pitch);
 void pocl_pthread_run (void *data, const char *bytecode,
 		      cl_kernel kernel,
 		      struct pocl_context *pc);
 
 extern size_t pocl_pthread_max_work_item_sizes[];
 
+/* Determine preferred vector sizes */
+#if defined(__AVX__)
+#  define POCL_DEVICES_PTHREAD_PREFERRED_VECTOR_WIDTH_CHAR   16
+#  define POCL_DEVICES_PTHREAD_PREFERRED_VECTOR_WIDTH_SHORT   8
+#  define POCL_DEVICES_PTHREAD_PREFERRED_VECTOR_WIDTH_INT     4
+#  define POCL_DEVICES_PTHREAD_PREFERRED_VECTOR_WIDTH_LONG    2
+#  define POCL_DEVICES_PTHREAD_PREFERRED_VECTOR_WIDTH_FLOAT   4
+#  define POCL_DEVICES_PTHREAD_PREFERRED_VECTOR_WIDTH_DOUBLE  2
+#  define POCL_DEVICES_PTHREAD_NATIVE_VECTOR_WIDTH_CHAR      16
+#  define POCL_DEVICES_PTHREAD_NATIVE_VECTOR_WIDTH_SHORT      8
+#  define POCL_DEVICES_PTHREAD_NATIVE_VECTOR_WIDTH_INT        4
+#  define POCL_DEVICES_PTHREAD_NATIVE_VECTOR_WIDTH_LONG       2
+#  define POCL_DEVICES_PTHREAD_NATIVE_VECTOR_WIDTH_FLOAT      8
+#  define POCL_DEVICES_PTHREAD_NATIVE_VECTOR_WIDTH_DOUBLE     4
+#elif defined(__SSE2__)
+#  define POCL_DEVICES_PTHREAD_PREFERRED_VECTOR_WIDTH_CHAR   16
+#  define POCL_DEVICES_PTHREAD_PREFERRED_VECTOR_WIDTH_SHORT   8
+#  define POCL_DEVICES_PTHREAD_PREFERRED_VECTOR_WIDTH_INT     4
+#  define POCL_DEVICES_PTHREAD_PREFERRED_VECTOR_WIDTH_LONG    2
+#  define POCL_DEVICES_PTHREAD_PREFERRED_VECTOR_WIDTH_FLOAT   4
+#  define POCL_DEVICES_PTHREAD_PREFERRED_VECTOR_WIDTH_DOUBLE  2
+#  define POCL_DEVICES_PTHREAD_NATIVE_VECTOR_WIDTH_CHAR      16
+#  define POCL_DEVICES_PTHREAD_NATIVE_VECTOR_WIDTH_SHORT      8
+#  define POCL_DEVICES_PTHREAD_NATIVE_VECTOR_WIDTH_INT        4
+#  define POCL_DEVICES_PTHREAD_NATIVE_VECTOR_WIDTH_LONG       2
+#  define POCL_DEVICES_PTHREAD_NATIVE_VECTOR_WIDTH_FLOAT      4
+#  define POCL_DEVICES_PTHREAD_NATIVE_VECTOR_WIDTH_DOUBLE     2
+#else
+#  define POCL_DEVICES_PTHREAD_PREFERRED_VECTOR_WIDTH_CHAR    1
+#  define POCL_DEVICES_PTHREAD_PREFERRED_VECTOR_WIDTH_SHORT   1
+#  define POCL_DEVICES_PTHREAD_PREFERRED_VECTOR_WIDTH_INT     1
+#  define POCL_DEVICES_PTHREAD_PREFERRED_VECTOR_WIDTH_LONG    1
+#  define POCL_DEVICES_PTHREAD_PREFERRED_VECTOR_WIDTH_FLOAT   1
+#  define POCL_DEVICES_PTHREAD_PREFERRED_VECTOR_WIDTH_DOUBLE  1
+#  define POCL_DEVICES_PTHREAD_NATIVE_VECTOR_WIDTH_CHAR       1
+#  define POCL_DEVICES_PTHREAD_NATIVE_VECTOR_WIDTH_SHORT      1
+#  define POCL_DEVICES_PTHREAD_NATIVE_VECTOR_WIDTH_INT        1
+#  define POCL_DEVICES_PTHREAD_NATIVE_VECTOR_WIDTH_LONG       1
+#  define POCL_DEVICES_PTHREAD_NATIVE_VECTOR_WIDTH_FLOAT      1
+#  define POCL_DEVICES_PTHREAD_NATIVE_VECTOR_WIDTH_DOUBLE     1
+#endif
+/* Half is internally represented as short */
+#define POCL_DEVICES_PTHREAD_PREFERRED_VECTOR_WIDTH_HALF POCL_DEVICES_PTHREAD_PREFERRED_VECTOR_WIDTH_SHORT
+
 #define POCL_DEVICES_PTHREAD {						\
   CL_DEVICE_TYPE_CPU, /* type */					\
   0, /* vendor_id */							\
@@ -45,12 +114,12 @@
   3, /* max_work_item_dimensions */					\
   pocl_pthread_max_work_item_sizes, /* max_work_item_sizes */		\
   CL_INT_MAX, /* max_work_group_size */					\
-  0, /* preferred_vector_width_char */					\
-  0, /* preferred_vector_width_short */					\
-  0, /* preferred_vector_width_int */					\
-  0, /* preferred_vector_width_long */					\
-  0, /* preferred_vector_width_float */					\
-  0, /* preferred_vector_width_double */				\
+  POCL_DEVICES_PTHREAD_PREFERRED_VECTOR_WIDTH_CHAR  , /* preferred_vector_width_char */ \
+  POCL_DEVICES_PTHREAD_PREFERRED_VECTOR_WIDTH_SHORT , /* preferred_vector_width_short */ \
+  POCL_DEVICES_PTHREAD_PREFERRED_VECTOR_WIDTH_INT   , /* preferred_vector_width_int */ \
+  POCL_DEVICES_PTHREAD_PREFERRED_VECTOR_WIDTH_LONG  , /* preferred_vector_width_long */ \
+  POCL_DEVICES_PTHREAD_PREFERRED_VECTOR_WIDTH_FLOAT , /* preferred_vector_width_float */ \
+  POCL_DEVICES_PTHREAD_PREFERRED_VECTOR_WIDTH_DOUBLE, /* preferred_vector_width_double */ \
   0, /* max_clock_frequency */						\
   0, /* address_bits */							\
   0, /* max_mem_alloc_size */						\
@@ -94,7 +163,11 @@
   pocl_pthread_malloc, /* malloc */					\
   pocl_pthread_free, /* free */						\
   pocl_pthread_read, /* read */						\
+  pocl_pthread_read_rect, /* read_rect */				\
   pocl_pthread_write, /* write */					\
+  pocl_pthread_write_rect, /* write_rect */				\
+  pocl_pthread_copy, /* copy */						\
+  pocl_pthread_copy_rect, /* copy_rect */				\
   pocl_pthread_run, /* run */						\
   NULL /* data */							\
 }

=== modified file 'lib/CL/pocl_cl.h'
--- lib/CL/pocl_cl.h	2011-12-08 14:06:07 +0000
+++ lib/CL/pocl_cl.h	2011-12-18 04:51:24 +0000
@@ -100,7 +100,32 @@
 		   size_t size, void *host_ptr);
   void (*free) (void *data, cl_mem_flags flags, void *ptr);
   void (*read) (void *data, void *host_ptr, void *device_ptr, size_t cb);
+  void (*read_rect) (void *data, void *host_ptr, void *device_ptr,
+                     const size_t *buffer_origin,
+                     const size_t *host_origin, 
+                     const size_t *region,
+                     size_t buffer_row_pitch,
+                     size_t buffer_slice_pitch,
+                     size_t host_row_pitch,
+                     size_t host_slice_pitch);
   void (*write) (void *data, const void *host_ptr, void *device_ptr, size_t cb);
+  void (*write_rect) (void *data, const void *host_ptr, void *device_ptr,
+                      const size_t *buffer_origin,
+                      const size_t *host_origin, 
+                      const size_t *region,
+                      size_t buffer_row_pitch,
+                      size_t buffer_slice_pitch,
+                      size_t host_row_pitch,
+                      size_t host_slice_pitch);
+  void (*copy) (void *data, const void *src_ptr, const void *dst_ptr, size_t cb);
+  void (*copy_rect) (void *data, const void *src_ptr, void *dst_ptr,
+                     const size_t *src_origin,
+                     const size_t *dst_origin, 
+                     const size_t *region,
+                     size_t src_row_pitch,
+                     size_t src_slice_pitch,
+                     size_t dst_row_pitch,
+                     size_t dst_slice_pitch);
   void (*run) (void *data, const char *bytecode,
 	       cl_kernel kernel,
 	       struct pocl_context *pc);

=== modified file 'lib/kernel/all.cl'
--- lib/kernel/all.cl	2011-11-24 15:56:10 +0000
+++ lib/kernel/all.cl	2011-12-18 04:51:24 +0000
@@ -21,123 +21,123 @@
    THE SOFTWARE.
 */
 
-int __attribute__((__overloadable__)) all(char a)
+int _cl_overloadable all(char a)
 {
   return a < (char)0;
 }
 
-int __attribute__((__overloadable__)) all(char2 a)
+int _cl_overloadable all(char2 a)
 {
   return all(a.lo) && all(a.hi);
 }
 
-int __attribute__((__overloadable__)) all(char3 a)
+int _cl_overloadable all(char3 a)
 {
   return all(a.s01) && all(a.s2);
 }
 
-int __attribute__((__overloadable__)) all(char4 a)
-{
-  return all(a.lo) && all(a.hi);
-}
-
-int __attribute__((__overloadable__)) all(char8 a)
-{
-  return all(a.lo) && all(a.hi);
-}
-
-int __attribute__((__overloadable__)) all(char16 a)
-{
-  return all(a.lo) && all(a.hi);
-}
-
-int __attribute__((__overloadable__)) all(short a)
+int _cl_overloadable all(char4 a)
+{
+  return all(a.lo) && all(a.hi);
+}
+
+int _cl_overloadable all(char8 a)
+{
+  return all(a.lo) && all(a.hi);
+}
+
+int _cl_overloadable all(char16 a)
+{
+  return all(a.lo) && all(a.hi);
+}
+
+int _cl_overloadable all(short a)
 {
   return a < (short)0;
 }
 
-int __attribute__((__overloadable__)) all(short2 a)
+int _cl_overloadable all(short2 a)
 {
   return all(a.lo) && all(a.hi);
 }
 
-int __attribute__((__overloadable__)) all(short3 a)
+int _cl_overloadable all(short3 a)
 {
   return all(a.s01) && all(a.s2);
 }
 
-int __attribute__((__overloadable__)) all(short4 a)
-{
-  return all(a.lo) && all(a.hi);
-}
-
-int __attribute__((__overloadable__)) all(short8 a)
-{
-  return all(a.lo) && all(a.hi);
-}
-
-int __attribute__((__overloadable__)) all(short16 a)
-{
-  return all(a.lo) && all(a.hi);
-}
-
-int __attribute__((__overloadable__)) all(int a)
+int _cl_overloadable all(short4 a)
+{
+  return all(a.lo) && all(a.hi);
+}
+
+int _cl_overloadable all(short8 a)
+{
+  return all(a.lo) && all(a.hi);
+}
+
+int _cl_overloadable all(short16 a)
+{
+  return all(a.lo) && all(a.hi);
+}
+
+int _cl_overloadable all(int a)
 {
   return a < 0;
 }
 
-int __attribute__((__overloadable__)) all(int2 a)
+int _cl_overloadable all(int2 a)
 {
   return all(a.lo) && all(a.hi);
 }
 
-int __attribute__((__overloadable__)) all(int3 a)
+int _cl_overloadable all(int3 a)
 {
   return all(a.s01) && all(a.s2);
 }
 
-int __attribute__((__overloadable__)) all(int4 a)
-{
-  return all(a.lo) && all(a.hi);
-}
-
-int __attribute__((__overloadable__)) all(int8 a)
-{
-  return all(a.lo) && all(a.hi);
-}
-
-int __attribute__((__overloadable__)) all(int16 a)
+int _cl_overloadable all(int4 a)
+{
+  return all(a.lo) && all(a.hi);
+}
+
+int _cl_overloadable all(int8 a)
+{
+  return all(a.lo) && all(a.hi);
+}
+
+int _cl_overloadable all(int16 a)
 {
   return all(a.lo) && all(a.hi);
 }
 
 #ifdef cles_khr_int64
-int __attribute__((__overloadable__)) all(long a)
+int _cl_overloadable all(long a)
 {
   return a < 0L;
 }
 
-int __attribute__((__overloadable__)) all(long2 a)
+int _cl_overloadable all(long2 a)
 {
   return all(a.lo) && all(a.hi);
 }
 
-int __attribute__((__overloadable__)) all(long3 a)
+int _cl_overloadable all(long3 a)
 {
   return all(a.s01) && all(a.s2);
 }
 
-int __attribute__((__overloadable__)) all(long4 a)
-{
-  return all(a.lo) && all(a.hi);
-}
-
-int __attribute__((__overloadable__)) all(long8 a)
-{
-  return all(a.lo) && all(a.hi);
-}
-
-int __attribute__((__overloadable__)) all(long16 a)
+int _cl_overloadable all(long4 a)
+{
+  return all(a.lo) && all(a.hi);
+}
+
+int _cl_overloadable all(long8 a)
+{
+  return all(a.lo) && all(a.hi);
+}
+
+int _cl_overloadable all(long16 a)
 {
   return all(a.lo) && all(a.hi);
 }

=== modified file 'lib/kernel/any.cl'
--- lib/kernel/any.cl	2011-11-24 15:56:10 +0000
+++ lib/kernel/any.cl	2011-12-18 04:51:24 +0000
@@ -21,123 +21,123 @@
    THE SOFTWARE.
 */
 
-int __attribute__((__overloadable__)) any(char a)
+int _cl_overloadable any(char a)
 {
   return a < (char)0;
 }
 
-int __attribute__((__overloadable__)) any(char2 a)
+int _cl_overloadable any(char2 a)
 {
   return any(a.lo) || any(a.hi);
 }
 
-int __attribute__((__overloadable__)) any(char3 a)
+int _cl_overloadable any(char3 a)
 {
   return any(a.s01) || any(a.s2);
 }
 
-int __attribute__((__overloadable__)) any(char4 a)
-{
-  return any(a.lo) || any(a.hi);
-}
-
-int __attribute__((__overloadable__)) any(char8 a)
-{
-  return any(a.lo) || any(a.hi);
-}
-
-int __attribute__((__overloadable__)) any(char16 a)
-{
-  return any(a.lo) || any(a.hi);
-}
-
-int __attribute__((__overloadable__)) any(short a)
+int _cl_overloadable any(char4 a)
+{
+  return any(a.lo) || any(a.hi);
+}
+
+int _cl_overloadable any(char8 a)
+{
+  return any(a.lo) || any(a.hi);
+}
+
+int _cl_overloadable any(char16 a)
+{
+  return any(a.lo) || any(a.hi);
+}
+
+int _cl_overloadable any(short a)
 {
   return a < (short)0;
 }
 
-int __attribute__((__overloadable__)) any(short2 a)
+int _cl_overloadable any(short2 a)
 {
   return any(a.lo) || any(a.hi);
 }
 
-int __attribute__((__overloadable__)) any(short3 a)
+int _cl_overloadable any(short3 a)
 {
   return any(a.s01) || any(a.s2);
 }
 
-int __attribute__((__overloadable__)) any(short4 a)
-{
-  return any(a.lo) || any(a.hi);
-}
-
-int __attribute__((__overloadable__)) any(short8 a)
-{
-  return any(a.lo) || any(a.hi);
-}
-
-int __attribute__((__overloadable__)) any(short16 a)
-{
-  return any(a.lo) || any(a.hi);
-}
-
-int __attribute__((__overloadable__)) any(int a)
+int _cl_overloadable any(short4 a)
+{
+  return any(a.lo) || any(a.hi);
+}
+
+int _cl_overloadable any(short8 a)
+{
+  return any(a.lo) || any(a.hi);
+}
+
+int _cl_overloadable any(short16 a)
+{
+  return any(a.lo) || any(a.hi);
+}
+
+int _cl_overloadable any(int a)
 {
   return a < 0;
 }
 
-int __attribute__((__overloadable__)) any(int2 a)
+int _cl_overloadable any(int2 a)
 {
   return any(a.lo) || any(a.hi);
 }
 
-int __attribute__((__overloadable__)) any(int3 a)
+int _cl_overloadable any(int3 a)
 {
   return any(a.s01) || any(a.s2);
 }
 
-int __attribute__((__overloadable__)) any(int4 a)
-{
-  return any(a.lo) || any(a.hi);
-}
-
-int __attribute__((__overloadable__)) any(int8 a)
-{
-  return any(a.lo) || any(a.hi);
-}
-
-int __attribute__((__overloadable__)) any(int16 a)
+int _cl_overloadable any(int4 a)
+{
+  return any(a.lo) || any(a.hi);
+}
+
+int _cl_overloadable any(int8 a)
+{
+  return any(a.lo) || any(a.hi);
+}
+
+int _cl_overloadable any(int16 a)
 {
   return any(a.lo) || any(a.hi);
 }
 
 #ifdef cles_khr_int64
-int __attribute__((__overloadable__)) any(long a)
+int _cl_overloadable any(long a)
 {
   return a < 0L;
 }
 
-int __attribute__((__overloadable__)) any(long2 a)
+int _cl_overloadable any(long2 a)
 {
   return any(a.lo) || any(a.hi);
 }
 
-int __attribute__((__overloadable__)) any(long3 a)
+int _cl_overloadable any(long3 a)
 {
   return any(a.s01) || any(a.s2);
 }
 
-int __attribute__((__overloadable__)) any(long4 a)
-{
-  return any(a.lo) || any(a.hi);
-}
-
-int __attribute__((__overloadable__)) any(long8 a)
-{
-  return any(a.lo) || any(a.hi);
-}
-
-int __attribute__((__overloadable__)) any(long16 a)
+int _cl_overloadable any(long4 a)
+{
+  return any(a.lo) || any(a.hi);
+}
+
+int _cl_overloadable any(long8 a)
+{
+  return any(a.lo) || any(a.hi);
+}
+
+int _cl_overloadable any(long16 a)
 {
   return any(a.lo) || any(a.hi);
 }

=== added file 'lib/kernel/arm/vload_half.cl'
--- lib/kernel/arm/vload_half.cl	1970-01-01 00:00:00 +0000
+++ lib/kernel/arm/vload_half.cl	2011-12-18 04:51:24 +0000
@@ -0,0 +1,109 @@
+/* OpenCL built-in library: vload_half()
+
+   Copyright (c) 2011 Universidad Rey Juan Carlos
+   
+   Permission is hereby granted, free of charge, to any person obtaining a copy
+   of this software and associated documentation files (the "Software"), to deal
+   in the Software without restriction, including without limitation the rights
+   to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+   copies of the Software, and to permit persons to whom the Software is
+   furnished to do so, subject to the following conditions:
+   
+   The above copyright notice and this permission notice shall be included in
+   all copies or substantial portions of the Software.
+   
+   THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+   IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+   FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+   AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+   LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+   OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+   THE SOFTWARE.
+*/
+
+
+
+#define IMPLEMENT_VLOAD_HALF(MOD)                       \
+                                                        \
+  float _cl_overloadable                                \
+  vload_half(size_t offset, const MOD half *p)          \
+  {                                                     \
+    return p[offset];                                   \
+  }                                                     \
+                                                        \
+  float2 _cl_overloadable                               \
+  vload_half2(size_t offset, const MOD half *p)         \
+  {                                                     \
+    return (float2)(vload_half(0, &p[offset*2]),        \
+                    vload_half(0, &p[offset*2+1]));     \
+  }                                                     \
+                                                        \
+  float3 _cl_overloadable                               \
+  vload_half3(size_t offset, const MOD half *p)         \
+  {                                                     \
+    return (float3)(vload_half2(0, &p[offset*3]),       \
+                    vload_half(0, &p[offset*3+2]));     \
+  }                                                     \
+                                                        \
+  float4 _cl_overloadable                               \
+  vload_half4(size_t offset, const MOD half *p)         \
+  {                                                     \
+    return (float4)(vload_half2(0, &p[offset*4]),       \
+                    vload_half2(0, &p[offset*4+2]));    \
+  }                                                     \
+                                                        \
+  float8 _cl_overloadable                               \
+  vload_half8(size_t offset, const MOD half *p)         \
+  {                                                     \
+    return (float8)(vload_half4(0, &p[offset*8]),       \
+                    vload_half4(0, &p[offset*8+4]));    \
+  }                                                     \
+                                                        \
+  float16 _cl_overloadable                              \
+  vload_half16(size_t offset, const MOD half *p)        \
+  {                                                     \
+    return (float16)(vload_half8(0, &p[offset*16]),     \
+                     vload_half8(0, &p[offset*16+8]));  \
+  }                                                     \
+                                                        \
+  float2 _cl_overloadable                               \
+  vloada_half2(size_t offset, const MOD half *p)        \
+  {                                                     \
+    return (float2)(vload_half(0, &p[offset*2]),        \
+                    vload_half(0, &p[offset*2+1]));     \
+  }                                                     \
+                                                        \
+  float3 _cl_overloadable                               \
+  vloada_half3(size_t offset, const MOD half *p)        \
+  {                                                     \
+    return (float3)(vloada_half2(0, &p[offset*4]),      \
+                    vload_half(0, &p[offset*4+2]));     \
+  }                                                     \
+                                                        \
+  float4 _cl_overloadable                               \
+  vloada_half4(size_t offset, const MOD half *p)        \
+  {                                                     \
+    return (float4)(vloada_half2(0, &p[offset*4]),      \
+                    vloada_half2(0, &p[offset*4+2]));   \
+  }                                                     \
+                                                        \
+  float8 _cl_overloadable                               \
+  vloada_half8(size_t offset, const MOD half *p)        \
+  {                                                     \
+    return (float8)(vloada_half4(0, &p[offset*8]),      \
+                    vloada_half4(0, &p[offset*8+4]));   \
+  }                                                     \
+                                                        \
+  float16 _cl_overloadable                              \
+  vloada_half16(size_t offset, const MOD half *p)       \
+  {                                                     \
+    return (float16)(vloada_half8(0, &p[offset*16]),    \
+                     vloada_half8(0, &p[offset*16+8])); \
+  }
+
+
+
+IMPLEMENT_VLOAD_HALF(__global)
+IMPLEMENT_VLOAD_HALF(__local)
+IMPLEMENT_VLOAD_HALF(__constant)
+/* IMPLEMENT_VLOAD_HALF(__private) */

=== added file 'lib/kernel/arm/vstore_half.cl'
--- lib/kernel/arm/vstore_half.cl	1970-01-01 00:00:00 +0000
+++ lib/kernel/arm/vstore_half.cl	2011-12-18 04:51:24 +0000
@@ -0,0 +1,120 @@
+/* OpenCL built-in library: vstore_half()
+
+   Copyright (c) 2011 Universidad Rey Juan Carlos
+   
+   Permission is hereby granted, free of charge, to any person obtaining a copy
+   of this software and associated documentation files (the "Software"), to deal
+   in the Software without restriction, including without limitation the rights
+   to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+   copies of the Software, and to permit persons to whom the Software is
+   furnished to do so, subject to the following conditions:
+   
+   The above copyright notice and this permission notice shall be included in
+   all copies or substantial portions of the Software.
+   
+   THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+   IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+   FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+   AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+   LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+   OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+   THE SOFTWARE.
+*/
+
+
+
+#define IMPLEMENT_VSTORE_HALF(MOD, SUFFIX)                              \
+                                                                        \
+  void _cl_overloadable                                                 \
+  vstore_half##SUFFIX(float data, size_t offset, MOD half *p)           \
+  {                                                                     \
+    p[offset] = data;                                                   \
+  }                                                                     \
+                                                                        \
+  void _cl_overloadable                                                 \
+  vstore_half2##SUFFIX(float2 data, size_t offset, MOD half *p)         \
+  {                                                                     \
+    vstore_half##SUFFIX(data.lo, 0, &p[offset*2]);                      \
+    vstore_half##SUFFIX(data.hi, 0, &p[offset*2+1]);                    \
+  }                                                                     \
+                                                                        \
+  void _cl_overloadable                                                 \
+  vstore_half3##SUFFIX(float3 data, size_t offset, MOD half *p)         \
+  {                                                                     \
+    vstore_half2##SUFFIX(data.lo, 0, &p[offset*3]);                     \
+    vstore_half##SUFFIX(data.s2, 0, &p[offset*3+2]);                    \
+  }                                                                     \
+                                                                        \
+  void _cl_overloadable                                                 \
+  vstore_half4##SUFFIX(float4 data, size_t offset, MOD half *p)         \
+  {                                                                     \
+    vstore_half2##SUFFIX(data.lo, 0, &p[offset*4]);                     \
+    vstore_half2##SUFFIX(data.hi, 0, &p[offset*4+2]);                   \
+  }                                                                     \
+                                                                        \
+  void _cl_overloadable                                                 \
+  vstore_half8##SUFFIX(float8 data, size_t offset, MOD half *p)         \
+  {                                                                     \
+    vstore_half4##SUFFIX(data.lo, 0, &p[offset*8]);                     \
+    vstore_half4##SUFFIX(data.hi, 0, &p[offset*8+4]);                   \
+  }                                                                     \
+                                                                        \
+  void _cl_overloadable                                                 \
+  vstore_half16##SUFFIX(float16 data, size_t offset, MOD half *p)       \
+  {                                                                     \
+    vstore_half8##SUFFIX(data.lo, 0, &p[offset*16]);                    \
+    vstore_half8##SUFFIX(data.hi, 0, &p[offset*16+8]);                  \
+  }                                                                     \
+                                                                        \
+  void _cl_overloadable                                                 \
+  vstorea_half2##SUFFIX(float2 data, size_t offset, MOD half *p)        \
+  {                                                                     \
+    vstore_half##SUFFIX(data.lo, 0, &p[offset*2]);                      \
+    vstore_half##SUFFIX(data.hi, 0, &p[offset*2+1]);                    \
+  }                                                                     \
+                                                                        \
+  void _cl_overloadable                                                 \
+  vstorea_half3##SUFFIX(float3 data, size_t offset, MOD half *p)        \
+  {                                                                     \
+    vstorea_half2##SUFFIX(data.lo, 0, &p[offset*3]);                    \
+    vstore_half##SUFFIX(data.s2, 0, &p[offset*3+2]);                    \
+  }                                                                     \
+                                                                        \
+  void _cl_overloadable                                                 \
+  vstorea_half4##SUFFIX(float4 data, size_t offset, MOD half *p)        \
+  {                                                                     \
+    vstorea_half2##SUFFIX(data.lo, 0, &p[offset*4]);                    \
+    vstorea_half2##SUFFIX(data.hi, 0, &p[offset*4+2]);                  \
+  }                                                                     \
+                                                                        \
+  void _cl_overloadable                                                 \
+  vstorea_half8##SUFFIX(float8 data, size_t offset, MOD half *p)        \
+  {                                                                     \
+    vstorea_half4##SUFFIX(data.lo, 0, &p[offset*8]);                    \
+    vstorea_half4##SUFFIX(data.hi, 0, &p[offset*8+4]);                  \
+  }                                                                     \
+                                                                        \
+  void _cl_overloadable                                                 \
+  vstorea_half16##SUFFIX(float16 data, size_t offset, MOD half *p)      \
+  {                                                                     \
+    vstorea_half8##SUFFIX(data.lo, 0, &p[offset*16]);                   \
+    vstorea_half8##SUFFIX(data.hi, 0, &p[offset*16+8]);                 \
+  }
+
+
+
+IMPLEMENT_VSTORE_HALF(__global  ,     )
+IMPLEMENT_VSTORE_HALF(__global  , _rte)
+IMPLEMENT_VSTORE_HALF(__global  , _rtz)
+IMPLEMENT_VSTORE_HALF(__global  , _rtp)
+IMPLEMENT_VSTORE_HALF(__global  , _rtn)
+IMPLEMENT_VSTORE_HALF(__local   ,     )
+IMPLEMENT_VSTORE_HALF(__local   , _rte)
+IMPLEMENT_VSTORE_HALF(__local   , _rtz)
+IMPLEMENT_VSTORE_HALF(__local   , _rtp)
+IMPLEMENT_VSTORE_HALF(__local   , _rtn)
+/* IMPLEMENT_VSTORE_HALF(__private ,     ) */
+/* IMPLEMENT_VSTORE_HALF(__private , _rte) */
+/* IMPLEMENT_VSTORE_HALF(__private , _rtz) */
+/* IMPLEMENT_VSTORE_HALF(__private , _rtp) */
+/* IMPLEMENT_VSTORE_HALF(__private , _rtn) */

=== modified file 'lib/kernel/as_type.cl'
--- lib/kernel/as_type.cl	2011-11-25 17:02:42 +0000
+++ lib/kernel/as_type.cl	2011-12-18 04:51:24 +0000
@@ -22,7 +22,7 @@
 */
 
 #define DEFINE_AS_TYPE(SRC, DST)                \
-  DST __attribute__ ((__overloadable__))        \
+  DST _cl_overloadable                          \
   as_##DST(SRC a)                               \
   {                                             \
     return *(DST*)&a;                           \

=== modified file 'lib/kernel/convert_type.cl'
--- lib/kernel/convert_type.cl	2011-11-25 16:47:53 +0000
+++ lib/kernel/convert_type.cl	2011-12-18 04:51:24 +0000
@@ -24,28 +24,28 @@
 #include "templates.h"
 
 /* SRC and DST must be scalars */
-#define DEFINE_CONVERT_TYPE(SRC, DST)                           \
-  DST __attribute__ ((__overloadable__)) convert_##DST(SRC a)   \
-  {                                                             \
-    return (DST)a;                                              \
+#define DEFINE_CONVERT_TYPE(SRC, DST)           \
+  DST _cl_overloadable convert_##DST(SRC a)     \
+  {                                             \
+    return (DST)a;                              \
   }
 
 /* implementing vector SRC and DST in terms of scalars */
 #define DEFINE_CONVERT_TYPE_HALF(SRC, DST, HALFDST)                     \
-  DST __attribute__ ((__overloadable__)) convert_##DST(SRC a)           \
+  DST _cl_overloadable convert_##DST(SRC a)                             \
   {                                                                     \
     return (DST)(convert_##HALFDST(a.lo), convert_##HALFDST(a.hi));     \
   }
 
 #define DEFINE_CONVERT_TYPE_012(SRC, DST, DST01, DST2)          \
-  DST __attribute__ ((__overloadable__)) convert_##DST(SRC a)   \
+  DST _cl_overloadable convert_##DST(SRC a)                     \
   {                                                             \
     return (DST)(convert_##DST01(a.s01), convert_##DST2(a.s2)); \
   }
 
 /* SRC and DST may be vectors */
 #define DEFINE_CONVERT_TYPE_SAT(SRC, DST, SIZE)                         \
-  DST##SIZE __attribute__ ((__overloadable__))                          \
+  DST##SIZE _cl_overloadable                                            \
     convert_##DST##SIZE##_sat(SRC##SIZE a)                              \
   {                                                                     \
     int const src_float    = (SRC)0.1f > (SRC)0;                        \

=== modified file 'lib/kernel/cross.cl'
--- lib/kernel/cross.cl	2011-11-01 16:33:17 +0000
+++ lib/kernel/cross.cl	2011-12-18 04:51:24 +0000
@@ -21,12 +21,12 @@
    THE SOFTWARE.
 */
 
-float4 __attribute__ ((__overloadable__)) cross(float4 a, float4 b)
+float4 _cl_overloadable cross(float4 a, float4 b)
 {
   return (float4)(cross(a.xyz, b.xyz), 0.0f);
 }
 
-float3 __attribute__ ((__overloadable__)) cross(float3 a, float3 b)
+float3 _cl_overloadable cross(float3 a, float3 b)
 {
   return (float3)(a.y * b.z - a.z * b.y,
                   a.z * b.x - a.x * b.z,
@@ -34,12 +34,12 @@
 }
 
 #ifdef cl_khr_fp64
-double4 __attribute__ ((__overloadable__)) cross(double4 a, double4 b)
+double4 _cl_overloadable cross(double4 a, double4 b)
 {
   return (double4)(cross(a.xyz, b.xyz), 0.0f);
 }
 
-double3 __attribute__ ((__overloadable__)) cross(double3 a, double3 b)
+double3 _cl_overloadable cross(double3 a, double3 b)
 {
   return (double3)(a.y * b.z - a.z * b.y,
                    a.z * b.x - a.x * b.z,

=== modified file 'lib/kernel/dot.cl'
--- lib/kernel/dot.cl	2011-11-01 16:33:17 +0000
+++ lib/kernel/dot.cl	2011-12-18 04:51:24 +0000
@@ -21,63 +21,63 @@
    THE SOFTWARE.
 */
 
-float __attribute__ ((__overloadable__)) dot(float a, float b)
+float _cl_overloadable dot(float a, float b)
 {
   return a * b;
 }
 
-float __attribute__ ((__overloadable__)) dot(float2 a, float2 b)
+float _cl_overloadable dot(float2 a, float2 b)
 {
   return a.lo * b.lo + a.hi * b.hi;
 }
 
-float __attribute__ ((__overloadable__)) dot(float3 a, float3 b)
+float _cl_overloadable dot(float3 a, float3 b)
 {
   return dot(a.s01, b.s01) + a.s2 * b.s2;
 }
 
-float __attribute__ ((__overloadable__)) dot(float4 a, float4 b)
-{
-  return dot(a.lo, b.lo) + dot(a.hi, b.hi);
-}
-
-float __attribute__ ((__overloadable__)) dot(float8 a, float8 b)
-{
-  return dot(a.lo, b.lo) + dot(a.hi, b.hi);
-}
-
-float __attribute__ ((__overloadable__)) dot(float16 a, float16 b)
+float _cl_overloadable dot(float4 a, float4 b)
+{
+  return dot(a.lo, b.lo) + dot(a.hi, b.hi);
+}
+
+float _cl_overloadable dot(float8 a, float8 b)
+{
+  return dot(a.lo, b.lo) + dot(a.hi, b.hi);
+}
+
+float _cl_overloadable dot(float16 a, float16 b)
 {
   return dot(a.lo, b.lo) + dot(a.hi, b.hi);
 }
 
 #ifdef cl_khr_fp64
-double __attribute__ ((__overloadable__)) dot(double a, double b)
+double _cl_overloadable dot(double a, double b)
 {
   return a * b;
 }
 
-double __attribute__ ((__overloadable__)) dot(double2 a, double2 b)
+double _cl_overloadable dot(double2 a, double2 b)
 {
   return a.lo * b.lo + a.hi * b.hi;
 }
 
-double __attribute__ ((__overloadable__)) dot(double3 a, double3 b)
+double _cl_overloadable dot(double3 a, double3 b)
 {
   return dot(a.s01, b.s01) + a.s2 * b.s2;
 }
 
-double __attribute__ ((__overloadable__)) dot(double4 a, double4 b)
-{
-  return dot(a.lo, b.lo) + dot(a.hi, b.hi);
-}
-
-double __attribute__ ((__overloadable__)) dot(double8 a, double8 b)
-{
-  return dot(a.lo, b.lo) + dot(a.hi, b.hi);
-}
-
-double __attribute__ ((__overloadable__)) dot(double16 a, double16 b)
+double _cl_overloadable dot(double4 a, double4 b)
+{
+  return dot(a.lo, b.lo) + dot(a.hi, b.hi);
+}
+
+double _cl_overloadable dot(double8 a, double8 b)
+{
+  return dot(a.lo, b.lo) + dot(a.hi, b.hi);
+}
+
+double _cl_overloadable dot(double16 a, double16 b)
 {
   return dot(a.lo, b.lo) + dot(a.hi, b.hi);
 }

=== modified file 'lib/kernel/upsample.cl'
--- lib/kernel/upsample.cl	2011-11-24 15:56:10 +0000
+++ lib/kernel/upsample.cl	2011-12-18 04:51:24 +0000
@@ -25,7 +25,7 @@
 // convert_* function calls
 
 #define IMPLEMENT_UPSAMPLE_LG_GUG(GTYPE, SGTYPE, UGTYPE, LGTYPE)        \
-  LGTYPE __attribute__ ((__overloadable__))                                 \
+  LGTYPE _cl_overloadable                                               \
   upsample(GTYPE a, UGTYPE b)                                           \
   {                                                                     \
     int bits = CHAR_BIT * sizeof(SGTYPE);                               \

=== modified file 'lib/kernel/vload.cl'
--- lib/kernel/vload.cl	2011-12-14 23:01:01 +0000
+++ lib/kernel/vload.cl	2011-12-18 04:51:24 +0000
@@ -27,31 +27,31 @@
 
 #define IMPLEMENT_VLOAD(TYPE, MOD)                                      \
                                                                         \
-  TYPE##2 __attribute__ ((__overloadable__))                            \
+  TYPE##2 _cl_overloadable                                              \
   vload2(size_t offset, const MOD TYPE *p)                              \
   {                                                                     \
     return (TYPE##2)(p[offset*2], p[offset*2+1]);                       \
   }                                                                     \
                                                                         \
-  TYPE##3 __attribute__ ((__overloadable__))                            \
+  TYPE##3 _cl_overloadable                                              \
   vload3(size_t offset, const MOD TYPE *p)                              \
   {                                                                     \
     return (TYPE##3)(vload2(0, &p[offset*3]), p[offset*3+2]);           \
   }                                                                     \
                                                                         \
-  TYPE##4 __attribute__ ((__overloadable__))                            \
+  TYPE##4 _cl_overloadable                                              \
   vload4(size_t offset, const MOD TYPE *p)                              \
   {                                                                     \
     return (TYPE##4)(vload2(0, &p[offset*4]), vload2(0, &p[offset*4+2])); \
   }                                                                     \
                                                                         \
-  TYPE##8 __attribute__ ((__overloadable__))                            \
+  TYPE##8 _cl_overloadable                                              \
   vload8(size_t offset, const MOD TYPE *p)                              \
   {                                                                     \
     return (TYPE##8)(vload4(0, &p[offset*8]), vload4(0, &p[offset*8+4])); \
   }                                                                     \
                                                                         \
-  TYPE##16 __attribute__ ((__overloadable__))                           \
+  TYPE##16 _cl_overloadable                                             \
   vload16(size_t offset, const MOD TYPE *p)                             \
   {                                                                     \
     return (TYPE##16)(vload8(0, &p[offset*16]), vload8(0, &p[offset*16+8])); \

=== modified file 'lib/kernel/vload_half.cl'
--- lib/kernel/vload_half.cl	2011-12-14 23:01:01 +0000
+++ lib/kernel/vload_half.cl	2011-12-18 04:51:24 +0000
@@ -25,78 +25,100 @@
 
 #ifdef cl_khr_fp16
 
-#define IMPLEMENT_VLOAD_HALF(MOD)                       \
-                                                        \
-  float    __attribute__ ((__overloadable__))           \
-  vload_half(size_t offset, const MOD half *p)          \
-  {                                                     \
-    return (float)p[offset];                            \
-  }                                                     \
-                                                        \
-  float2 __attribute__ ((__overloadable__))             \
+/*
+  half:   1 sign bit,  5 exponent bits, 10 mantissa bits
+  float:  1 sign bit,  8 exponent bits, 23 mantissa bits
+  double: 1 sign bit, 10 exponent bits, 53 mantissa bits
+*/
+
+
+
+#define IMPLEMENT_VLOAD_HALF(MOD)                                       \
+                                                                        \
+  float _cl_overloadable                                                \
+  vload_half(size_t offset, const MOD half *p)                          \
+  {                                                                     \
+    /* This conversion always succeeds */                               \
+    short hval = ((const MOD short*)p)[offset];                         \
+    short hsign = (hval & (short)0x8000) >> (short)15;                  \
+    short hexp = (hval & (short)0x7c00) >> (short)10;                   \
+    short hmant = hval & (short)0x03ff;                                 \
+    bool isdenorm = hexp == (short)0;                                   \
+    bool isinfnan = hexp == (short)31;                                  \
+    hexp -= (short)15;                                                  \
+    int fsign = (int)hsign << 31;                                       \
+    int fexp = (__builtin_expect(isdenorm, false) ? 0 :                 \
+                __builtin_expect(isinfnan, false) ? 255 : (int)hexp + 127); \
+    fexp <<= 23;                                                        \
+    int fmant = (int)hmant << 13;                                       \
+    int fval = fsign | fexp | fmant;                                    \
+    return as_float(fval);                                              \
+  }                                                                     \
+                                                        \
+  float2 _cl_overloadable                               \
   vload_half2(size_t offset, const MOD half *p)         \
   {                                                     \
     return (float2)(vload_half(0, &p[offset*2]),        \
                     vload_half(0, &p[offset*2+1]));     \
   }                                                     \
                                                         \
-  float3 __attribute__ ((__overloadable__))             \
+  float3 _cl_overloadable                               \
   vload_half3(size_t offset, const MOD half *p)         \
   {                                                     \
     return (float3)(vload_half2(0, &p[offset*3]),       \
                     vload_half(0, &p[offset*3+2]));     \
   }                                                     \
                                                         \
-  float4 __attribute__ ((__overloadable__))             \
+  float4 _cl_overloadable                               \
   vload_half4(size_t offset, const MOD half *p)         \
   {                                                     \
     return (float4)(vload_half2(0, &p[offset*4]),       \
                     vload_half2(0, &p[offset*4+2]));    \
   }                                                     \
                                                         \
-  float8 __attribute__ ((__overloadable__))             \
+  float8 _cl_overloadable                               \
   vload_half8(size_t offset, const MOD half *p)         \
   {                                                     \
     return (float8)(vload_half4(0, &p[offset*8]),       \
                     vload_half4(0, &p[offset*8+4]));    \
   }                                                     \
                                                         \
-  float16 __attribute__ ((__overloadable__))            \
+  float16 _cl_overloadable                              \
   vload_half16(size_t offset, const MOD half *p)        \
   {                                                     \
     return (float16)(vload_half8(0, &p[offset*16]),     \
                      vload_half8(0, &p[offset*16+8]));  \
   }                                                     \
                                                         \
-  float2 __attribute__ ((__overloadable__))             \
+  float2 _cl_overloadable                               \
   vloada_half2(size_t offset, const MOD half *p)        \
   {                                                     \
     return (float2)(vload_half(0, &p[offset*2]),        \
                     vload_half(0, &p[offset*2+1]));     \
   }                                                     \
                                                         \
-  float3 __attribute__ ((__overloadable__))             \
+  float3 _cl_overloadable                               \
   vloada_half3(size_t offset, const MOD half *p)        \
   {                                                     \
     return (float3)(vloada_half2(0, &p[offset*4]),      \
                     vload_half(0, &p[offset*4+2]));     \
   }                                                     \
                                                         \
-  float4 __attribute__ ((__overloadable__))             \
+  float4 _cl_overloadable                               \
   vloada_half4(size_t offset, const MOD half *p)        \
   {                                                     \
     return (float4)(vloada_half2(0, &p[offset*4]),      \
                     vloada_half2(0, &p[offset*4+2]));   \
   }                                                     \
                                                         \
-  float8 __attribute__ ((__overloadable__))             \
+  float8 _cl_overloadable                               \
   vloada_half8(size_t offset, const MOD half *p)        \
   {                                                     \
     return (float8)(vloada_half4(0, &p[offset*8]),      \
                     vloada_half4(0, &p[offset*8+4]));   \
   }                                                     \
                                                         \
-  float16 __attribute__ ((__overloadable__))            \
+  float16 _cl_overloadable                              \
   vloada_half16(size_t offset, const MOD half *p)       \
   {                                                     \
     return (float16)(vloada_half8(0, &p[offset*16]),    \

=== modified file 'lib/kernel/vstore.cl'
--- lib/kernel/vstore.cl	2011-11-25 17:02:42 +0000
+++ lib/kernel/vstore.cl	2011-12-18 04:51:24 +0000
@@ -27,35 +27,35 @@
 
 #define IMPLEMENT_VSTORE(TYPE, MOD)                     \
                                                         \
-  void __attribute__ ((__overloadable__))               \
+  void _cl_overloadable                                 \
   vstore2(TYPE##2 data, size_t offset, MOD TYPE *p)     \
   {                                                     \
     p[offset*2] = data.lo;                              \
     p[offset*2+1] = data.hi;                            \
   }                                                     \
                                                         \
-  void __attribute__ ((__overloadable__))               \
+  void _cl_overloadable                                 \
   vstore3(TYPE##3 data, size_t offset, MOD TYPE *p)     \
   {                                                     \
     vstore2(data.lo, 0, &p[offset*3]);                  \
     p[offset*3+2] = data.s2;                            \
   }                                                     \
                                                         \
-  void __attribute__ ((__overloadable__))               \
+  void _cl_overloadable                                 \
   vstore4(TYPE##4 data, size_t offset, MOD TYPE *p)     \
   {                                                     \
     vstore2(data.lo, 0, &p[offset*4]);                  \
     vstore2(data.hi, 0, &p[offset*4+2]);                \
   }                                                     \
                                                         \
-  void __attribute__ ((__overloadable__))               \
+  void _cl_overloadable                                 \
   vstore8(TYPE##8 data, size_t offset, MOD TYPE *p)     \
   {                                                     \
     vstore4(data.lo, 0, &p[offset*8]);                  \
     vstore4(data.hi, 0, &p[offset*8+4]);                \
   }                                                     \
                                                         \
-  void __attribute__ ((__overloadable__))               \
+  void _cl_overloadable                                 \
   vstore16(TYPE##16 data, size_t offset, MOD TYPE *p)   \
   {                                                     \
     vstore8(data.lo, 0, &p[offset*16]);                 \

=== modified file 'lib/kernel/vstore_half.cl'
--- lib/kernel/vstore_half.cl	2011-12-14 23:01:01 +0000
+++ lib/kernel/vstore_half.cl	2011-12-18 04:51:24 +0000
@@ -25,78 +25,114 @@
 
 #ifdef cl_khr_fp16
 
+/*
+  half:   1 sign bit,  5 exponent bits, 10 mantissa bits
+  float:  1 sign bit,  7 exponent bits, 23 mantissa bits
+  double: 1 sign bit, 10 exponent bits, 53 mantissa bits
+*/
+
+#define HALF_MAXPLUS 0x1.ffdp15f /* "one more" than HALF_MAX */
+#define HALF_MIN     0x1.0p-14f
+#define HALF_ZERO    ((short)0x0000) /* zero */
+#define HALF_INF     ((short)0x4000) /* infinity */
+#define HALF_SIGN    ((short)0x8000) /* sign bit */
+
+
+
 #define IMPLEMENT_VSTORE_HALF(MOD, SUFFIX)                              \
                                                                         \
-  void __attribute__ ((__overloadable__))                               \
+  void _cl_overloadable                                                 \
   vstore_half##SUFFIX(float data, size_t offset, MOD half *p)           \
   {                                                                     \
-    p[offset] = data;                                                   \
+    /* IDEA: modify data (e.g. add "1/2") to round correctly */         \
+    int fval = as_int(data);                                            \
+    int fsign = (fval & 0x80000000) >> 31;                              \
+    int fexp = (fval & 0x7f800000) >> 23;                               \
+    int fmant = fval & 0x007fffff;                                      \
+    bool isdenorm = fexp == 0;                                          \
+    bool isinfnan = fexp == 255;                                        \
+    fexp -= 127;                                                        \
+    short hsign = (short)fsign << (short)15;                            \
+    short hexp = (__builtin_expect(isdenorm, false) ? (short)0 :        \
+                  __builtin_expect(isinfnan, false) ? (short)31 :       \
+                  (short)fexp + (short)15);                             \
+    /* TODO: this always truncates */                                   \
+    short hmant = (short)(fmant >> 13);                                 \
+    short hval;                                                         \
+    if (__builtin_expect(fabs(data) >= HALF_MAXPLUS, false)) {          \
+      hval = data > 0.0f ? HALF_INF : HALF_INF | HALF_SIGN;             \
+    } else if (__builtin_expect(fabs(data) < HALF_MIN, false)) {        \
+      hval = signbit(data)==0 ? HALF_ZERO : HALF_ZERO | HALF_SIGN;      \
+    } else {                                                            \
+      hval = hsign | hexp | hmant;                                      \
+    }                                                                   \
+    ((MOD short*)p)[offset] = hval;                                     \
   }                                                                     \
                                                                         \
-  void __attribute__ ((__overloadable__))                               \
+  void _cl_overloadable                                                 \
   vstore_half2##SUFFIX(float2 data, size_t offset, MOD half *p)         \
   {                                                                     \
     vstore_half##SUFFIX(data.lo, 0, &p[offset*2]);                      \
     vstore_half##SUFFIX(data.hi, 0, &p[offset*2+1]);                    \
   }                                                                     \
                                                                         \
-  void __attribute__ ((__overloadable__))                               \
+  void _cl_overloadable                                                 \
   vstore_half3##SUFFIX(float3 data, size_t offset, MOD half *p)         \
   {                                                                     \
     vstore_half2##SUFFIX(data.lo, 0, &p[offset*3]);                     \
     vstore_half##SUFFIX(data.s2, 0, &p[offset*3+2]);                    \
   }                                                                     \
                                                                         \
-  void __attribute__ ((__overloadable__))                               \
+  void _cl_overloadable                                                 \
   vstore_half4##SUFFIX(float4 data, size_t offset, MOD half *p)         \
   {                                                                     \
     vstore_half2##SUFFIX(data.lo, 0, &p[offset*4]);                     \
     vstore_half2##SUFFIX(data.hi, 0, &p[offset*4+2]);                   \
   }                                                                     \
                                                                         \
-  void __attribute__ ((__overloadable__))                               \
+  void _cl_overloadable                                                 \
   vstore_half8##SUFFIX(float8 data, size_t offset, MOD half *p)         \
   {                                                                     \
     vstore_half4##SUFFIX(data.lo, 0, &p[offset*8]);                     \
     vstore_half4##SUFFIX(data.hi, 0, &p[offset*8+4]);                   \
   }                                                                     \
                                                                         \
-  void __attribute__ ((__overloadable__))                               \
+  void _cl_overloadable                                                 \
   vstore_half16##SUFFIX(float16 data, size_t offset, MOD half *p)       \
   {                                                                     \
     vstore_half8##SUFFIX(data.lo, 0, &p[offset*16]);                    \
     vstore_half8##SUFFIX(data.hi, 0, &p[offset*16+8]);                  \
   }                                                                     \
                                                                         \
-  void __attribute__ ((__overloadable__))                               \
+  void _cl_overloadable                                                 \
   vstorea_half2##SUFFIX(float2 data, size_t offset, MOD half *p)        \
   {                                                                     \
     vstore_half##SUFFIX(data.lo, 0, &p[offset*2]);                      \
     vstore_half##SUFFIX(data.hi, 0, &p[offset*2+1]);                    \
   }                                                                     \
                                                                         \
-  void __attribute__ ((__overloadable__))                               \
+  void _cl_overloadable                                                 \
   vstorea_half3##SUFFIX(float3 data, size_t offset, MOD half *p)        \
   {                                                                     \
     vstorea_half2##SUFFIX(data.lo, 0, &p[offset*3]);                    \
     vstore_half##SUFFIX(data.s2, 0, &p[offset*3+2]);                    \
   }                                                                     \
                                                                         \
-  void __attribute__ ((__overloadable__))                               \
+  void _cl_overloadable                                                 \
   vstorea_half4##SUFFIX(float4 data, size_t offset, MOD half *p)        \
   {                                                                     \
     vstorea_half2##SUFFIX(data.lo, 0, &p[offset*4]);                    \
     vstorea_half2##SUFFIX(data.hi, 0, &p[offset*4+2]);                  \
   }                                                                     \
                                                                         \
-  void __attribute__ ((__overloadable__))                               \
+  void _cl_overloadable                                                 \
   vstorea_half8##SUFFIX(float8 data, size_t offset, MOD half *p)        \
   {                                                                     \
     vstorea_half4##SUFFIX(data.lo, 0, &p[offset*8]);                    \
     vstorea_half4##SUFFIX(data.hi, 0, &p[offset*8+4]);                  \
   }                                                                     \
                                                                         \
-  void __attribute__ ((__overloadable__))                               \
+  void _cl_overloadable                                                 \
   vstorea_half16##SUFFIX(float16 data, size_t offset, MOD half *p)      \
   {                                                                     \
     vstorea_half8##SUFFIX(data.lo, 0, &p[offset*16]);                   \
@@ -105,20 +141,20 @@
 
 
 
-IMPLEMENT_VSTORE_HALF(__global  ,     )
-IMPLEMENT_VSTORE_HALF(__global  , _rte)
-IMPLEMENT_VSTORE_HALF(__global  , _rtz)
-IMPLEMENT_VSTORE_HALF(__global  , _rtp)
-IMPLEMENT_VSTORE_HALF(__global  , _rtn)
-IMPLEMENT_VSTORE_HALF(__local   ,     )
-IMPLEMENT_VSTORE_HALF(__local   , _rte)
-IMPLEMENT_VSTORE_HALF(__local   , _rtz)
-IMPLEMENT_VSTORE_HALF(__local   , _rtp)
-IMPLEMENT_VSTORE_HALF(__local   , _rtn)
-/* IMPLEMENT_VSTORE_HALF(__private ,     ) */
-/* IMPLEMENT_VSTORE_HALF(__private , _rte) */
-/* IMPLEMENT_VSTORE_HALF(__private , _rtz) */
-/* IMPLEMENT_VSTORE_HALF(__private , _rtp) */
-/* IMPLEMENT_VSTORE_HALF(__private , _rtn) */
+IMPLEMENT_VSTORE_HALF(__global ,     )
+IMPLEMENT_VSTORE_HALF(__global , _rte)
+IMPLEMENT_VSTORE_HALF(__global , _rtz)
+IMPLEMENT_VSTORE_HALF(__global , _rtp)
+IMPLEMENT_VSTORE_HALF(__global , _rtn)
+IMPLEMENT_VSTORE_HALF(__local  ,     )
+IMPLEMENT_VSTORE_HALF(__local  , _rte)
+IMPLEMENT_VSTORE_HALF(__local  , _rtz)
+IMPLEMENT_VSTORE_HALF(__local  , _rtp)
+IMPLEMENT_VSTORE_HALF(__local  , _rtn)
+/* IMPLEMENT_VSTORE_HALF(__private,     ) */
+/* IMPLEMENT_VSTORE_HALF(__private, _rte) */
+/* IMPLEMENT_VSTORE_HALF(__private, _rtz) */
+/* IMPLEMENT_VSTORE_HALF(__private, _rtp) */
+/* IMPLEMENT_VSTORE_HALF(__private, _rtn) */
 
 #endif