Merge lp:~viitanet/pocl/ImageSupport into lp:~pocl/pocl/trunk

[Timo Viitanen]

Implemented parts of OpenCL Image API:
clCreateImage2D
clCreateSampler
clReleaseSampler
clEnqueueCopyBufferToImage
clEnqueueWriteImage
clEnqueueReadImage
clGetSupportedImageFormats

read_imagef
write_imagef
get_image_width
get_image_height

Supported data types: CL_UNORM_INT8, CL_FLOAT
Supported channel orders: CL_RGBA, CL_R
Sampler always defaults to CLK_NORMALIZED_COORDINATES_FALSE | CLK_ADDRESS_NONE | CLK_FILTER_NEAREST.

Added a 'make check' test from opencl-book-samples, chapter 8, ImageFilter2D.
Other possible tests:
Chapter 14, histogram, requires atomics
Chapter 19, oclFlow, requires OpenCV
Rodinia-leukocyte crashes on compiler bug (https://bugs.launchpad.net/pocl/+bug/966992)
Rodinia-nn doesn't actually use images, but runs, TODO
Rodinia-gaussian doesn't actually use images; crashes
Rodinia-particlefilter crashes on another compiler bug

[lp:~viitanet/pocl/ImageSupport]

306. By Pekka Jääskeläinen on 2012-05-30

Some image APIs from Timo.

=== modified file 'TODO'
--- TODO	2012-05-24 21:26:09 +0000
+++ TODO	2012-05-30 08:34:23 +0000
@@ -49,15 +49,12 @@
 * 5.2.1 Creating buffer objects
 * 5.2.4 Mapping buffer objects
 * 5.3 Image objects
-  * clCreateImage2D (deprecated in OpenCL 1.2) (*, R[gaussian, leukocyte, nn])
 * 5.3.3 Reading, Writing and Copying Image Objects
-  * clEnqueueReadImage (*, R[nn])
 * 5.4 Querying, Umapping, Migrating, ... Mem objects
 * 5.4.1 Retaining and Releasing Memory Objects
 * 5.4.2 Unmapping Mapped Memory Objects
 * 5.5 Sampler objects
 * 5.5.1 Creating Sampler Objects
-  * clCreateSampler (*)
 * 5.6.1 Creating Program Objects
 * 5.7.1 Creating Kernel Objects
 * 5.9 Event objects
 
=== added file 'examples/opencl-book-samples/ImageFilter2D.cl'
--- examples/opencl-book-samples/ImageFilter2D.cl	1970-01-01 00:00:00 +0000
+++ examples/opencl-book-samples/ImageFilter2D.cl	2012-05-30 08:34:23 +0000
@@ -0,0 +1,34 @@
+
+    // Gaussian Kernel is:
+    // 1  2  1
+    // 2  4  2
+    // 1  2  1
+float kernelWeights[9] = { 1.0f, 2.0f, 1.0f, 2.0f, 4.0f, 2.0f,1.0f, 2.0f, 1.0f };  
+
+// Gaussian filter of image
+__kernel void gaussian_filter(__read_only image2d_t srcImg,
+                              __write_only image2d_t dstImg,
+                              sampler_t sampler,
+                              int width, int height)
+{
+    int2 startImageCoord = (int2) (get_global_id(0) - 1, get_global_id(1) - 1);
+    int2 endImageCoord   = (int2) (get_global_id(0) + 1, get_global_id(1) + 1);
+    int2 outImageCoord = (int2) (get_global_id(0), get_global_id(1));
+
+    if (outImageCoord.x < width && outImageCoord.y < height)
+    {
+        int weight = 0;
+        float4 outColor = (float4)(0.0f, 0.0f, 0.0f, 0.0f);
+        for( int y = startImageCoord.y; y <= endImageCoord.y; y++)
+        {
+            for( int x = startImageCoord.x; x <= endImageCoord.x; x++)
+            {
+                outColor += (read_imagef(srcImg, sampler, (int2)(x, y)) * (kernelWeights[weight] / 16.0f));
+                weight += 1;
+            }
+        }
+
+        // Write the output value to image
+        write_imagef(dstImg, outImageCoord, outColor);
+    }
+}
 
=== modified file 'examples/opencl-book-samples/Makefile.am'
--- examples/opencl-book-samples/Makefile.am	2012-05-21 08:51:32 +0000
+++ examples/opencl-book-samples/Makefile.am	2012-05-30 08:34:23 +0000
@@ -39,6 +39,7 @@
 	sed -i 's/size_t localWorkSize = maxWorkGroupSize/size_t localWorkSize = 2/g' src/Chapter_16/Dijkstra/oclDijkstraKernel.cpp; \
 	sed -i 's/device.j..type == CL_DEVICE_TYPE_GPU/device[j].type \& CL_DEVICE_TYPE_GPU/g' src/Chapter_22/spmv.c; \
 	sed -i 's/context.CL_DEVICE_TYPE_GPU/context\(CL_DEVICE_TYPE_CPU/g' src/Chapter_12/VectorAdd/vecadd.cpp; \
+	rm src/Chapter_12/Sinewave/CMakeLists.txt; \
 	mkdir -p build; cd build; \
 	cmake \
 	-D CMAKE_CXX_FLAGS:STRING='-g -O3 -DCL_USE_DEPRECATED_OPENCL_1_1_APIS -lGL -lglut ${PTHREAD_CFLAGS}' \
 
=== modified file 'include/_kernel.h'
--- include/_kernel.h	2012-03-06 23:17:03 +0000
+++ include/_kernel.h	2012-05-30 08:34:23 +0000
@@ -1772,3 +1772,39 @@
 _CL_DECLARE_ASYNC_COPY_FUNCS(float);
 __IF_FP64(_CL_DECLARE_ASYNC_COPY_FUNCS(double));
 
+// Image support
+
+typedef int sampler_t;
+
+#define CLK_ADDRESS_NONE                0x00
+#define CLK_ADDRESS_MIRRORED_REPEAT     0x01
+#define CLK_ADDRESS_REPEAT              0x02
+#define CLK_ADDRESS_CLAMP_TO_EDGE       0x03
+#define CLK_ADDRESS_CLAMP               0x04
+
+#define CLK_NORMALIZED_COORDS_FALSE     0x00
+#define CLK_NORMALIZED_COORDS_TRUE      0x08
+
+#define CLK_FILTER_NEAREST              0x00
+#define CLK_FILTER_LINEAR               0x10
+
+typedef struct image2d_t_* image2d_t;
+
+float4 _cl_overloadable read_imagef( image2d_t image,
+        sampler_t sampler,
+        int2 coord);
+
+float4 _cl_overloadable read_imagef( image2d_t image,
+        sampler_t sampler,
+        float2 coord);
+
+void _cl_overloadable write_imagef( image2d_t image,
+        int2 coord,
+        float4 color);
+
+void _cl_overloadable write_imagei( image2d_t image,
+        int2 coord,
+        int4 color);
+
+int get_image_width (image2d_t image);
+int get_image_height (image2d_t image);
 
=== modified file 'include/pocl_device.h'
--- include/pocl_device.h	2012-05-08 17:37:35 +0000
+++ include/pocl_device.h	2012-05-30 08:34:23 +0000
@@ -51,6 +51,8 @@
     unsigned short alocal_sizes[MAX_KERNEL_ARGS];
     char arg_is_local[MAX_KERNEL_ARGS];
     char arg_is_pointer[MAX_KERNEL_ARGS];
+    char arg_is_image[MAX_KERNEL_ARGS];
+    char arg_is_sampler[MAX_KERNEL_ARGS];
     pocl_workgroup work_group_func;
 } __kernel_metadata;
 
 
=== modified file 'lib/CL/Makefile.am'
--- lib/CL/Makefile.am	2012-05-28 12:03:51 +0000
+++ lib/CL/Makefile.am	2012-05-30 08:34:23 +0000
@@ -38,9 +38,9 @@
                    clGetCommandQueueInfo.c	\
                    clCreateBuffer.c		\
                    clCreateSubBuffer.c		\
-		       	   clEnqueueReadBuffer.c	\
-				   clEnqueueReadBufferRect.c	\
-	    	   	   clEnqueueMapBuffer.c	\
+                   clEnqueueReadBuffer.c	\
+                   clEnqueueReadBufferRect.c	\
+                   clEnqueueMapBuffer.c	\
                    clEnqueueUnmapMemObject.c	\
                    clReleaseMemObject.c		\
                    clRetainMemObject.c		\
@@ -89,23 +89,24 @@
                    clEnqueueTask.c		\
                    clCreateImage2D.c		\
                    clCreateImage3D.c		\
-                   clEnqueueReadImage.c		\
-                   clEnqueueWriteImage.c		\
+                   clEnqueueReadImage.c	\
+                   clEnqueueWriteImage.c	\
+                   clCreateSampler.c	\
+                   clReleaseSampler.c	\
+                   clRetainSampler.c \
+                   clGetSamplerInfo.c \
                    clEnqueueCopyImage.c		\
-                   clEnqueueMapImage.c		\
-                   clGetSupportedImageFormats.c		\
+                   clEnqueueMapImage.c		\	\
                    clGetImageInfo.c		\
                    clCreateFromGLTexture2D.c	\
                    clCreateFromGLTexture3D.c	\
                    clUnloadCompiler.c		\
+                   clGetSupportedImageFormats.c	\
                    clGetExtensionFunctionAddress.c \
                    clIcdGetPlatformIDsKHR.c \
-                   clCreateSampler.c \
-                   clReleaseSampler.c \
-                   clRetainSampler.c \
-                   clGetSamplerInfo.c \
-                   pocl_cl.h			\
-                   pocl_util.c pocl_util.h 
+                   pocl_cl.h \
+                   pocl_util.c pocl_util.h \
+                   pocl_image_util.c pocl_image_util.h 
 
 
 
 
=== modified file 'lib/CL/clCreateBuffer.c'
--- lib/CL/clCreateBuffer.c	2012-05-24 20:46:33 +0000
+++ lib/CL/clCreateBuffer.c	2012-05-30 08:34:23 +0000
@@ -49,6 +49,7 @@
   mem->map_count = 0;
   mem->mappings = NULL;
   mem->flags = flags;
+  mem->is_image = CL_FALSE;
   POCL_INIT_ICD_OBJECT(mem);
 
   /* Store the per device buffer pointers always to a known
 
=== modified file 'lib/CL/clCreateImage2D.c'
--- lib/CL/clCreateImage2D.c	2012-01-30 12:21:12 +0000
+++ lib/CL/clCreateImage2D.c	2012-05-30 08:34:23 +0000
@@ -1,17 +1,121 @@
 #include "pocl_cl.h"
+#include "assert.h"
+#include "pocl_image_util.h"
+
 CL_API_ENTRY cl_mem CL_API_CALL
 clCreateImage2D(cl_context              context,
                 cl_mem_flags            flags,
                 const cl_image_format * image_format,
-                size_t                  image_width,
-                size_t                  image_height,
+                size_t                  width,
+                size_t                  height,
                 size_t                  image_row_pitch, 
                 void *                  host_ptr,
                 cl_int *                errcode_ret)
 CL_API_SUFFIX__VERSION_1_0
 {
-  POCL_ABORT_UNIMPLEMENTED();
-  return CL_SUCCESS;
+  cl_mem mem;
+  cl_device_id device_id;
+  void *device_ptr;
+  unsigned i, j;
+  int size;
+
+  if (context == NULL)
+    POCL_ERROR(CL_INVALID_CONTEXT);
+
+  mem = (cl_mem) malloc(sizeof(struct _cl_mem));
+  if (mem == NULL)
+    POCL_ERROR(CL_OUT_OF_HOST_MEMORY);
+
+  POCL_INIT_OBJECT(mem);
+  mem->parent = NULL;
+  mem->map_count = 0;
+  mem->mappings = NULL;
+  mem->flags = flags;
+  mem->is_image = CL_TRUE;
+  
+  cl_channel_order order = image_format->image_channel_order;
+  cl_channel_type type = image_format->image_channel_data_type;
+  
+  int dev_elem_size = sizeof(cl_float); //TODO
+  int dev_channels = 4;
+  
+  if (image_row_pitch == 0)
+    image_row_pitch = width;
+  
+  if (image_row_pitch != width)
+    POCL_ABORT_UNIMPLEMENTED();
+  
+  size = width * height * dev_elem_size * dev_channels;
+  
+  mem->device_ptrs = (void **) malloc(context->num_devices * sizeof(void *));
+  if (mem->device_ptrs == NULL)
+    {
+      free(mem);
+      POCL_ERROR(CL_OUT_OF_HOST_MEMORY);
+    }  
+  
+  int host_channels;
+  
+  if (order == CL_RGBA)
+    host_channels=4;
+  else if (order == CL_R)
+    host_channels=1;
+  else
+    POCL_ABORT_UNIMPLEMENTED();
+  
+  mem->size = size;
+  mem->context = context;
+  
+  mem->image_width = width;
+  mem->image_height = height;
+  mem->image_row_pitch = image_row_pitch;
+  mem->image_channel_data_type = type;
+  mem->image_channel_order = order;
+    
+  for (i = 0; i < context->num_devices; ++i)
+    {
+      if (i > 0)
+        clRetainMemObject (mem);
+      device_id = context->devices[i];
+      device_ptr = device_id->malloc(device_id->data, 0, size, NULL);
+      
+      if (device_ptr == NULL)
+        {
+          for (j = 0; j < i; ++j)
+            {
+              device_id = context->devices[j];
+              device_id->free(device_id->data, 0, mem->device_ptrs[j]);
+            }
+          free(mem);
+          POCL_ERROR(CL_MEM_OBJECT_ALLOCATION_FAILURE);
+        }
+      mem->device_ptrs[i] = device_ptr;
+      /* The device allocator allocated from a device-host shared memory. */
+      if (flags & CL_MEM_ALLOC_HOST_PTR ||
+          flags & CL_MEM_USE_HOST_PTR)
+        POCL_ABORT_UNIMPLEMENTED();
+      
+      if (flags & CL_MEM_COPY_HOST_PTR)  
+        {
+          size_t origin[3] = { 0, 0, 0 };
+          size_t region[3] = { width, height, 1 };
+          pocl_write_image( mem,
+                            context->devices[i],
+                            origin,
+                            region,
+                            0,
+                            1,
+                            host_ptr );
+        }
+    }
+
+  POCL_RETAIN_OBJECT(context);
+  
+  
+  if (errcode_ret != NULL)
+    *errcode_ret = CL_SUCCESS;
+  
+  return mem;
 }
     
  
 
=== modified file 'lib/CL/clCreateKernel.c'
--- lib/CL/clCreateKernel.c	2012-05-23 08:38:12 +0000
+++ lib/CL/clCreateKernel.c	2012-05-30 08:34:23 +0000
@@ -143,6 +143,8 @@
   kernel->dlhandle = dlhandle; /* TODO: why is this stored? */
   kernel->arg_is_pointer = lt_dlsym(dlhandle, "_arg_is_pointer");
   kernel->arg_is_local = lt_dlsym(dlhandle, "_arg_is_local");
+  kernel->arg_is_image = lt_dlsym(dlhandle, "_arg_is_image");
+  kernel->arg_is_sampler = lt_dlsym(dlhandle, "_arg_is_sampler");
   kernel->num_locals = *(cl_uint *) lt_dlsym(dlhandle, "_num_locals");
   kernel->arguments =
     (struct pocl_argument *) malloc ((kernel->num_args + kernel->num_locals) *
 
=== modified file 'lib/CL/clCreateSampler.c'
--- lib/CL/clCreateSampler.c	2012-05-22 21:26:29 +0000
+++ lib/CL/clCreateSampler.c	2012-05-30 08:34:23 +0000
@@ -1,13 +1,33 @@
 #include "pocl_cl.h"
-
-CL_API_ENTRY cl_sampler CL_API_CALL
-clCreateSampler(cl_context           context ,
-                cl_bool              normalized_coords , 
-                cl_addressing_mode   addressing_mode , 
-                cl_filter_mode       filter_mode ,
-                cl_int *             errcode_ret ) CL_API_SUFFIX__VERSION_1_0
+extern CL_API_ENTRY cl_sampler CL_API_CALL
+clCreateSampler(cl_context          context,
+                cl_bool             normalized_coords, 
+                cl_addressing_mode  addressing_mode, 
+                cl_filter_mode      filter_mode,
+                cl_int *            errcode_ret)
+CL_API_SUFFIX__VERSION_1_0
 {
-  POCL_ABORT_UNIMPLEMENTED();
-  return CL_SUCCESS;
+  cl_sampler sampler;
+
+  if (context == NULL)
+    POCL_ERROR(CL_INVALID_CONTEXT);
+  
+  sampler = (cl_sampler) malloc(sizeof(struct _cl_sampler));
+  if (sampler == NULL)
+    POCL_ERROR(CL_OUT_OF_HOST_MEMORY);
+  
+  if (normalized_coords == CL_TRUE)
+    POCL_ABORT_UNIMPLEMENTED();
+  
+  if (addressing_mode != CL_ADDRESS_CLAMP_TO_EDGE)
+    POCL_ABORT_UNIMPLEMENTED();
+  
+  if (filter_mode != CL_FILTER_NEAREST)
+    POCL_ABORT_UNIMPLEMENTED();
+  
+  sampler->normalized_coords = normalized_coords;
+  sampler->addressing_mode = addressing_mode;
+  sampler->filter_mode = filter_mode;
+  
+  return sampler;
 }
-
 
=== modified file 'lib/CL/clEnqueueCopyBufferToImage.c'
--- lib/CL/clEnqueueCopyBufferToImage.c	2012-05-22 21:26:29 +0000
+++ lib/CL/clEnqueueCopyBufferToImage.c	2012-05-30 08:34:23 +0000
@@ -1,17 +1,50 @@
 #include "pocl_cl.h"
 
-CL_API_ENTRY cl_int CL_API_CALL
-clEnqueueCopyBufferToImage(cl_command_queue  command_queue ,
-                           cl_mem            src_buffer ,
-                           cl_mem            dst_image , 
-                           size_t            src_offset ,
-                           const size_t *    dst_origin ,
-                           const size_t *    region , 
-                           cl_uint           num_events_in_wait_list ,
-                           const cl_event *  event_wait_list ,
+extern CL_API_ENTRY cl_int CL_API_CALL
+clEnqueueCopyBufferToImage(cl_command_queue  command_queue,
+                           cl_mem            buffer,
+                           cl_mem            image, 
+                           size_t            src_offset,
+                           const size_t *    dst_origin, /*[3]*/
+                           const size_t *    region,  /*[3]*/
+                           cl_uint           num_events_in_wait_list,
+                           const cl_event *  event_wait_list,
                            cl_event *        event ) CL_API_SUFFIX__VERSION_1_0
-{
-  POCL_ABORT_UNIMPLEMENTED();
-  return CL_SUCCESS;
-}
+  {
+    if (region == NULL)
+      return CL_INVALID_VALUE;
+    
+    if (region[2] != 1) //3D image
+      POCL_ABORT_UNIMPLEMENTED();
+    
+    int dev_elem_size = sizeof(cl_float);
+    int dev_channels = 4;
+    
+    int host_elem_size;    
+    int host_channels;
+    pocl_get_image_information (image, &host_channels, &host_elem_size);
+    
+    void* temp = malloc (image->size);
+    
+    cl_device_id device_id = command_queue->device;
 
+    device_id->read
+      (device_id->data, 
+        temp, 
+        image->device_ptrs[device_id->dev_id], 
+        image->size); 
+            
+    cl_int ret_code = pocl_write_image   
+      (image,
+        command_queue->device,
+        dst_origin,
+        region,
+        0,
+        0, 
+        temp+src_offset);
+    
+    free(temp);
+    
+    return ret_code;
+  }
+  
 \ No newline at end of file
 
=== modified file 'lib/CL/clEnqueueReadBufferRect.c'
--- lib/CL/clEnqueueReadBufferRect.c	2012-05-14 11:45:48 +0000
+++ lib/CL/clEnqueueReadBufferRect.c	2012-05-30 08:34:23 +0000
@@ -23,6 +23,7 @@
 
 #include "pocl_cl.h"
 #include <assert.h>
+#include <stdio.h>
 
 CL_API_ENTRY cl_int CL_API_CALL
 clEnqueueReadBufferRect(cl_command_queue command_queue,
@@ -78,6 +79,7 @@
                        buffer_origin, host_origin, region,
                        buffer_row_pitch, buffer_slice_pitch,
                        host_row_pitch, host_slice_pitch);
+  
 
   return CL_SUCCESS;
 }
 
=== modified file 'lib/CL/clEnqueueReadImage.c'
--- lib/CL/clEnqueueReadImage.c	2012-05-22 21:26:29 +0000
+++ lib/CL/clEnqueueReadImage.c	2012-05-30 08:34:23 +0000
@@ -1,19 +1,26 @@
 #include "pocl_cl.h"
+#include "assert.h"
+#include "pocl_image_util.h"
 
-CL_API_ENTRY cl_int CL_API_CALL
-clEnqueueReadImage(cl_command_queue      command_queue ,
-                   cl_mem                image ,
-                   cl_bool               blocking_read , 
-                   const size_t *        origin ,
-                   const size_t *        region ,
-                   size_t                row_pitch ,
-                   size_t                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_0
+extern CL_API_ENTRY cl_int CL_API_CALL
+clEnqueueReadImage(cl_command_queue     command_queue,
+                   cl_mem               image,
+                   cl_bool              blocking_read, 
+                   const size_t *       origin, /* [3] */
+                   const size_t *       region, /* [3] */
+                   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_0 
 {
-  POCL_ABORT_UNIMPLEMENTED();
-  return CL_SUCCESS;
-}
-
+  return pocl_read_image(image,
+                         command_queue->device,
+                         origin,
+                         region,
+                         host_row_pitch,
+                         host_slice_pitch, 
+                         ptr);
+}
 \ No newline at end of file
 
=== modified file 'lib/CL/clEnqueueWriteImage.c'
--- lib/CL/clEnqueueWriteImage.c	2012-05-22 21:26:29 +0000
+++ lib/CL/clEnqueueWriteImage.c	2012-05-30 08:34:23 +0000
@@ -1,19 +1,25 @@
 #include "pocl_cl.h"
+#include "pocl_image_util.h"
 
-CL_API_ENTRY cl_int CL_API_CALL
-clEnqueueWriteImage(cl_command_queue     command_queue ,
-                    cl_mem               image ,
-                    cl_bool              blocking_write , 
-                    const size_t *       origin,
-                    const size_t *       region,
-                    size_t               input_row_pitch ,
-                    size_t               input_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_0
-{
-  POCL_ABORT_UNIMPLEMENTED();
-  return CL_SUCCESS;
-}
+extern CL_API_ENTRY cl_int CL_API_CALL
+clEnqueueWriteImage(cl_command_queue    command_queue,
+                    cl_mem              image,
+                    cl_bool             blocking_write, 
+                    const size_t *      origin, /*[3]*/
+                    const size_t *      region, /*[3]*/
+                    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_0
+  {
+    return pocl_write_image(image,
+                            command_queue->device,
+                            origin,
+                            region,
+                            host_row_pitch,
+                            host_slice_pitch, 
+                            ptr);
+  }
 
 
=== modified file 'lib/CL/clGetDeviceInfo.c'
--- lib/CL/clGetDeviceInfo.c	2012-05-15 22:38:06 +0000
+++ lib/CL/clGetDeviceInfo.c	2012-05-30 08:34:23 +0000
@@ -62,8 +62,7 @@
   switch (param_name)
   {
   case CL_DEVICE_IMAGE_SUPPORT: 
-    /* Return CL_FALSE until the APIs are implemented. */
-    POCL_RETURN_DEVICE_INFO(cl_bool, CL_FALSE);
+    POCL_RETURN_DEVICE_INFO(cl_bool, CL_TRUE);
   case CL_DEVICE_TYPE:
     POCL_RETURN_DEVICE_INFO(cl_device_type, device->type);   
   case CL_DEVICE_VENDOR_ID:
 
=== modified file 'lib/CL/clGetSupportedImageFormats.c'
--- lib/CL/clGetSupportedImageFormats.c	2012-05-22 21:26:29 +0000
+++ lib/CL/clGetSupportedImageFormats.c	2012-05-30 08:34:23 +0000
@@ -1,14 +1,53 @@
 #include "pocl_cl.h"
 
-CL_API_ENTRY cl_int CL_API_CALL
-clGetSupportedImageFormats(cl_context            context ,
-                           cl_mem_flags          flags ,
-                           cl_mem_object_type    image_type ,
-                           cl_uint               num_entries ,
-                           cl_image_format *     image_formats ,
-                           cl_uint *             num_image_formats ) CL_API_SUFFIX__VERSION_1_0
+extern CL_API_ENTRY cl_int CL_API_CALL
+clGetSupportedImageFormats(cl_context           context,
+                           cl_mem_flags         flags,
+                           cl_mem_object_type   image_type,
+                           cl_uint              num_entries,
+                           cl_image_format *    image_formats,
+                           cl_uint *            num_image_formats) CL_API_SUFFIX__VERSION_1_0
 {
-  POCL_ABORT_UNIMPLEMENTED();
-  return CL_SUCCESS;
+  if (context == NULL)
+    return CL_INVALID_CONTEXT;
+  
+  if (image_type != CL_MEM_OBJECT_IMAGE2D)
+    return CL_INVALID_VALUE;
+  
+  if (num_entries==0 && image_formats!=NULL)
+    return CL_INVALID_VALUE;
+  
+  int idx=0;
+  
+  const int supported_order_count = 2;
+  cl_channel_order supported_orders[] = 
+  {
+    CL_RGBA,
+    CL_R
+  };
+  
+  const int supported_type_count = 2;
+  cl_channel_type  supported_types[] =
+  {
+    CL_UNORM_INT8,
+    CL_FLOAT
+  };
+  
+  int i, j;
+  for (i=0; i<supported_order_count; i++)
+    for (j=0; j<supported_type_count; j++)
+      {
+        if (idx >= num_entries)
+          return CL_SUCCESS;
+        
+        image_formats[idx].image_channel_order = supported_orders[i];
+        image_formats[idx].image_channel_data_type = supported_types[j];
+        
+        idx++;
+      }
+      
+   // Add special cases here if a channel order is supported with only some types or vice versa.
+   *num_image_formats = idx;
+   
+   return CL_SUCCESS;
 }
-
 
=== modified file 'lib/CL/clReleaseSampler.c'
--- lib/CL/clReleaseSampler.c	2012-05-22 21:26:29 +0000
+++ lib/CL/clReleaseSampler.c	2012-05-30 08:34:23 +0000
@@ -1,9 +1,7 @@
 #include "pocl_cl.h"
-
-CL_API_ENTRY cl_int CL_API_CALL
-clReleaseSampler(cl_sampler sampler) CL_API_SUFFIX__VERSION_1_0
+extern CL_API_ENTRY cl_int CL_API_CALL
+clReleaseSampler(cl_sampler sampler)
+CL_API_SUFFIX__VERSION_1_0
 {
-  POCL_ABORT_UNIMPLEMENTED();
   return CL_SUCCESS;
 }
-
 
=== modified file 'lib/CL/devices/basic/basic.c'
--- lib/CL/devices/basic/basic.c	2012-05-28 16:17:57 +0000
+++ lib/CL/devices/basic/basic.c	2012-05-30 08:34:23 +0000
@@ -27,6 +27,7 @@
 #include <string.h>
 #include <stdlib.h>
 #include <unistd.h>
+#include <../dev_image.h>
 #include <sys/time.h>
 
 #define max(a,b) (((a) > (b)) ? (a) : (b))
@@ -235,6 +236,30 @@
         {
           arguments[i] = &((*(cl_mem *) (p->value))->device_ptrs[device]);
         }
+      else if (kernel->arg_is_image[i])
+        {
+          dev_image2d_t di;      
+          cl_mem mem = *(cl_mem*)p->value;
+          di.data = &((*(cl_mem *) (p->value))->device_ptrs[device]);
+          di.data = ((*(cl_mem *) (p->value))->device_ptrs[device]);
+          di.width = mem->image_width;
+          di.height = mem->image_height;
+          di.rowpitch = mem->image_row_pitch;
+          di.order = mem->image_channel_order;
+          di.data_type = mem->image_channel_data_type;
+          void* devptr = pocl_basic_malloc(data, 0, sizeof(dev_image2d_t), NULL);
+          arguments[i] = malloc (sizeof (void *));
+          *(void **)(arguments[i]) = devptr; 
+          pocl_basic_write( data, &di, devptr, sizeof(dev_image2d_t) );
+        }
+      else if (kernel->arg_is_sampler[i])
+        {
+          dev_sampler_t ds;
+          
+          arguments[i] = malloc (sizeof (void *));
+          *(void **)(arguments[i]) = pocl_basic_malloc(data, 0, sizeof(dev_sampler_t), NULL);
+          pocl_basic_write( data, &ds, *(void**)arguments[i], sizeof(dev_sampler_t) );
+        }
       else
         {
           arguments[i] = p->value;
 
=== added file 'lib/CL/devices/dev_image.h'
--- lib/CL/devices/dev_image.h	1970-01-01 00:00:00 +0000
+++ lib/CL/devices/dev_image.h	2012-05-30 08:34:23 +0000
@@ -0,0 +1,17 @@
+#ifndef __X86_IMAGE_H__
+#define __X86_IMAGE_H__
+
+//Definition of the image datatype used on basic and pthread (and probably tce?)
+
+typedef cl_int dev_sampler_t;
+
+typedef struct dev_image2d_t {
+  void* data;
+  cl_int width;
+  cl_int height;
+  cl_int rowpitch;
+  cl_int order;
+  cl_int data_type;
+} dev_image2d_t;
+
+#endif
 
=== modified file 'lib/CL/devices/pthread/pthread.c'
--- lib/CL/devices/pthread/pthread.c	2012-05-29 12:16:56 +0000
+++ lib/CL/devices/pthread/pthread.c	2012-05-30 08:34:23 +0000
@@ -36,6 +36,7 @@
 #ifdef CUSTOM_BUFFER_ALLOCATOR
 
 #include "bufalloc.h"
+#include <../dev_image.h>
 
 /* Instead of mallocing a buffer size for a region, try to allocate 
    this many times the buffer size to hopefully avoid mallocs for 
@@ -651,6 +652,15 @@
 }
 
 void *
+pocl_pthread_map_mem (void *data, void *buf_ptr, 
+                      size_t offset, size_t size, void* host_ptr) 
+{
+  /* All global pointers of the pthread/CPU device are in 
+     the host address space already, and up to date. */     
+  return buf_ptr + offset;
+}
+
+void *
 workgroup_thread (void *p)
 {
   struct thread_arguments *ta = (struct thread_arguments *) p;
@@ -669,6 +679,30 @@
         }
       else if (kernel->arg_is_pointer[i])
         arguments[i] = &((*(cl_mem *) (al->value))->device_ptrs[ta->device]);
+      else if (kernel->arg_is_image[i])
+        {
+          dev_image2d_t di;      
+          cl_mem mem = *(cl_mem*)al->value;
+          di.data = &((*(cl_mem *) (al->value))->device_ptrs[ta->device]);
+          di.data = ((*(cl_mem *) (al->value))->device_ptrs[ta->device]);
+          di.width = mem->image_width;
+          di.height = mem->image_height;
+          di.rowpitch = mem->image_row_pitch;
+          di.order = mem->image_channel_order;
+          di.data_type = mem->image_channel_data_type;
+          void* devptr = pocl_pthread_malloc(ta->data, 0, sizeof(dev_image2d_t), NULL);
+          arguments[i] = malloc (sizeof (void *));
+          *(void **)(arguments[i]) = devptr; 
+          pocl_pthread_write( ta->data, &di, devptr, sizeof(dev_image2d_t) );
+        }
+      else if (kernel->arg_is_sampler[i])
+        {
+          dev_sampler_t ds;
+          
+          arguments[i] = malloc (sizeof (void *));
+          *(void **)(arguments[i]) = pocl_pthread_malloc(ta->data, 0, sizeof(dev_sampler_t), NULL);
+          pocl_pthread_write( ta->data, &ds, *(void**)arguments[i], sizeof(dev_sampler_t) );
+        }
       else
         arguments[i] = al->value;
     }
@@ -699,11 +733,15 @@
 
   for (i = 0; i < kernel->num_args; ++i)
     {
-      if (kernel->arg_is_local[i])
+      if (kernel->arg_is_local[i] )
         {
           pocl_pthread_free(ta->data, 0, *(void **)(arguments[i]));
           free(arguments[i]);
         }
+      else if( kernel->arg_is_sampler[i] || kernel->arg_is_image[i] ) 
+        {
+          free(arguments[i]);
+        }
     }
   for (i = kernel->num_args;
        i < kernel->num_args + kernel->num_locals;
 
=== modified file 'lib/CL/pocl_cl.h'
--- lib/CL/pocl_cl.h	2012-05-29 12:16:56 +0000
+++ lib/CL/pocl_cl.h	2012-05-30 08:34:23 +0000
@@ -303,6 +303,15 @@
   /* in case this is a sub buffer, this points to the parent
      buffer */
   cl_mem_t *parent;
+  /* Image flags */
+  cl_bool                 is_image;
+  cl_channel_order        image_channel_order;
+  cl_channel_type         image_channel_data_type;
+  size_t                  image_width;
+  size_t                  image_height;
+  /*size_t                  image_depth;*/
+  size_t                  image_row_pitch;
+  /*size_t                  image_slice_pitch;*/
 };
 
 struct _cl_program {
@@ -337,6 +346,8 @@
   lt_dlhandle dlhandle;
   cl_int *arg_is_pointer;
   cl_int *arg_is_local;
+  cl_int *arg_is_image;
+  cl_int *arg_is_sampler;
   cl_uint num_locals;
   struct pocl_argument *arguments;
   struct _cl_kernel *next;
@@ -359,8 +370,13 @@
 
 };
 
+typedef struct _cl_sampler cl_sampler_t;
+
 struct _cl_sampler {
-  POCL_ICD_OBJECT
+  POCL_ICD_OBJECT;
+  cl_bool             normalized_coords;
+  cl_addressing_mode  addressing_mode;
+  cl_filter_mode      filter_mode;
 };
 
 #define POCL_PROFILE_QUEUED                                             \
 
=== added file 'lib/CL/pocl_image_util.c'
--- lib/CL/pocl_image_util.c	1970-01-01 00:00:00 +0000
+++ lib/CL/pocl_image_util.c	2012-05-30 08:34:23 +0000
@@ -0,0 +1,228 @@
+#include "pocl_cl.h"
+#include "pocl_image_util.h"
+#include "assert.h"
+
+extern void
+pocl_get_image_information (cl_mem        image,
+                          int* channels_out,
+                          int* elem_size_out)
+  {
+    cl_channel_order order = image->image_channel_order;
+    cl_channel_type type = image->image_channel_data_type;
+    
+    int host_elem_size;
+    if (type == CL_FLOAT)
+      host_elem_size=4;
+    else if (type==CL_UNORM_INT8)
+      host_elem_size=1;
+    else
+      POCL_ABORT_UNIMPLEMENTED();
+    if (elem_size_out != NULL) 
+      *elem_size_out = host_elem_size;
+    
+    int host_channels;
+    if (order == CL_RGBA)
+      host_channels=4;
+    else if (order == CL_R) 
+      host_channels=1;
+    else
+      POCL_ABORT_UNIMPLEMENTED();
+    if (channels_out != NULL) 
+      *channels_out = host_channels;
+  }
+extern cl_int
+pocl_write_image    (cl_mem               image,
+                   cl_device_id         device_id,
+                   const size_t *       origin_, /*[3]*/
+                   const size_t *       region_, /*[3]*/
+                   size_t               host_row_pitch,
+                   size_t               host_slice_pitch, 
+                   const void *         ptr)
+  {
+    if (image == NULL)
+      return CL_INVALID_MEM_OBJECT;
+
+    if ((ptr == NULL) ||
+        (region_ == NULL))
+      return CL_INVALID_VALUE;
+    
+    int width = image->image_width;
+    int height = image->image_height;
+    cl_channel_order order = image->image_channel_order;
+    cl_channel_type type = image->image_channel_data_type;
+    
+    size_t dev_elem_size = sizeof(cl_float);
+    int dev_channels = 4;
+    
+    int host_elem_size;
+    int host_channels;
+    pocl_get_image_information (image, &host_channels, &host_elem_size);
+    
+    size_t origin[3] = { origin_[0]*dev_elem_size*dev_channels, origin_[1], origin_[2] };
+    size_t region[3] = { region_[0]*dev_elem_size*dev_channels, region_[1], region_[2] };
+    
+    size_t image_row_pitch = width*dev_elem_size*dev_channels;
+    size_t image_slice_pitch = 0;
+    
+    if ((region[0]*region[1]*region[2] > 0) &&
+        (region[0]-1 +
+        image_row_pitch * (region[1]-1) +
+        image_slice_pitch * (region[2]-1) >= image->size))
+      return CL_INVALID_VALUE;
+    
+    cl_float* temp = malloc( width*height*dev_channels*dev_elem_size );
+    
+    if (temp == NULL) 
+      return CL_OUT_OF_HOST_MEMORY;
+    
+    int x, y, k;
+    
+    for (y=0; y<height; y++)
+      for (x=0; x<width*dev_channels; x++)
+        temp[x+y*width*dev_channels] = 0.f;
+    
+    for (y=0; y<height; y++)
+      {
+        for (x=0; x<width; x++)
+          {
+            cl_float elem[4]; //TODO 0,0,0,0 for some modes?
+            
+            for (k=0; k<host_channels; k++) 
+              {
+                if (type == CL_FLOAT)
+                  elem[k] = ((float*)ptr)[k+(x+y*width)*host_channels];
+                else if (type==CL_UNORM_INT8) 
+                  {
+                    cl_uchar foo = ((cl_uchar*)ptr)[k+(x+y*width)*host_channels];
+                    elem[k] = (float)(foo) * (1.f/255.f);
+                  }
+                else
+                  POCL_ABORT_UNIMPLEMENTED();
+              }
+          
+            if (order == CL_RGBA) 
+              for (k=0; k<4; k++)
+                temp[(x+y*width)*dev_channels+k] = elem[k];
+            else if (order == CL_R) 
+              {
+                temp[(x+y*width)*dev_channels+0] = elem[0];
+                temp[(x+y*width)*dev_channels+1] = 0.f;
+                temp[(x+y*width)*dev_channels+2] = 0.f;
+                temp[(x+y*width)*dev_channels+3] = 1.f;
+              }
+          }
+      }
+      
+    
+    device_id->write_rect(device_id->data, temp, 
+                        image->device_ptrs[device_id->dev_id],
+                        origin, origin, region,
+                        image_row_pitch, image_slice_pitch,
+                        image_row_pitch, image_slice_pitch);
+    
+    free (temp);
+    return CL_SUCCESS;
+  }
+           
+extern cl_int         
+pocl_read_image   (cl_mem               image,
+                   cl_device_id         device_id,
+                   const size_t *       origin_, /*[3]*/
+                   const size_t *       region_, /*[3]*/
+                   size_t               host_row_pitch,
+                   size_t               host_slice_pitch, 
+                   void *               ptr) 
+  {
+    
+    if (image == NULL)
+      return CL_INVALID_MEM_OBJECT;
+
+    if ((ptr == NULL) ||
+        (region_ == NULL))
+      return CL_INVALID_VALUE;
+    
+    int width = image->image_width;
+    int height = image->image_height;
+    int dev_elem_size = sizeof(cl_float);
+    int dev_channels = 4;
+    size_t origin[3] = { origin_[0]*dev_elem_size*dev_channels, origin_[1], origin_[2] };
+    size_t region[3] = { region_[0]*dev_elem_size*dev_channels, region_[1], region_[2] };
+    
+    size_t image_row_pitch = width*dev_elem_size*dev_channels;
+    size_t image_slice_pitch = 0;
+    
+    if ((region[0]*region[1]*region[2] > 0) &&
+        (region[0]-1 +
+        image_row_pitch * (region[1]-1) +
+        image_slice_pitch * (region[2]-1) >= image->size))
+      return CL_INVALID_VALUE;
+
+  
+    int i, j, k;
+  
+    cl_channel_order order = image->image_channel_order;
+    cl_channel_type type = image->image_channel_data_type;
+    
+    cl_float* temp = malloc( width*height*dev_channels*dev_elem_size );
+    
+    if (temp == NULL)
+      return CL_OUT_OF_HOST_MEMORY;
+      
+    int host_channels, host_elem_size;
+      
+    pocl_get_image_information(image, &host_channels, &host_elem_size);
+      
+    if (host_row_pitch == 0) {
+      host_row_pitch = width*host_channels;
+    }
+    
+    size_t buffer_origin[3] = { 0, 0, 0 };
+    
+    device_id->read_rect(device_id->data, temp, 
+                        image->device_ptrs[device_id->dev_id],
+                        origin, origin, region,
+                        image_row_pitch, image_slice_pitch,
+                        image_row_pitch, image_slice_pitch);
+    
+    for (j=0; j<height; j++) {
+      for (i=0; i<width; i++) {
+        cl_float elem[4];
+        
+        for (k=0; k<4; k++)
+          elem[k]=0;
+        
+        if (order == CL_RGBA) {
+          for (k=0; k<4; k++) 
+            elem[k] = temp[i*dev_channels + j*width*dev_channels + k];
+        }
+        else if (order == CL_R) { // host_channels == 1
+          elem[0] = temp[i*dev_channels + j*width*dev_channels + 0];
+        }
+        
+        if (type == CL_UNORM_INT8) 
+          { // host_channels == 4
+            for (k=0; k<host_channels; k++)
+              {
+                ((cl_uchar*)ptr)[i*host_channels + j*host_row_pitch + k] 
+                  = (unsigned char)(255*elem[k]);
+              }
+          }
+        else if (type == CL_FLOAT) 
+          {
+            for (k=0; k<host_channels; k++)
+              {
+                POCL_ABORT_UNIMPLEMENTED();
+                ((cl_float*)ptr)[i*host_channels + j*host_row_pitch + k] 
+                  = elem[k];
+              }
+          }
+        else
+          POCL_ABORT_UNIMPLEMENTED();
+      }
+    }
+    
+    free (temp);
+    
+    return CL_SUCCESS;
+  }
+  
 \ No newline at end of file
 
=== added file 'lib/CL/pocl_image_util.h'
--- lib/CL/pocl_image_util.h	1970-01-01 00:00:00 +0000
+++ lib/CL/pocl_image_util.h	2012-05-30 08:34:23 +0000
@@ -0,0 +1,50 @@
+/* OpenCL runtime library: pocl_image_util image utility functions
+
+   Copyright (c) 2012 Timo Viitanen / Tampere University of Technology
+   
+   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.
+*/
+
+#ifndef POCL_IMAGE_UTIL_H
+#define POCL_IMAGE_UTIL_H
+
+extern cl_int
+pocl_write_image    (cl_mem               image,
+                   cl_device_id         device_id,
+                   const size_t *       origin_, /*[3]*/
+                   const size_t *       region_, /*[3]*/
+                   size_t               host_row_pitch,
+                   size_t               host_slice_pitch, 
+                   const void *         ptr);
+
+extern cl_int         
+pocl_read_image     (cl_mem               image,
+                   cl_device_id         device,
+                   const size_t *       origin, /*[3]*/
+                   const size_t *       region, /*[3]*/
+                   size_t               host_row_pitch,
+                   size_t               host_slice_pitch, 
+                   void *               ptr);
+
+extern void
+pocl_get_image_information (cl_mem        image,
+                          int* host_channels,
+                          int* host_elem_size);
+                   
+#endif
 
=== added file 'lib/kernel/get_image_height.cl'
--- lib/kernel/get_image_height.cl	1970-01-01 00:00:00 +0000
+++ lib/kernel/get_image_height.cl	2012-05-30 08:34:23 +0000
@@ -0,0 +1,7 @@
+#include "templates.h"
+#include "image.h"
+
+int get_image_width (image2d_t image)
+{
+  return image->height;
+}
 \ No newline at end of file
 
=== added file 'lib/kernel/get_image_width.cl'
--- lib/kernel/get_image_width.cl	1970-01-01 00:00:00 +0000
+++ lib/kernel/get_image_width.cl	2012-05-30 08:34:23 +0000
@@ -0,0 +1,7 @@
+#include "templates.h"
+#include "image.h"
+
+int get_image_width (image2d_t image)
+{
+  return image->width;
+}
 \ No newline at end of file
 
=== added file 'lib/kernel/image.h'
--- lib/kernel/image.h	1970-01-01 00:00:00 +0000
+++ lib/kernel/image.h	2012-05-30 08:34:23 +0000
@@ -0,0 +1,15 @@
+#ifndef __IMAGE_H__
+#define __IMAGE_H__
+
+#include "templates.h"
+
+typedef struct image2d_t_ {
+  uchar4* data;
+  int width;
+  int height;
+  int rowpitch;
+  int order;
+  int data_type;
+} image2d_t_;
+
+#endif
 
=== added file 'lib/kernel/read_image.cl'
--- lib/kernel/read_image.cl	1970-01-01 00:00:00 +0000
+++ lib/kernel/read_image.cl	2012-05-30 08:34:23 +0000
@@ -0,0 +1,30 @@
+#include "templates.h"
+
+#include "image.h"
+
+float4 _cl_overloadable read_imagef ( image2d_t image,
+        sampler_t sampler,
+        int2 coord) {
+  //TODO: Sampler options
+  if( coord.x<0 )
+    coord.x=0;
+  if( coord.y<0 )
+    coord.y=0;
+  if( coord.x>=image->width )
+    coord.x=image->width-1;
+  if( coord.y>=image->height )
+    coord.y=image->height-1;
+
+  float4 color = ((float4*)image->data)[ coord.x + coord.y*image->rowpitch ];
+
+  return color;
+}
+
+float4 _cl_overloadable read_imagef ( image2d_t image,
+        sampler_t sampler,
+        float2 coord) {
+  
+  float4 color = ((float4*)image->data)[ (int)coord.x + (int)coord.y*image->rowpitch ];
+
+  return color;
+}
 
=== modified file 'lib/kernel/sources.mk'
--- lib/kernel/sources.mk	2012-03-05 15:33:38 +0000
+++ lib/kernel/sources.mk	2012-05-30 08:34:23 +0000
@@ -130,7 +130,11 @@
                       vload_half.cl		\
                       vstore_half.cl 	\
                       async_work_group_copy.cl \
-                      wait_group_events.cl
+                      wait_group_events.cl      \
+                      read_image.cl      \
+                      write_image.cl     \
+                      get_image_width.cl      \
+                      get_image_height.cl     
 
 
 
 
=== added file 'lib/kernel/write_image.cl'
--- lib/kernel/write_image.cl	1970-01-01 00:00:00 +0000
+++ lib/kernel/write_image.cl	2012-05-30 08:34:23 +0000
@@ -0,0 +1,17 @@
+#include "templates.h"
+
+#include "image.h"
+
+//typedef struct image2d_t_* image2d_t;
+
+void _cl_overloadable write_imagef (     image2d_t image,
+        int2 coord,
+        float4 color) {
+  ((float4*)image->data)[ coord.x + coord.y*image->rowpitch ] = color;
+}
+
+void _cl_overloadable write_imagei (     image2d_t image,
+        int2 coord,
+        int4 color) {
+  ((float4*)image->data)[ coord.x + coord.y*image->rowpitch ] = (float4)(color.x,color.y,color.z,color.w);
+}
 
=== modified file 'lib/llvmopencl/GenerateHeader.cc'
--- lib/llvmopencl/GenerateHeader.cc	2012-02-24 17:50:01 +0000
+++ lib/llvmopencl/GenerateHeader.cc	2012-05-30 08:34:23 +0000
@@ -123,6 +123,8 @@
   return changed;
 }
 
+#include <iostream>
+
 
 
 void
 GenerateHeader::ProcessPointers(Function *F,
@@ -134,12 +136,17 @@
       
   bool is_pointer[num_args];
   bool is_local[num_args];
+  bool is_image[num_args];
+  bool is_sampler[num_args];
     
   int i = 0;
   for (Function::const_arg_iterator ii = F->arg_begin(),
          ee = F->arg_end();
        ii != ee; ++ii) {
     Type *t = ii->getType();
+  
+    is_image[i] = false;
+    is_sampler[i] = false;
     
     if (const PointerType *p = dyn_cast<PointerType> (t)) {
       is_pointer[i] = true;
@@ -153,6 +160,22 @@
       is_pointer[i] = false;
       is_local[i] = false;
     }
+  
+    if( t->isPointerTy() ) {
+      if( t->getPointerElementType()->isStructTy() ) {
+        string name = t->getPointerElementType()->getStructName().str();
+        if( name == "struct.image2d_t_" ) { // TODO image3d?
+          is_image[i] = true;
+          is_pointer[i] = false;
+          is_local[i] = false;
+        }
+        if( name == "struct.sampler_t_" ) {
+          is_sampler[i] = true;
+          is_pointer[i] = false;
+          is_local[i] = false;
+        }
+      }
+    }
     
     ++i;
   }
@@ -172,6 +195,22 @@
       out << ", " << is_local[i];
   }
   out << "}\n";
+  
+  out << "#define _" << F->getName() << "_ARG_IS_IMAGE {";
+  if (num_args != 0) {
+    out << is_image[0];
+    for (i = 1; i < num_args; ++i)
+      out << ", " << is_image[i];
+  }
+  out << "}\n";
+  
+  out << "#define _" << F->getName() << "_ARG_IS_SAMPLER {";
+  if (num_args != 0) {
+    out << is_sampler[0];
+    for (i = 1; i < num_args; ++i)
+      out << ", " << is_sampler[i];
+  }
+  out << "}\n";
 }
 
 
 
 
=== modified file 'scripts/pocl-kernel.in'
--- scripts/pocl-kernel.in	2012-05-14 17:44:09 +0000
+++ scripts/pocl-kernel.in	2012-05-30 08:34:23 +0000
@@ -60,6 +60,8 @@
 unsigned _num_args = _${kernel}_NUM_ARGS;
 int _arg_is_pointer[] = _${kernel}_ARG_IS_POINTER;
 int _arg_is_local[] = _${kernel}_ARG_IS_LOCAL;
+int _arg_is_image[] = _${kernel}_ARG_IS_IMAGE;
+int _arg_is_sampler[] = _${kernel}_ARG_IS_SAMPLER;
 unsigned _num_locals = _${kernel}_NUM_LOCALS;
 #if _${kernel}_NUM_LOCALS != 0
 unsigned _local_sizes[_${kernel}_NUM_LOCALS] = _${kernel}_LOCAL_SIZE;
@@ -93,6 +95,8 @@
 #endif
     _${kernel}_ARG_IS_LOCAL,
     _${kernel}_ARG_IS_POINTER,
+    _${kernel}_ARG_IS_IMAGE,
+    _${kernel}_ARG_IS_SAMPLER,
     _${kernel}_workgroup_fast
 };
 EOF
 
=== modified file 'tests/testsuite-samples.at'
--- tests/testsuite-samples.at	2012-05-24 21:26:09 +0000
+++ tests/testsuite-samples.at	2012-05-30 08:34:23 +0000
@@ -73,6 +73,32 @@
           
 AT_CLEANUP
 
+AT_SETUP([Run Chapter 8: ImageFilter2D])
+AT_KEYWORDS([booksamples imagefilter2d])
+AT_SKIP_IF([! test -e $abs_top_srcdir/examples/opencl-book-samples/checkout])
+AT_CHECK_UNQUOTED([ 
+cd ${abs_top_srcdir}/examples/opencl-book-samples/checkout/src/Chapter_8/ImageFilter2D ;
+#sed '13c\ ' -i ImageFilter2D.cl ;
+#sed '14c\ ' -i ImageFilter2D.cl ;
+#sed '15c\ ' -i ImageFilter2D.cl ;
+#sed "3c float kernelWeights[9] = { 1.0f, 2.0f, 1.0f, 2.0f, 4.0f, 2.0f,1.0f, 2.0f, 1.0f }; \ " -i ImageFilter2D.cl ;
+cp ${abs_top_srcdir}/examples/opencl-book-samples/ImageFilter2D.cl ./
+sed '418cclFinish(commandQueue);' -i ImageFilter2D.cpp ;
+cd ${abs_top_srcdir}/examples/opencl-book-samples/checkout/build/src/Chapter_8/ImageFilter2D ;
+make]
+, 0, [ignore], [ignore])
+
+AT_CHECK_UNQUOTED([
+cd $abs_top_srcdir/examples/opencl-book-samples/checkout/build/src/Chapter_8/ImageFilter2D
+./ImageFilter2D ../../../../src/Chapter_19/oclFlow/data/minicooper/frame10.png output.png
+], 0, 
+[Could not create GPU context, trying CPU...
+
+Executed program succesfully.
+], [ignore])
+          
+AT_CLEANUP
+
 AT_SETUP([Run Chapter 12: VectorAdd (C++ bindings)])
 AT_KEYWORDS([booksamples])
 AT_SKIP_IF([! test -e $abs_top_srcdir/examples/opencl-book-samples/checkout])