Updated the logic to choose the local size automatically for
embarrassingly parallel kernel launches:
- Split first at the different dimensions while trying to retain
  the optimal WG width (for WG vectorization).
- Try not to create too big local size due to the WI context data
  overheads it causes.
- Fixed a couple of test cases which erroneously left the WG size
  to be automatically set and broke down when it was set to something
  unexpected
diff --git a/lib/CL/clEnqueueNDRangeKernel.c b/lib/CL/clEnqueueNDRangeKernel.c
index a0951cb..1e15848 100644
--- a/lib/CL/clEnqueueNDRangeKernel.c
+++ b/lib/CL/clEnqueueNDRangeKernel.c
@@ -112,6 +112,15 @@ POname(clEnqueueNDRangeKernel)(cl_command_queue command_queue,
     }
   else
     {
+      /* Embarrassingly parallel kernel with a free work-group
+         size. Try to figure out one which utilizes all the
+         resources efficiently. Assume work-groups are scheduled
+         to compute units, so try to split it to a number of
+         work groups at the equal to the number of CUs, while still
+         trying to respect the preferred WG size multiple (for better
+         SIMD instruction utilization).
+      */
+
       size_t preferred_wg_multiple;
       cl_int retval =
         POname(clGetKernelWorkGroupInfo)
@@ -119,30 +128,83 @@ POname(clEnqueueNDRangeKernel)(cl_command_queue command_queue,
          CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE,
          sizeof (size_t), &preferred_wg_multiple, NULL);

-      local_x = local_y = local_z = 1;
-      if (retval == CL_SUCCESS)
-        {
-          /* Find the largest multiple of the preferred wg multiple.
-             E.g. if the preferred is 8 it doesn't work with a
-             global size of 20. However, 4 is better than 1 in that
-             case because it still enables wi-parallelization. */
-          while (preferred_wg_multiple >= 1)
-            {
-              if (global_x % preferred_wg_multiple == 0 &&
-                  preferred_wg_multiple <= global_x)
-                {
-                  local_x = preferred_wg_multiple;
-                  break;
-                }
-              preferred_wg_multiple /= 2;
-            }
-        }
+      POCL_MSG_PRINT_INFO("Preferred WG size multiple %zu\n",
+                          preferred_wg_multiple);
+
+      local_x = global_x;
+      local_y = global_y;
+      local_z = global_z;
+
+      /* First try to split a dimension with the WG multiple
+         to make it still be divisible with the WG multiple. */
+      do {
+        /* Split the dimension, but avoid ending up with a dimension that
+           is not multiple of the wanted size. */
+        if (local_x > 1 && local_x % 2 == 0 &&
+            (local_x / 2) % preferred_wg_multiple == 0)
+          {
+            local_x /= 2;
+            continue;
+          }
+        else if (local_y > 1 && local_y % 2 == 0 &&
+                 (local_y / 2) % preferred_wg_multiple == 0)
+          {
+            local_y /= 2;
+            continue;
+          }
+        else if (local_z > 1 && local_z % 2 == 0 &&
+                 (local_z / 2) % preferred_wg_multiple == 0)
+          {
+            local_z /= 2;
+            continue;
+          }
+
+        /* Next find out a dimension that is not a multiple anyways,
+           so one cannot nicely vectorize over it, and set it to one. */
+        if (local_z > 1 && local_z % preferred_wg_multiple != 0)
+          {
+            local_z = 1;
+            continue;
+          }
+        else if (local_y > 1 && local_y % preferred_wg_multiple != 0)
+          {
+            local_y = 1;
+            continue;
+          }
+        else if (local_z > 1 && local_z % preferred_wg_multiple != 0)
+          {
+            local_z = 1;
+            continue;
+          }
+
+        /* Finally, start setting them to zero starting from the Z
+           dimension. */
+        if (local_z > 1)
+          {
+            local_z = 1;
+            continue;
+          }
+        else if (local_y > 1)
+          {
+            local_y = 1;
+            continue;
+          }
+        else if (local_x > 1)
+          {
+            local_x = 1;
+            continue;
+          }
+      }
+      while (local_x * local_y * local_z >
+             command_queue->device->max_work_group_size);
     }

-#ifdef DEBUG_NDRANGE
-  printf("### queueing kernel %s for dimensions %zu x %zu x %zu...",
-         kernel->function_name, local_x, local_y, local_z);
-#endif
+  POCL_MSG_PRINT_INFO("Qeueing kernel %s with local size %u x %u x %u...\n",
+                      kernel->function_name, local_x, local_y, local_z);
+
+  POCL_MSG_PRINT_INFO("Qeueing kernel %s with groups %u x %u x %u...\n",
+                      kernel->function_name, global_x / local_x,
+                      global_y / local_y, global_z / local_z);

   POCL_RETURN_ERROR_ON((local_x * local_y * local_z > command_queue->device->max_work_group_size),
     CL_INVALID_WORK_GROUP_SIZE, "Local worksize dimensions exceed device's max workgroup size\n");
diff --git a/lib/CL/devices/basic/basic.c b/lib/CL/devices/basic/basic.c
index 216118f..35f70a0 100644
--- a/lib/CL/devices/basic/basic.c
+++ b/lib/CL/devices/basic/basic.c
@@ -221,10 +221,17 @@ pocl_basic_init_device_infos(struct _cl_device_id* dev)
   dev->vendor_id = 0;
   dev->max_compute_units = 0;
   dev->max_work_item_dimensions = 3;
-  dev->max_work_item_sizes[0] = CL_INT_MAX;
-  dev->max_work_item_sizes[1] = CL_INT_MAX;
-  dev->max_work_item_sizes[2] = CL_INT_MAX;
-  dev->max_work_group_size = 1024;
+  dev->max_work_item_sizes[0] = SIZE_MAX;
+  dev->max_work_item_sizes[1] = SIZE_MAX;
+  dev->max_work_item_sizes[2] = SIZE_MAX;
+  /*
+    The hard restriction will be the context data which is
+    stored in stack that can be as small as 8K in Linux.
+    Thus, there should be enough work-items alive to fill up
+    the SIMD lanes times the vector units, but not more than
+    that to avoid stack overflow and cache trashing.
+  */
+  dev->max_work_group_size = 1024*4;
   dev->preferred_wg_size_multiple = 8;
   dev->preferred_vector_width_char = POCL_DEVICES_PREFERRED_VECTOR_WIDTH_CHAR;
   dev->preferred_vector_width_short = POCL_DEVICES_PREFERRED_VECTOR_WIDTH_SHORT;
diff --git a/lib/CL/devices/pthread/pthread.c b/lib/CL/devices/pthread/pthread.c
index 18e6c9b..20b7234 100644
--- a/lib/CL/devices/pthread/pthread.c
+++ b/lib/CL/devices/pthread/pthread.c
@@ -199,11 +199,9 @@ pocl_pthread_init_device_infos(struct _cl_device_id* dev)
   pocl_basic_init_device_infos(dev);

   dev->type = CL_DEVICE_TYPE_CPU;
-  /* This could be SIZE_T_MAX, but setting it to INT_MAX should suffice, */
-  /* and may avoid errors in user code that uses int instead of size_t */
-  dev->max_work_item_sizes[0] = 1024;
-  dev->max_work_item_sizes[1] = 1024;
-  dev->max_work_item_sizes[2] = 1024;
+  dev->max_work_item_sizes[0] = SIZE_MAX;
+  dev->max_work_item_sizes[1] = SIZE_MAX;
+  dev->max_work_item_sizes[2] = SIZE_MAX;

 }

diff --git a/lib/CL/pocl_cl.h b/lib/CL/pocl_cl.h
index f120b7c..f7e05d3 100644
--- a/lib/CL/pocl_cl.h
+++ b/lib/CL/pocl_cl.h
@@ -111,6 +111,17 @@ extern int pocl_debug_messages;
   #ifdef HAVE_CLOCK_GETTIME

   extern struct timespec pocl_debug_timespec;
+  #define POCL_MSG_PRINT_INFO(...)                                            \
+    do {                                                                      \
+    if (pocl_debug_messages) {                                                \
+      clock_gettime(CLOCK_REALTIME, &pocl_debug_timespec);                    \
+      fprintf(stderr, "[%li.%li] POCL: in function %s"                  \
+      " at line %u:", (long)pocl_debug_timespec.tv_sec, (long)pocl_debug_timespec.tv_nsec, \
+        __func__, __LINE__);                                                  \
+      fprintf(stderr, __VA_ARGS__);                                           \
+    }                                                                         \
+  } while(0)
+
   #define POCL_MSG_PRINT(TYPE, ERRCODE, ...)                                  \
     do {                                                                      \
     if (pocl_debug_messages) {                                                \
diff --git a/tests/regression/test_barrier_before_return.cpp b/tests/regression/test_barrier_before_return.cpp
index edac92d..5edaeef 100644
--- a/tests/regression/test_barrier_before_return.cpp
+++ b/tests/regression/test_barrier_before_return.cpp
@@ -113,8 +113,7 @@ main(void)
             kernel,
             cl::NullRange,
             cl::NDRange(WORK_ITEMS),
-            cl::NullRange);
-
+            cl::NDRange(WORK_ITEMS));

         // Map cBuffer to host pointer. This enforces a sync with
         // the host backing space, remember we choose GPU device.
diff --git a/tests/regression/test_locals.cpp b/tests/regression/test_locals.cpp
index a1701a3..be41eb3 100644
--- a/tests/regression/test_locals.cpp
+++ b/tests/regression/test_locals.cpp
@@ -113,7 +113,7 @@ main(void)
             kernel,
             cl::NullRange,
             cl::NDRange(WORK_ITEMS),
-            cl::NullRange);
+            cl::NDRange(WORK_ITEMS));

         // Map aBuffer to host pointer. This enforces a sync with
         // the host backing space, remember we choose GPU device.
diff --git a/tests/tce/fp16/host.cpp b/tests/tce/fp16/host.cpp
index b0e6e1f..fd8196e 100644
--- a/tests/tce/fp16/host.cpp
+++ b/tests/tce/fp16/host.cpp
@@ -109,7 +109,7 @@ main(void)
             kernel,
             cl::NullRange,
             cl::NDRange(8),
-            cl::NullRange,
+            cl::NDRange(8),
             NULL, &enqEvent);
 	queue.finish();