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

[Erik Schnetter]

I corrected several build errors (and some cosmetic issues) for Mac OS X.

[Pekka Jääskeläinen (pekka-jaaskelainen)]

Fixes should be committed to the pocl/0.6 branch first and then 0.6 merged to trunk so we get fixes to the 0.6 release. Otherwise, looks good to me. Please commit to the lp:pocl/0.6 branch directly (which we can merge to trunk later). I'm probably going to roll a yet another 0.6 release candidate package later this week.

[Erik Schnetter (schnetter)]

I'm still relatively unfamiliar with bzr. How exactly would I "commit directly to 0.6"? I assume I check out the 0.6 branch first, and then what? Would "bzr merge -i" do the trick that allows me to select the fixes?

Should I commit my cosmetic changes as well, or should I commit those to the trunk only?

[Pekka Jääskeläinen (pekka-jaaskelainen)]

Yes, check out 0.6 (e.g. 'bzr co lp:pocl/0.6 pocl-0.6') then do a cherry pick merge from your branch (bzr merge -c revno ../yourbranchlocation or bzr merge -rFIRSTREV..LASTREV ../yourbranchlocation or similar) then commit (if a bound branch, just bzr commit in the 0.6 tree). Though, cherry picking book keeping is not the best in Bazaar, and IIRC the '-c' that selects a single rev is the most robust in "backporting" revisions. I cannot recall the exact problems involved but I've bumped into them in the past.

IMO code base cleanups (cosmetic) are OK to commit to 0.6 as long as we do not introduce instability/regressions.

[Erik Schnetter (schnetter)]

I cherry-picked the changes into a local checkout of the 0.6 branch.
However, it appears as if I only received the changes, not the
commits? I now have a large number of "locally modified" files... Is
this to be expected? How should I handle this?

-erik

On Mon, Aug 13, 2012 at 11:33 PM, Pekka Jääskeläinen
<email address hidden> wrote:
> Yes, check out 0.6 (e.g. 'bzr co lp:pocl/0.6 pocl-0.6') then do a cherry pick merge from your branch (bzr merge -c revno ../yourbranchlocation or bzr merge -rFIRSTREV..LASTREV ../yourbranchlocation or similar) then commit (if a bound branch, just bzr commit in the 0.6 tree). Though, cherry picking book keeping is not the best in Bazaar, and IIRC the '-c' that selects a single rev is the most robust in "backporting" revisions. I cannot recall the exact problems involved but I've bumped into them in the past.
>
> IMO code base cleanups (cosmetic) are OK to commit to 0.6 as long as we do not introduce instability/regressions.
> --
> https://code.launchpad.net/~schnetter/pocl/pocl/+merge/119393
> You are the owner of lp:~schnetter/pocl/pocl.

--
Erik Schnetter <email address hidden>
http://www.perimeterinstitute.ca/personal/eschnetter/

[Pekka Jääskeläinen (pekka-jaaskelainen)]

That's normal. The "cherry picking" of Bazaar is not perfect in this regard. http://doc.bazaar.canonical.com/latest/en/user-guide/adv_merging.html

[lp:~schnetter/pocl/pocl]

334. By Erik Schnetter on 2012-08-15

Correct several build errors for Mac OS X

[Erik Schnetter (schnetter)]

I committed my changes (except adding the atomics) to the 0.6 branch, then merged the branch into the trunk.

[Pekka Jääskeläinen (pekka-jaaskelainen)]

I can't see them in 0.6. Yep, better put the atomics to trunk, let's try to keep 0.6 in feature freeze.

[lp:~schnetter/pocl/pocl]

335. By Erik Schnetter on 2012-08-16

Implement atomics

=== modified file 'INSTALL'
--- INSTALL	2012-08-08 12:45:17 +0000
+++ INSTALL	2012-08-15 20:36:22 +0000
@@ -36,10 +36,14 @@
   ./configure --enable-shared --prefix=YOUR_INSTALLATION_PREFIX_HERE 
   make REQUIRES_RTTI=1 && make install
 
-
 !!!NOTE: A 64-bit host is at the moment required. See
 https://bugs.launchpad.net/pocl/+bug/911911
 
+!!!NOTE: On Mac OS X, you cannot build LLVM with LLVM; you need to
+build using gcc. Use e.g. this command to configure instead:
+
+  ./configure --enable-shared CC=gcc CXX=g++ --prefix=YOUR_INSTALLATION_PREFIX_HERE
+
 After all the requirements are installed. The installation procedure
 follows the usual autotools build+install. If you are using a development
 source tree, you need to generate the autotool build files with 

=== modified file 'include/_kernel.h'
--- include/_kernel.h	2012-05-25 16:05:51 +0000
+++ include/_kernel.h	2012-08-15 20:36:22 +0000
@@ -1608,6 +1608,41 @@
 #endif
 
 

+/* Atomic operations */
+
+#define _CL_DECLARE_ATOMICS(MOD, TYPE)                                  \
+  _cl_overloadable TYPE atomic_add    (volatile MOD TYPE *p, TYPE val); \
+  _cl_overloadable TYPE atomic_sub    (volatile MOD TYPE *p, TYPE val); \
+  _cl_overloadable TYPE atomic_xchg   (volatile MOD TYPE *p, TYPE val); \
+  _cl_overloadable TYPE atomic_inc    (volatile MOD TYPE *p);           \
+  _cl_overloadable TYPE atomic_dec    (volatile MOD TYPE *p);           \
+  _cl_overloadable TYPE atomic_cmpxchg(volatile MOD TYPE *p, TYPE cmp, TYPE val); \
+  _cl_overloadable TYPE atomic_min    (volatile MOD TYPE *p, TYPE val); \
+  _cl_overloadable TYPE atomic_max    (volatile MOD TYPE *p, TYPE val); \
+  _cl_overloadable TYPE atomic_and    (volatile MOD TYPE *p, TYPE val); \
+  _cl_overloadable TYPE atomic_or     (volatile MOD TYPE *p, TYPE val); \
+  _cl_overloadable TYPE atomic_xor    (volatile MOD TYPE *p, TYPE val);
+_CL_DECLARE_ATOMICS(__global, int )
+_CL_DECLARE_ATOMICS(__global, uint)
+_CL_DECLARE_ATOMICS(__local , int )
+_CL_DECLARE_ATOMICS(__local , uint)
+
+_cl_overloadable float atomic_xchg(volatile __global float *p, float val);
+_cl_overloadable float atomic_xchg(volatile __local  float *p, float val);
+
+#define atom_add     atomic_add
+#define atom_sub     atomic_sub
+#define atom_xchg    atomic_xchg
+#define atom_inc     atomic_inc
+#define atom_dec     atomic_dec
+#define atom_cmpxchg atomic_cmpxchg
+#define atom_min     atomic_min
+#define atom_max     atomic_max
+#define atom_and     atomic_and
+#define atom_or      atomic_or
+#define atom_xor     atomic_xor
+
+

 /* Miscellaneous Vector Functions */
 
 // This code leads to an ICE in Clang

=== modified file 'lib/CL/clEnqueueCopyBuffer.c'
--- lib/CL/clEnqueueCopyBuffer.c	2012-05-14 11:45:48 +0000
+++ lib/CL/clEnqueueCopyBuffer.c	2012-08-15 20:36:22 +0000
@@ -23,6 +23,7 @@
 */
 
 #include "pocl_cl.h"
+#include "pocl_icd.h"
 #include "utlist.h"
 #include <assert.h>
 
@@ -55,14 +56,28 @@
     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);
 
+  if (event != NULL)
+    {
+      *event = (cl_event)malloc(sizeof(struct _cl_event));
+      if (*event == NULL)
+        return CL_OUT_OF_HOST_MEMORY; 
+      POCL_INIT_OBJECT(*event);
+      (*event)->queue = command_queue;
+      POCL_INIT_ICD_OBJECT(*event);
+      clRetainCommandQueue (command_queue);
+
+      POCL_PROFILE_QUEUED;
+    }
+
+
   _cl_command_node * cmd = malloc(sizeof(_cl_command_node));
   if (cmd == NULL)
     return CL_OUT_OF_HOST_MEMORY;
@@ -81,8 +96,9 @@
   cmd->command.copy.dst_ptr = dst_buffer->device_ptrs[device_id->dev_id] + dst_offset;
   cmd->command.copy.cb = cb;
   cmd->next = NULL;
+  cmd->event = event ? *event : NULL;
 
-  LL_APPEND(command_queue->root, cmd );
+  LL_APPEND(command_queue->root, cmd);
 
   return CL_SUCCESS;
 }

=== modified file 'lib/CL/clEnqueueCopyBufferRect.c'
--- lib/CL/clEnqueueCopyBufferRect.c	2012-05-14 11:45:48 +0000
+++ lib/CL/clEnqueueCopyBufferRect.c	2012-08-15 20:36:22 +0000
@@ -22,6 +22,7 @@
 */
 
 #include "pocl_cl.h"
+#include "pocl_icd.h"
 #include <assert.h>
 
 CL_API_ENTRY cl_int CL_API_CALL
@@ -69,14 +70,51 @@
     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);
 
+  if (event != NULL)
+    {
+      *event = (cl_event)malloc(sizeof(struct _cl_event));
+      if (*event == NULL)
+        return CL_OUT_OF_HOST_MEMORY; 
+      POCL_INIT_OBJECT(*event);
+      (*event)->queue = command_queue;
+      POCL_INIT_ICD_OBJECT(*event);
+
+      clRetainCommandQueue (command_queue);
+
+      POCL_PROFILE_QUEUED;
+    }
+
+
+  /* execute directly */
+  /* TODO: enqueue the read_rect if this is a non-blocking read (see
+     clEnqueueReadBuffer) */
+  if (command_queue->properties & CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE)
+    {
+      /* wait for the event in event_wait_list to finish */
+      POCL_ABORT_UNIMPLEMENTED();
+    }
+  else
+    {
+      /* in-order queue - all previously enqueued commands must 
+       * finish before this read */
+      // ensure our buffer is not freed yet
+      clRetainMemObject (src_buffer);
+      clRetainMemObject (dst_buffer);
+      clFinish(command_queue);
+    }
+  POCL_PROFILE_SUBMITTED;
+  POCL_PROFILE_RUNNING;
+
+  /* TODO: offset computation doesn't work in case the ptr is not 
+     a direct pointer */
   device_id->copy_rect(device_id->data,
                        src_buffer->device_ptrs[device_id->dev_id], 
                        dst_buffer->device_ptrs[device_id->dev_id],
@@ -84,5 +122,10 @@
                        src_row_pitch, src_slice_pitch,
                        dst_row_pitch, dst_slice_pitch);
 
+  POCL_PROFILE_COMPLETE;
+
+  clReleaseMemObject (src_buffer);
+  clReleaseMemObject (dst_buffer);
+
   return CL_SUCCESS;
 }

=== modified file 'lib/CL/clEnqueueNDRangeKernel.c'
--- lib/CL/clEnqueueNDRangeKernel.c	2012-05-30 14:10:44 +0000
+++ lib/CL/clEnqueueNDRangeKernel.c	2012-08-15 20:36:22 +0000
@@ -314,6 +314,8 @@
       }
   }
 
+  command_node->event = event ? *event : NULL;
+
   LL_APPEND(command_queue->root, command_node);
 
   return CL_SUCCESS;

=== modified file 'lib/CL/clEnqueueReadBuffer.c'
--- lib/CL/clEnqueueReadBuffer.c	2012-05-29 12:16:56 +0000
+++ lib/CL/clEnqueueReadBuffer.c	2012-08-15 20:36:22 +0000
@@ -38,6 +38,7 @@
                     cl_event *event) CL_API_SUFFIX__VERSION_1_0
 {
   cl_device_id device;
+  unsigned i;
 
   if (command_queue == NULL)
     return CL_INVALID_COMMAND_QUEUE;
@@ -54,6 +55,13 @@
 
   device = command_queue->device;
 
+  for (i = 0; i < command_queue->context->num_devices; ++i)
+    {
+        if (command_queue->context->devices[i] == device)
+            break;
+    }
+  assert(i < command_queue->context->num_devices);
+
   if (event != NULL)
     {
       *event = (cl_event)malloc(sizeof(struct _cl_event));
@@ -62,14 +70,15 @@
       POCL_INIT_OBJECT(*event);
       (*event)->queue = command_queue;
       POCL_INIT_ICD_OBJECT(*event);
-      
       clRetainCommandQueue (command_queue);
 
-      POCL_PROFILE_QUEUED;      
+      POCL_PROFILE_QUEUED;
     }
 
 
   /* enqueue the read, or execute directly */
+  /* TODO: why do we implement both? direct execution seems
+     unnecessary. */
   if (blocking_read)
     {
       if (command_queue->properties & CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE)
@@ -101,16 +110,17 @@
     _cl_command_node * cmd = malloc(sizeof(_cl_command_node));
     if (cmd == NULL)
       return CL_OUT_OF_HOST_MEMORY;
-    
+
     cmd->type = CL_COMMAND_TYPE_READ;
     cmd->command.read.data = device->data;
     cmd->command.read.host_ptr = ptr;
     cmd->command.read.device_ptr = buffer->device_ptrs[device->dev_id]+offset;
     cmd->command.read.cb = cb;
+    cmd->command.read.buffer = buffer;
     cmd->next = NULL;
-    cmd->command.read.buffer = buffer;
+    cmd->event = event ? *event : NULL;
     clRetainMemObject (buffer);
-    LL_APPEND(command_queue->root, cmd );
+    LL_APPEND(command_queue->root, cmd);
   }
 
   return CL_SUCCESS;

=== modified file 'lib/CL/clEnqueueReadBufferRect.c'
--- lib/CL/clEnqueueReadBufferRect.c	2012-05-25 16:05:51 +0000
+++ lib/CL/clEnqueueReadBufferRect.c	2012-08-15 20:36:22 +0000
@@ -22,6 +22,7 @@
 */
 
 #include "pocl_cl.h"
+#include "pocl_icd.h"
 #include <assert.h>
 #include <stdio.h>
 
@@ -41,7 +42,7 @@
                         const cl_event *event_wait_list,
                         cl_event *event) CL_API_SUFFIX__VERSION_1_1
 {
-  cl_device_id device_id;
+  cl_device_id device;
   unsigned i;
 
   if (command_queue == NULL)
@@ -58,28 +59,67 @@
       (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;
+  device = command_queue->device;
+
   for (i = 0; i < command_queue->context->num_devices; ++i)
     {
-      if (command_queue->context->devices[i] == device_id)
-        break;
+        if (command_queue->context->devices[i] == device)
+            break;
     }
-
   assert(i < command_queue->context->num_devices);
 
-  device_id->read_rect(device_id->data, ptr, 
-                       buffer->device_ptrs[device_id->dev_id],
-                       buffer_origin, host_origin, region,
-                       buffer_row_pitch, buffer_slice_pitch,
-                       host_row_pitch, host_slice_pitch);
-  
+  if (event != NULL)
+    {
+      *event = (cl_event)malloc(sizeof(struct _cl_event));
+      if (*event == NULL)
+        return CL_OUT_OF_HOST_MEMORY; 
+      POCL_INIT_OBJECT(*event);
+      (*event)->queue = command_queue;
+      POCL_INIT_ICD_OBJECT(*event);
+
+      clRetainCommandQueue (command_queue);
+
+      POCL_PROFILE_QUEUED;
+    }
+
+
+  /* execute directly */
+  /* TODO: enqueue the read_rect if this is a non-blocking read (see
+     clEnqueueReadBuffer) */
+  if (command_queue->properties & CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE)
+    {
+      /* wait for the event in event_wait_list to finish */
+      POCL_ABORT_UNIMPLEMENTED();
+    }
+  else
+    {
+      /* in-order queue - all previously enqueued commands must 
+       * finish before this read */
+      // ensure our buffer is not freed yet
+      clRetainMemObject (buffer);
+      clFinish(command_queue);
+    }
+  POCL_PROFILE_SUBMITTED;
+  POCL_PROFILE_RUNNING;
+
+  /* TODO: offset computation doesn't work in case the ptr is not 
+     a direct pointer */
+  device->read_rect(device->data, ptr, 
+                    buffer->device_ptrs[device->dev_id],
+                    buffer_origin, host_origin, region,
+                    buffer_row_pitch, buffer_slice_pitch,
+                    host_row_pitch, host_slice_pitch);
+
+  POCL_PROFILE_COMPLETE;
+
+  clReleaseMemObject (buffer);
 
   return CL_SUCCESS;
 }

=== modified file 'lib/CL/clEnqueueUnmapMemObject.c'
--- lib/CL/clEnqueueUnmapMemObject.c	2012-05-29 12:16:56 +0000
+++ lib/CL/clEnqueueUnmapMemObject.c	2012-08-15 20:36:22 +0000
@@ -104,6 +104,6 @@
   DL_DELETE(memobj->mappings, mapping);
   memobj->map_count--;
   clReleaseMemObject (memobj);
-  event = NULL;
+
   return CL_SUCCESS;
 }

=== modified file 'lib/CL/clEnqueueWriteBuffer.c'
--- lib/CL/clEnqueueWriteBuffer.c	2012-05-29 12:16:56 +0000
+++ lib/CL/clEnqueueWriteBuffer.c	2012-08-15 20:36:22 +0000
@@ -37,7 +37,7 @@
                      const cl_event *event_wait_list,
                      cl_event *event) CL_API_SUFFIX__VERSION_1_0
 {
-  cl_device_id device_id;
+  cl_device_id device;
   unsigned i;
 
   if (command_queue == NULL)
@@ -53,13 +53,13 @@
       (offset + cb > buffer->size))
     return CL_INVALID_VALUE;
 
-  device_id = command_queue->device;
+  device = command_queue->device;
+
   for (i = 0; i < command_queue->context->num_devices; ++i)
     {
-      if (command_queue->context->devices[i] == device_id)
-	break;
+        if (command_queue->context->devices[i] == device)
+            break;
     }
-
   assert(i < command_queue->context->num_devices);
 
   if (event != NULL)
@@ -76,6 +76,8 @@
     }
 
   /* enqueue the write, or execute directly */
+  /* TODO: why do we implement both? direct execution seems
+     unnecessary. */
   if (blocking_write)
     {
       if (command_queue->properties & CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE)
@@ -94,17 +96,9 @@
       POCL_PROFILE_SUBMITTED;
       POCL_PROFILE_RUNNING;
       /* TODO: fixme. The offset computation must be done at the device driver. */
-      device_id->write(device_id->data, ptr, buffer->device_ptrs[device_id->dev_id]+offset, cb);
+      device->write(device->data, ptr, buffer->device_ptrs[device->dev_id]+offset, cb);
       POCL_PROFILE_COMPLETE;
 
-      if (command_queue->properties & CL_QUEUE_PROFILING_ENABLE && 
-          event != NULL)
-      {
-          (*event)->status = CL_COMPLETE;
-          (*event)->time_end = 
-              command_queue->device->get_timer_value(command_queue->device->data);
-      }
-
       clReleaseMemObject (buffer);
     }
   else
@@ -112,14 +106,15 @@
     _cl_command_node * cmd = malloc(sizeof(_cl_command_node));
     if (cmd == NULL)
       return CL_OUT_OF_HOST_MEMORY;
-    
+
     cmd->type = CL_COMMAND_TYPE_WRITE;
-    cmd->command.write.data = device_id->data;
+    cmd->command.write.data = device->data;
     cmd->command.write.host_ptr = ptr;
     cmd->command.write.device_ptr = buffer->device_ptrs[i]+offset;
     cmd->command.write.cb = cb;
+    cmd->command.write.buffer = buffer;
     cmd->next = NULL;
-    cmd->command.write.buffer = buffer;
+    cmd->event = event ? *event : NULL;
     clRetainMemObject (buffer);
 
     LL_APPEND(command_queue->root, cmd);

=== modified file 'lib/CL/clEnqueueWriteBufferRect.c'
--- lib/CL/clEnqueueWriteBufferRect.c	2012-05-14 11:45:48 +0000
+++ lib/CL/clEnqueueWriteBufferRect.c	2012-08-15 20:36:22 +0000
@@ -40,7 +40,7 @@
                          const cl_event *event_wait_list,
                          cl_event *event) CL_API_SUFFIX__VERSION_1_1
 {
-  cl_device_id device_id;
+  cl_device_id device;
   unsigned i;
 
   if (command_queue == NULL)
@@ -57,43 +57,55 @@
       (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))
   {
       POCL_ABORT_UNIMPLEMENTED();
-#if 0
-    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);
-#endif
     return CL_INVALID_VALUE;
   }
 
-  device_id = command_queue->device;
+  device = command_queue->device;
+
   for (i = 0; i < command_queue->context->num_devices; ++i)
     {
-        if (command_queue->context->devices[i] == device_id)
+        if (command_queue->context->devices[i] == device)
             break;
     }
-
   assert(i < command_queue->context->num_devices);
 
-  device_id->write_rect(device_id->data, ptr, buffer->device_ptrs[device_id->dev_id],
+
+  /* execute directly */
+  /* TODO: enqueue the write_rect if this is a non-blocking read (see
+     clEnqueueWriteBuffer) */
+  if (command_queue->properties & CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE)
+    {
+      /* wait for the event in event_wait_list to finish */
+      POCL_ABORT_UNIMPLEMENTED();
+    }
+  else
+    {
+      /* in-order queue - all previously enqueued commands must 
+       * finish before this read */
+      // ensure our buffer is not freed yet
+      clRetainMemObject (buffer);
+      clFinish(command_queue);
+    }
+  POCL_PROFILE_SUBMITTED;
+  POCL_PROFILE_RUNNING;
+
+  /* TODO: offset computation doesn't work in case the ptr is not 
+     a direct pointer */
+  device->write_rect(device->data, ptr, buffer->device_ptrs[device->dev_id],
                         buffer_origin, host_origin, region,
                         buffer_row_pitch, buffer_slice_pitch,
                         host_row_pitch, host_slice_pitch);
 
+  POCL_PROFILE_COMPLETE;
+
+  clReleaseMemObject (buffer);
+
   return CL_SUCCESS;
 }

=== modified file 'lib/CL/clFinish.c'
--- lib/CL/clFinish.c	2012-05-31 12:01:27 +0000
+++ lib/CL/clFinish.c	2012-08-15 20:36:22 +0000
@@ -41,6 +41,7 @@
         {
         case CL_COMMAND_TYPE_READ:
           POCL_PROFILE_SUBMITTED;
+          POCL_PROFILE_RUNNING;
           command_queue->device->read
             (node->command.read.data, 
              node->command.read.host_ptr, 
@@ -51,6 +52,7 @@
           break;
         case CL_COMMAND_TYPE_WRITE:
           POCL_PROFILE_SUBMITTED;
+          POCL_PROFILE_RUNNING;
           command_queue->device->write
             (node->command.write.data, 
              node->command.write.host_ptr, 
@@ -61,6 +63,7 @@
           break;
         case CL_COMMAND_TYPE_COPY:
           POCL_PROFILE_SUBMITTED;
+          POCL_PROFILE_RUNNING;
           command_queue->device->copy
             (node->command.copy.data, 
              node->command.copy.src_ptr, 
@@ -71,9 +74,11 @@
           clReleaseMemObject (node->command.copy.dst_buffer);
           break;
         case CL_COMMAND_TYPE_RUN:
-          POCL_PROFILE_SUBMITTED;
           assert (*event == node->event);
+          POCL_PROFILE_SUBMITTED;
+          POCL_PROFILE_RUNNING;
           command_queue->device->run(node->command.run.data, node);
+          POCL_PROFILE_COMPLETE;
           for (i = 0; i < node->command.run.arg_buffer_count; ++i)
             {
               cl_mem buf = node->command.run.arg_buffers[i];

=== modified file 'lib/CL/clGetDeviceInfo.c'
--- lib/CL/clGetDeviceInfo.c	2012-08-08 12:45:17 +0000
+++ lib/CL/clGetDeviceInfo.c	2012-08-15 20:36:22 +0000
@@ -126,23 +126,17 @@
     }
     
   case CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR:
-    POCL_RETURN_DEVICE_INFO_WITH_IMPL_CHECK(cl_uint, device->preferred_vector_width_char);   
-    
+    POCL_RETURN_DEVICE_INFO_WITH_IMPL_CHECK(cl_uint, device->preferred_vector_width_char);
   case CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT:
     POCL_RETURN_DEVICE_INFO_WITH_IMPL_CHECK(cl_uint, device->preferred_vector_width_short);
-    
   case CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT:
     POCL_RETURN_DEVICE_INFO_WITH_IMPL_CHECK(cl_uint, device->preferred_vector_width_int);
-
   case CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG:
     POCL_RETURN_DEVICE_INFO_WITH_IMPL_CHECK(cl_uint, device->preferred_vector_width_long);
-    
   case CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT:
     POCL_RETURN_DEVICE_INFO_WITH_IMPL_CHECK(cl_uint, device->preferred_vector_width_float);
-    
   case CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE:
     POCL_RETURN_DEVICE_INFO_WITH_IMPL_CHECK(cl_uint, device->preferred_vector_width_double);
-
   case CL_DEVICE_MAX_CLOCK_FREQUENCY               :
     POCL_RETURN_DEVICE_INFO_WITH_IMPL_CHECK(cl_uint, device->max_clock_frequency);
   case CL_DEVICE_ADDRESS_BITS                      :
@@ -186,7 +180,7 @@
   case CL_DEVICE_MAX_CONSTANT_ARGS                 : 
     POCL_RETURN_DEVICE_INFO_WITH_IMPL_CHECK(cl_uint, device->max_constant_args);
   case CL_DEVICE_LOCAL_MEM_TYPE                    :
-    POCL_RETURN_DEVICE_INFO_WITH_IMPL_CHECK(cl_uint, device->local_mem_size);
+    POCL_RETURN_DEVICE_INFO_WITH_IMPL_CHECK(cl_uint, device->local_mem_type);
   case CL_DEVICE_LOCAL_MEM_SIZE:
     POCL_RETURN_DEVICE_INFO_WITH_IMPL_CHECK(cl_ulong, device->local_mem_size);
   case CL_DEVICE_ERROR_CORRECTION_SUPPORT          :
@@ -218,28 +212,34 @@
 
   case CL_DEVICE_EXTENSIONS                        : 
     POCL_RETURN_DEVICE_INFO_STR("cl_khr_fp16");
-  case CL_DEVICE_PLATFORM                          : break;
-  case CL_DEVICE_DOUBLE_FP_CONFIG                  : break;
+  case CL_DEVICE_PLATFORM                          :
+    {
+      /* Return the first platform id, assuming this is the only
+         platform id (which is currently always the case for pocl) */
+      cl_platform_id platform_id;
+      clGetPlatformIDs(1, &platform_id, NULL);
+      POCL_RETURN_DEVICE_INFO(cl_platform_id, platform_id);
+    }
+  case CL_DEVICE_DOUBLE_FP_CONFIG                  :
+    POCL_RETURN_DEVICE_INFO_WITH_IMPL_CHECK(cl_ulong, device->double_fp_config);
   case CL_DEVICE_HALF_FP_CONFIG                    : break;
-  case CL_DEVICE_PREFERRED_VECTOR_WIDTH_HALF       : break;
+  case CL_DEVICE_PREFERRED_VECTOR_WIDTH_HALF       :
+    POCL_RETURN_DEVICE_INFO_WITH_IMPL_CHECK(cl_uint, device->preferred_vector_width_half);
   case CL_DEVICE_HOST_UNIFIED_MEMORY               : break;
-    /* TODO: figure out what the difference between preferred and native
-       widths are. And why there is no struct fields 'native_vector...' */
   case CL_DEVICE_NATIVE_VECTOR_WIDTH_CHAR          : 
-    POCL_RETURN_DEVICE_INFO_WITH_IMPL_CHECK(cl_uint, device->preferred_vector_width_char);
+    POCL_RETURN_DEVICE_INFO_WITH_IMPL_CHECK(cl_uint, device->native_vector_width_char);
   case CL_DEVICE_NATIVE_VECTOR_WIDTH_SHORT         :
-    POCL_RETURN_DEVICE_INFO_WITH_IMPL_CHECK(cl_uint, device->preferred_vector_width_short);
+    POCL_RETURN_DEVICE_INFO_WITH_IMPL_CHECK(cl_uint, device->native_vector_width_short);
   case CL_DEVICE_NATIVE_VECTOR_WIDTH_INT           : 
-    POCL_RETURN_DEVICE_INFO_WITH_IMPL_CHECK(cl_uint, device->preferred_vector_width_int);
+    POCL_RETURN_DEVICE_INFO_WITH_IMPL_CHECK(cl_uint, device->native_vector_width_int);
   case CL_DEVICE_NATIVE_VECTOR_WIDTH_LONG          : 
-    POCL_RETURN_DEVICE_INFO_WITH_IMPL_CHECK(cl_uint, device->preferred_vector_width_long);
+    POCL_RETURN_DEVICE_INFO_WITH_IMPL_CHECK(cl_uint, device->native_vector_width_long);
   case CL_DEVICE_NATIVE_VECTOR_WIDTH_FLOAT         : 
-    POCL_RETURN_DEVICE_INFO_WITH_IMPL_CHECK(cl_uint, device->preferred_vector_width_float);
+    POCL_RETURN_DEVICE_INFO_WITH_IMPL_CHECK(cl_uint, device->native_vector_width_float);
   case CL_DEVICE_NATIVE_VECTOR_WIDTH_DOUBLE        : 
-    POCL_RETURN_DEVICE_INFO_WITH_IMPL_CHECK(cl_uint, device->preferred_vector_width_double);
+    POCL_RETURN_DEVICE_INFO_WITH_IMPL_CHECK(cl_uint, device->native_vector_width_double);
   case CL_DEVICE_NATIVE_VECTOR_WIDTH_HALF          : 
-    /* TODO: why is there no preferred_vector_width_half? */
-    POCL_RETURN_DEVICE_INFO(cl_uint, 0);
+    POCL_RETURN_DEVICE_INFO_WITH_IMPL_CHECK(cl_uint, device->native_vector_width_half);
   case CL_DEVICE_OPENCL_C_VERSION                  :
     POCL_RETURN_DEVICE_INFO_STR("1.2");
   }

=== modified file 'lib/CL/clGetKernelWorkGroupInfo.c'
--- lib/CL/clGetKernelWorkGroupInfo.c	2012-03-28 11:25:45 +0000
+++ lib/CL/clGetKernelWorkGroupInfo.c	2012-08-15 20:36:22 +0000
@@ -65,7 +65,7 @@
       POCL_ABORT_UNIMPLEMENTED();
       
     case CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE:
-      POCL_RETURN_KERNEL_WG_INFO(size_t, device->preferred_wg_size_multiple);        
+      POCL_RETURN_KERNEL_WG_INFO(size_t, device->preferred_wg_size_multiple);
       
     case CL_KERNEL_LOCAL_MEM_SIZE:
       POCL_RETURN_KERNEL_WG_INFO(cl_ulong, device->local_mem_size);

=== modified file 'lib/CL/clReleaseProgram.c'
--- lib/CL/clReleaseProgram.c	2012-05-31 12:01:27 +0000
+++ lib/CL/clReleaseProgram.c	2012-08-15 20:36:22 +0000
@@ -58,7 +58,7 @@
       free (program->binary_sizes);
 
       env = getenv ("POCL_LEAVE_TEMP_DIRS");
-      if (!(env != NULL && strlen (env) == 1 && env[0] == '1') &&
+      if (!(env != NULL && strcmp (env, "1") == 0) &&
           getenv("POCL_TEMP_DIR") == NULL)
         {
           remove_directory (program->temp_dir);

=== modified file 'lib/CL/devices/basic/basic.h'
--- lib/CL/devices/basic/basic.h	2012-08-08 14:54:32 +0000
+++ lib/CL/devices/basic/basic.h	2012-08-15 20:36:22 +0000
@@ -46,6 +46,15 @@
   POCL_DEVICES_PREFERRED_VECTOR_WIDTH_LONG  , /* preferred_vector_width_long */ \
   POCL_DEVICES_PREFERRED_VECTOR_WIDTH_FLOAT , /* preferred_vector_width_float */ \
   POCL_DEVICES_PREFERRED_VECTOR_WIDTH_DOUBLE, /* preferred_vector_width_double */ \
+  POCL_DEVICES_PREFERRED_VECTOR_WIDTH_HALF  , /* preferred_vector_width_half */ \
+  /* TODO: figure out what the difference between preferred and native widths are. */ \
+  POCL_DEVICES_PREFERRED_VECTOR_WIDTH_CHAR  , /* preferred_vector_width_char */ \
+  POCL_DEVICES_PREFERRED_VECTOR_WIDTH_SHORT , /* preferred_vector_width_short */ \
+  POCL_DEVICES_PREFERRED_VECTOR_WIDTH_INT   , /* preferred_vector_width_int */ \
+  POCL_DEVICES_PREFERRED_VECTOR_WIDTH_LONG  , /* preferred_vector_width_long */ \
+  POCL_DEVICES_PREFERRED_VECTOR_WIDTH_FLOAT , /* preferred_vector_width_float */ \
+  POCL_DEVICES_PREFERRED_VECTOR_WIDTH_DOUBLE, /* preferred_vector_width_double */ \
+  POCL_DEVICES_PREFERRED_VECTOR_WIDTH_HALF  , /* preferred_vector_width_half */ \
   0, /* max_clock_frequency */						\
   0, /* address_bits */							\
   0, /* max_mem_alloc_size */						\
@@ -62,6 +71,7 @@
   0, /* mem_base_addr_align */						\
   0, /* min_data_type_align_size */					\
   CL_FP_ROUND_TO_NEAREST | CL_FP_INF_NAN, /* single_fp_config */	\
+  CL_FP_ROUND_TO_NEAREST | CL_FP_INF_NAN, /* double_fp_config */	\
   CL_NONE, /* global_mem_cache_type */					\
   0, /* global_mem_cacheline_size */					\
   0, /* global_mem_cache_size */					\

=== modified file 'lib/CL/devices/common.h'
--- lib/CL/devices/common.h	2012-04-24 13:12:25 +0000
+++ lib/CL/devices/common.h	2012-08-15 20:36:22 +0000
@@ -70,5 +70,6 @@
 #endif
 /* Half is internally represented as short */
 #define POCL_DEVICES_PREFERRED_VECTOR_WIDTH_HALF POCL_DEVICES_PREFERRED_VECTOR_WIDTH_SHORT
+#define POCL_DEVICES_NATIVE_VECTOR_WIDTH_HALF POCL_DEVICES_NATIVE_VECTOR_WIDTH_SHORT
 
 #endif

=== modified file 'lib/CL/devices/devices.c'
--- lib/CL/devices/devices.c	2012-05-31 12:01:27 +0000
+++ lib/CL/devices/devices.c	2012-08-15 20:36:22 +0000
@@ -57,7 +57,7 @@
 pocl_init_devices()
 {
   const char *device_list;
-  char *ptr, *tofree, *token, *saveptr, *saveptr2;
+  char *ptr, *tofree, *token, *saveptr;
   int i, devcount;
   if (pocl_num_devices > 0)
     return;
@@ -83,7 +83,7 @@
 
   ptr = tofree = strdup(device_list);
   devcount = 0;
-  while ((token = strtok_r (ptr, " ", &saveptr2)) != NULL)
+  while ((token = strtok_r (ptr, " ", &saveptr)) != NULL)
     {
       struct _cl_device_id* device_type = NULL;
 

=== modified file 'lib/CL/devices/pthread/pocl-pthread.h'
--- lib/CL/devices/pthread/pocl-pthread.h	2012-08-08 14:54:32 +0000
+++ lib/CL/devices/pthread/pocl-pthread.h	2012-08-15 20:36:22 +0000
@@ -50,6 +50,15 @@
   POCL_DEVICES_PREFERRED_VECTOR_WIDTH_LONG  , /* preferred_vector_width_long */ \
   POCL_DEVICES_PREFERRED_VECTOR_WIDTH_FLOAT , /* preferred_vector_width_float */ \
   POCL_DEVICES_PREFERRED_VECTOR_WIDTH_DOUBLE, /* preferred_vector_width_double */ \
+  POCL_DEVICES_PREFERRED_VECTOR_WIDTH_HALF,   /* preferred_vector_width_half */ \
+  /* TODO: figure out what the difference between preferred and native widths are. */ \
+  POCL_DEVICES_PREFERRED_VECTOR_WIDTH_CHAR  , /* preferred_vector_width_char */ \
+  POCL_DEVICES_PREFERRED_VECTOR_WIDTH_SHORT , /* preferred_vector_width_short */ \
+  POCL_DEVICES_PREFERRED_VECTOR_WIDTH_INT   , /* preferred_vector_width_int */ \
+  POCL_DEVICES_PREFERRED_VECTOR_WIDTH_LONG  , /* preferred_vector_width_long */ \
+  POCL_DEVICES_PREFERRED_VECTOR_WIDTH_FLOAT , /* preferred_vector_width_float */ \
+  POCL_DEVICES_PREFERRED_VECTOR_WIDTH_DOUBLE, /* preferred_vector_width_double */ \
+  POCL_DEVICES_PREFERRED_VECTOR_WIDTH_HALF  , /* preferred_vector_width_half */ \
   0, /* max_clock_frequency */						\
   0, /* address_bits */							\
   0, /* max_mem_alloc_size */						\
@@ -66,6 +75,7 @@
   0, /* mem_base_addr_align */						\
   0, /* min_data_type_align_size */					\
   CL_FP_ROUND_TO_NEAREST | CL_FP_INF_NAN, /* single_fp_config */	\
+  CL_FP_ROUND_TO_NEAREST | CL_FP_INF_NAN, /* double_fp_config */	\
   CL_NONE, /* global_mem_cache_type */					\
   0, /* global_mem_cacheline_size */					\
   0, /* global_mem_cache_size */					\

=== modified file 'lib/CL/pocl_cl.h'
--- lib/CL/pocl_cl.h	2012-08-08 14:54:32 +0000
+++ lib/CL/pocl_cl.h	2012-08-15 20:36:22 +0000
@@ -25,6 +25,7 @@
 #define POCL_CL_H
 
 #include "config.h"
+#include <assert.h>
 #include <stdio.h>
 #include <ltdl.h>
 #include <pthread.h>
@@ -150,6 +151,14 @@
   cl_uint preferred_vector_width_long;
   cl_uint preferred_vector_width_float;
   cl_uint preferred_vector_width_double;
+  cl_uint preferred_vector_width_half;
+  cl_uint native_vector_width_char;
+  cl_uint native_vector_width_short;
+  cl_uint native_vector_width_int;
+  cl_uint native_vector_width_long;
+  cl_uint native_vector_width_float;
+  cl_uint native_vector_width_double;
+  cl_uint native_vector_width_half;
   cl_uint max_clock_frequency;
   cl_uint address_bits;
   cl_ulong max_mem_alloc_size;
@@ -166,6 +175,7 @@
   cl_uint mem_base_addr_align;
   cl_uint min_data_type_align_size;
   cl_device_fp_config single_fp_config;
+  cl_device_fp_config double_fp_config;
   cl_device_mem_cache_type global_mem_cache_type;
   cl_uint global_mem_cacheline_size;
   cl_ulong global_mem_cache_size;
@@ -394,12 +404,12 @@
       }                                                                 \
   } while (0)                                                           \
 
-
 #define POCL_PROFILE_SUBMITTED                                          \
   do {                                                                  \
     if (command_queue->properties & CL_QUEUE_PROFILING_ENABLE &&        \
         event != NULL && (*event) != NULL)                              \
       {                                                                 \
+        assert((*event)->status = CL_QUEUED);                           \
         (*event)->status = CL_SUBMITTED;                                \
         (*event)->time_submit =                                         \
           command_queue->device->get_timer_value(command_queue->device->data); \
@@ -411,6 +421,7 @@
     if (command_queue->properties & CL_QUEUE_PROFILING_ENABLE &&        \
         event != NULL && (*event) != NULL)                              \
       {                                                                 \
+        assert((*event)->status = CL_SUBMITTED);                        \
         (*event)->status = CL_RUNNING;                                  \
         (*event)->time_start =                                          \
           command_queue->device->get_timer_value(command_queue->device->data); \
@@ -422,6 +433,7 @@
     if (command_queue->properties & CL_QUEUE_PROFILING_ENABLE &&        \
         event != NULL && (*event) != NULL)                              \
       {                                                                 \
+        assert((*event)->status = CL_RUNNING);                          \
         (*event)->status = CL_COMPLETE;                                 \
         (*event)->time_end =                                            \
           command_queue->device->get_timer_value(command_queue->device->data); \

=== added file 'lib/kernel/atomics.cl'
--- lib/kernel/atomics.cl	1970-01-01 00:00:00 +0000
+++ lib/kernel/atomics.cl	2012-08-15 20:36:22 +0000
@@ -0,0 +1,149 @@
+/* OpenCL built-in library: atomic operations
+
+   Copyright (c) 2012 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.
+*/
+
+
+
+// Repeat the content of this file several times with different values
+// for Q, T, and U:
+#if !defined(Q)
+
+#  define Q __global
+#  include "atomics.cl"
+#  undef Q
+
+#  define Q __local
+#  include "atomics.cl"
+#  undef Q
+
+#elif !defined(T)
+
+#  define T int
+#  define MIN __sync_fetch_and_min
+#  define MAX __sync_fetch_and_max
+#  include "atomics.cl"
+#  undef T
+#  undef MIN
+#  undef MAX
+
+#  define T uint
+#  define MIN __sync_fetch_and_umin
+#  define MAX __sync_fetch_and_umax
+#  include "atomics.cl"
+#  undef T
+#  undef MIN
+#  undef MAX
+
+
+// xchg is also supported for float as a special case
+__attribute__((overloadable))
+float atomic_xchg(volatile Q float *p, float val)
+{
+  // NOTE: We compare the float as int here...
+  return __atomic_exchange_n((volatile int*)p, val, __ATOMIC_RELAXED);
+}
+
+#else
+
+
+
+// basic
+
+// read, add, store
+__attribute__((overloadable))
+T atomic_add(volatile Q T *p, T val)
+{
+  return __sync_fetch_and_add((volatile T*)p, val, __ATOMIC_RELAXED);
+}
+
+// read, subtract, store
+__attribute__((overloadable))
+T atomic_sub(volatile Q T *p, T val)
+{
+  return __sync_fetch_and_sub(p, val, __ATOMIC_RELAXED);
+}
+
+// read, swap, store
+__attribute__((overloadable))
+T atomic_xchg(volatile Q T *p, T val)
+{
+  return __atomic_exchange_n(p, val, __ATOMIC_RELAXED);
+}
+
+// read, increment, store
+__attribute__((overloadable))
+T atomic_inc(volatile Q T *p)
+{
+  return atomic_add(p, (T)1);
+}
+
+// read, decrement, store
+__attribute__((overloadable))
+T atomic_dec(volatile Q T *p)
+{
+  return atomic_sub(p, (T)1);
+}
+
+// read, store
+__attribute__((overloadable))
+T atomic_cmpxchg(volatile Q T *p, T cmp, T val)
+{
+  __atomic_compare_exchange_n(p, &cmp, val, false,
+                              __ATOMIC_RELAXED, __ATOMIC_RELAXED);
+  return cmp;
+}
+
+// extended
+
+__attribute__((overloadable))
+T atomic_min(volatile Q T *p, T val)
+{
+  return MIN((volatile T*)p, val);
+}
+
+__attribute__((overloadable))
+T atomic_max(volatile Q T *p, T val)
+{
+  return MAX((volatile T*)p, val);
+}
+
+__attribute__((overloadable))
+T atomic_and(volatile Q T *p, T val)
+{
+  return __sync_fetch_and_and(p, val, __ATOMIC_RELAXED);
+}
+
+__attribute__((overloadable))
+T atomic_or(volatile Q T *p, T val)
+{
+  return __sync_fetch_and_or(p, val, __ATOMIC_RELAXED);
+}
+
+__attribute__((overloadable))
+T atomic_xor(volatile Q T *p, T val)
+{
+  return __sync_fetch_and_xor(p, val, __ATOMIC_RELAXED);
+}
+
+
+
+#endif

=== modified file 'lib/kernel/hadd.cl'
--- lib/kernel/hadd.cl	2011-10-26 13:03:37 +0000
+++ lib/kernel/hadd.cl	2012-08-15 20:36:22 +0000
@@ -23,7 +23,5 @@
 
 #include "templates.h"
 
-// This could do with some testing
-// This could probably also be optimised (i.e. the ?: operators eliminated)
 DEFINE_EXPR_G_GG(hadd,
                  (a >> (sgtype)1) + (b >> (sgtype)1) + (a & b & (gtype)1))

=== modified file 'lib/kernel/rhadd.cl'
--- lib/kernel/rhadd.cl	2011-10-26 13:03:37 +0000
+++ lib/kernel/rhadd.cl	2012-08-15 20:36:22 +0000
@@ -23,7 +23,5 @@
 
 #include "templates.h"
 
-// This could do with some testing
-// This could probably also be optimised (i.e. the ?: operators eliminated)
 DEFINE_EXPR_G_GG(rhadd,
                  (a >> (sgtype)1) + (b >> (sgtype)1) + ((a | b) & (gtype)1))

=== modified file 'lib/kernel/sources.mk'
--- lib/kernel/sources.mk	2012-06-04 12:15:18 +0000
+++ lib/kernel/sources.mk	2012-08-15 20:36:22 +0000
@@ -1,144 +1,146 @@
 # Nodist here because these files should be included
 # to the distribution only once, from the root kernel
 # makefile.
-nodist_libkernel_a_SOURCES = templates.h		\
-                      barrier.ll	                \
-					  image.h \
-                      get_work_dim.c		\
-                      get_global_size.c		\
-                      get_global_id.c		\
-                      get_local_size.c		\
-                      get_local_id.c		\
-                      get_num_groups.c		\
-                      get_group_id.c		\
-                      get_global_offset.c	\
-                      as_type.cl		\
-                      convert_type.cl		\
-                      acos.cl			\
-                      acosh.cl			\
-                      acospi.cl			\
-                      asin.cl			\
-                      asinh.cl			\
-                      asinpi.cl			\
-                      atan.cl			\
-                      atan2.cl			\
-                      atan2pi.cl		\
-                      atanh.cl			\
-                      atanpi.cl			\
-                      cbrt.cl			\
-                      ceil.cl			\
-                      copysign.cl		\
-                      cos.cl			\
-                      cosh.cl			\
-                      cospi.cl			\
-                      erfc.cl			\
-                      erf.cl			\
-                      exp.cl			\
-                      exp2.cl			\
-                      exp10.cl			\
-                      expm1.cl			\
-                      fabs.cl			\
-                      fdim.cl			\
-                      floor.cl			\
-                      fma.cl			\
-                      fmax.cl			\
-                      fmin.cl			\
-                      fmod.cl			\
-                      fract.cl			\
-                      hypot.cl			\
-                      ilogb.cl			\
-                      ldexp.cl			\
-                      lgamma.cl			\
-                      log.cl			\
-                      log2.cl			\
-                      log10.cl			\
-                      log1p.cl			\
-                      logb.cl			\
-                      mad.cl			\
-                      maxmag.cl			\
-                      minmag.cl			\
-                      nan.cl			\
-                      nextafter.cl		\
-                      pow.cl			\
-                      pown.cl			\
-                      powr.cl			\
-                      remainder.cl		\
-                      rint.cl			\
-                      rootn.cl			\
-                      round.cl			\
-                      rsqrt.cl			\
-                      sin.cl			\
-                      sincos.cl			\
-                      sinh.cl			\
-                      sinpi.cl			\
-                      sqrt.cl			\
-                      tan.cl			\
-                      tanh.cl			\
-                      tanpi.cl			\
-                      tgamma.cl			\
-                      trunc.cl			\
-                      divide.cl			\
-                      recip.cl			\
-                      abs.cl			\
-                      abs_diff.cl		\
-                      add_sat.cl		\
-                      hadd.cl			\
-                      rhadd.cl			\
-                      clamp.cl			\
-                      clz.cl			\
-                      mad_hi.cl			\
-                      mad_sat.cl		\
-                      max.cl			\
-                      min.cl			\
-                      mul_hi.cl			\
-                      rotate.cl			\
-                      sub_sat.cl		\
-                      upsample.cl		\
-                      popcount.cl		\
-                      mad24.cl			\
-                      mul24.cl			\
-                      degrees.cl		\
-                      mix.cl			\
-                      radians.cl		\
-                      step.cl			\
-                      smoothstep.cl		\
-                      sign.cl			\
-                      cross.cl			\
-                      dot.cl			\
-                      distance.cl		\
-                      length.cl			\
-                      normalize.cl		\
-                      fast_distance.cl		\
-                      fast_length.cl		\
-                      fast_normalize.cl		\
-                      isequal.cl		\
-                      isnotequal.cl		\
-                      isgreater.cl		\
-                      isgreaterequal.cl		\
-                      isless.cl			\
-                      islessequal.cl		\
-                      islessgreater.cl		\
-                      isfinite.cl		\
-                      isinf.cl			\
-                      isnan.cl			\
-                      isnormal.cl		\
-                      isordered.cl		\
-                      isunordered.cl		\
-                      signbit.cl		\
-                      any.cl			\
-                      all.cl			\
-                      bitselect.cl		\
-                      select.cl			\
-                      vload.cl			\
-                      vstore.cl			\
-                      vload_half.cl		\
-                      vstore_half.cl 	\
-                      async_work_group_copy.cl \
-                      wait_group_events.cl      \
-                      read_image.cl      \
-                      write_image.cl     \
-                      get_image_width.cl      \
-                      get_image_height.cl     
+nodist_libkernel_a_SOURCES =			\
+	templates.h				\
+	barrier.ll				\
+	image.h					\
+	get_work_dim.c				\
+	get_global_size.c			\
+	get_global_id.c				\
+	get_local_size.c			\
+	get_local_id.c				\
+	get_num_groups.c			\
+	get_group_id.c				\
+	get_global_offset.c			\
+	as_type.cl				\
+	atomics.cl				\
+	convert_type.cl				\
+	acos.cl					\
+	acosh.cl				\
+	acospi.cl				\
+	asin.cl					\
+	asinh.cl				\
+	asinpi.cl				\
+	atan.cl					\
+	atan2.cl				\
+	atan2pi.cl				\
+	atanh.cl				\
+	atanpi.cl				\
+	cbrt.cl					\
+	ceil.cl					\
+	copysign.cl				\
+	cos.cl					\
+	cosh.cl					\
+	cospi.cl				\
+	erfc.cl					\
+	erf.cl					\
+	exp.cl					\
+	exp2.cl					\
+	exp10.cl				\
+	expm1.cl				\
+	fabs.cl					\
+	fdim.cl					\
+	floor.cl				\
+	fma.cl					\
+	fmax.cl					\
+	fmin.cl					\
+	fmod.cl					\
+	fract.cl				\
+	hypot.cl				\
+	ilogb.cl				\
+	ldexp.cl				\
+	lgamma.cl				\
+	log.cl					\
+	log2.cl					\
+	log10.cl				\
+	log1p.cl				\
+	logb.cl					\
+	mad.cl					\
+	maxmag.cl				\
+	minmag.cl				\
+	nan.cl					\
+	nextafter.cl				\
+	pow.cl					\
+	pown.cl					\
+	powr.cl					\
+	remainder.cl				\
+	rint.cl					\
+	rootn.cl				\
+	round.cl				\
+	rsqrt.cl				\
+	sin.cl					\
+	sincos.cl				\
+	sinh.cl					\
+	sinpi.cl				\
+	sqrt.cl					\
+	tan.cl					\
+	tanh.cl					\
+	tanpi.cl				\
+	tgamma.cl				\
+	trunc.cl				\
+	divide.cl				\
+	recip.cl				\
+	abs.cl					\
+	abs_diff.cl				\
+	add_sat.cl				\
+	hadd.cl					\
+	rhadd.cl				\
+	clamp.cl				\
+	clz.cl					\
+	mad_hi.cl				\
+	mad_sat.cl				\
+	max.cl					\
+	min.cl					\
+	mul_hi.cl				\
+	rotate.cl				\
+	sub_sat.cl				\
+	upsample.cl				\
+	popcount.cl				\
+	mad24.cl				\
+	mul24.cl				\
+	degrees.cl				\
+	mix.cl					\
+	radians.cl				\
+	step.cl					\
+	smoothstep.cl				\
+	sign.cl					\
+	cross.cl				\
+	dot.cl					\
+	distance.cl				\
+	length.cl				\
+	normalize.cl				\
+	fast_distance.cl			\
+	fast_length.cl				\
+	fast_normalize.cl			\
+	isequal.cl				\
+	isnotequal.cl				\
+	isgreater.cl				\
+	isgreaterequal.cl			\
+	isless.cl				\
+	islessequal.cl				\
+	islessgreater.cl			\
+	isfinite.cl				\
+	isinf.cl				\
+	isnan.cl				\
+	isnormal.cl				\
+	isordered.cl				\
+	isunordered.cl				\
+	signbit.cl				\
+	any.cl					\
+	all.cl					\
+	bitselect.cl				\
+	select.cl				\
+	vload.cl				\
+	vstore.cl				\
+	vload_half.cl				\
+	vstore_half.cl				\
+	async_work_group_copy.cl		\
+	wait_group_events.cl			\
+	read_image.cl				\
+	write_image.cl				\
+	get_image_width.cl			\
+	get_image_height.cl     
 
 barrier.o: barrier.ll
 	$(LLVM_AS) -o $@ $<

=== modified file 'lib/kernel/sub_sat.cl'
--- lib/kernel/sub_sat.cl	2011-10-26 19:49:23 +0000
+++ lib/kernel/sub_sat.cl	2012-08-15 20:36:22 +0000
@@ -30,8 +30,6 @@
 //    ushort   __builtin_ia32_psubusw128
 // Other types don't seem to be supported.
 
-// This could do with some testing
-// This could probably also be optimised (i.e. the ?: operators eliminated)
 DEFINE_EXPR_G_GG(sub_sat,
                  (sgtype)-1 < (sgtype)0 ?
                  /* signed */

=== modified file 'lib/llvmopencl/Makefile.am'
--- lib/llvmopencl/Makefile.am	2012-06-15 12:26:40 +0000
+++ lib/llvmopencl/Makefile.am	2012-08-15 20:36:22 +0000
@@ -40,4 +40,4 @@
 						ImplicitLoopBarriers.h ImplicitLoopBarriers.cc \
 						WorkItemAliasAnalysis.cc WIVectorize.cc
 
-#llvmopencl_la_LIBADD = @LIBS_LLVMTRANSFORMUTILS@
\ No newline at end of file
+#llvmopencl_la_LIBADD = @LIBS_LLVMTRANSFORMUTILS@

=== modified file 'tests/regression/test_constant_array.cpp'
--- tests/regression/test_constant_array.cpp	2012-08-02 14:18:02 +0000
+++ tests/regression/test_constant_array.cpp	2012-08-15 20:36:22 +0000
@@ -29,6 +29,7 @@
 #include <cstdio>
 #include <cstdlib>
 #include <iostream>
+#include <unistd.h>
 
 #define WORK_ITEMS 1
 

=== modified file 'tests/regression/test_infinite_loop.cpp'
--- tests/regression/test_infinite_loop.cpp	2012-08-02 14:18:02 +0000
+++ tests/regression/test_infinite_loop.cpp	2012-08-15 20:36:22 +0000
@@ -28,6 +28,7 @@
 #include <cstdio>
 #include <cstdlib>
 #include <iostream>
+#include <unistd.h>
 
 #define WORK_ITEMS 1