commit f1e182bce33e42bee6e534ef35e763fd7d3bdda7
author Wind Yuan <feng.yuan@intel.com> Thu Jun 25 16:08:04 2015 +0800
committer Wind Yuan <feng.yuan@intel.com> Wed Jul 01 13:46:33 2015 +0800
tree 61830c57caeb76d367c80d8a6aeec1f41289df58
parent 09de60fc9332106d0d83bb6b5fef7d74f11801ec

cl: allow default local size if user doesn't set

 * developer can ignore local_size in prepare_arguments

xcore/cl_context.cpp

diff --git a/xcore/cl_context.cpp b/xcore/cl_context.cpp
index c15a05f..3f32911 100644
--- a/xcore/cl_context.cpp
+++ b/xcore/cl_context.cpp
@@ -228,6 +228,7 @@
     cl_event *event_out_id = NULL;
     cl_event events_id_wait[XCAM_CL_MAX_EVENT_SIZE];
     uint32_t num_of_events_wait = 0;
+    uint32_t work_group_size = 0;
 
     XCAM_ASSERT (kernel);
 
@@ -242,6 +243,12 @@
     if (event_out.ptr ())
         event_out_id = &event_out->get_event_id ();
 
+    for (uint32_t i = 0; i < work_dims; ++i) {
+        work_group_size *= local_sizes[i];
+    }
+    if (!work_group_size)
+        local_sizes = NULL;
+
     error_code =
         clEnqueueNDRangeKernel (
             cmd_queue_id, kernel_id,

xcore/cl_image_handler.cpp

diff --git a/xcore/cl_image_handler.cpp b/xcore/cl_image_handler.cpp
index aed885b..63dd0b6 100644
--- a/xcore/cl_image_handler.cpp
+++ b/xcore/cl_image_handler.cpp
@@ -20,6 +20,8 @@
 
 #include "cl_image_handler.h"
 #include "drm_display.h"
+#include "cl_device.h"
+
 
 namespace XCam {
 
@@ -29,8 +31,7 @@
     : dim (XCAM_DEFAULT_IMAGE_DIM)
 {
     xcam_mem_clear (global);
-    for (uint32_t i = 0; i < XCAM_CL_KERNEL_MAX_WORK_DIM; ++i)
-        local [i] = 1;
+    xcam_mem_clear (local);
 }
 
 CLArgument::CLArgument()
@@ -121,8 +122,8 @@
         work_size.global[0] = out_info.width;
         work_size.global[1] = out_info.height;
     }
-    work_size.local[0] = 4;
-    work_size.local[1] = 4;
+    work_size.local[0] = 0;
+    work_size.local[1] = 0;
 
     return XCAM_RETURN_NO_ERROR;
 }
@@ -331,6 +332,8 @@
             XCAM_STR (_name), kernel->get_kernel_name ());
     }
 
+    //CLDevice::instance()->get_context ()->finish ();
+
     XCAM_OBJ_PROFILING_END (XCAM_STR (_name), 30);
 
     return XCAM_RETURN_NO_ERROR;

xcore/cl_kernel.cpp

diff --git a/xcore/cl_kernel.cpp b/xcore/cl_kernel.cpp
index 3d20389..b0da967 100644
--- a/xcore/cl_kernel.cpp
+++ b/xcore/cl_kernel.cpp
@@ -23,7 +23,7 @@
 #include "cl_device.h"
 
 #define XCAM_CL_KERNEL_DEFAULT_WORK_DIM 2
-#define XCAM_CL_KERNEL_DEFAULT_LOCAL_WORK_SIZE 1
+#define XCAM_CL_KERNEL_DEFAULT_LOCAL_WORK_SIZE 0
 
 namespace XCam {
 
@@ -162,7 +162,7 @@
 
     XCAM_FAIL_RETURN (
         WARNING,
-        work_group_size <= dev_info.max_work_group_size,
+        work_group_size == 0 || work_group_size <= dev_info.max_work_group_size,
         XCAM_RETURN_ERROR_PARAM,
         "kernel(%s) work-group-size:%d is greater than device max work-group-size(%d)",
         _name, work_group_size, dev_info.max_work_group_size);