cl bayer: enable normalization and black level correction
* Shift 10bit bayer to 16bits to normalize image process
* Support black level correction on default value
Signed-off-by: Wind Yuan <feng.yuan@intel.com>
diff --git a/cl_kernel/kernel_blc.cl b/cl_kernel/kernel_blc.cl
new file mode 100644
index 0000000..8168d30
--- /dev/null
+++ b/cl_kernel/kernel_blc.cl
@@ -0,0 +1,53 @@
+/*
+ * function: kernel_blc
+ * black level correction for sensor data input
+ * input: image2d_t as read only
+ * output: image2d_t as write only
+ * blc_config: black level correction configuration
+ * param:
+ */
+
+"typedef struct "
+"{ "
+" float level_gr; /* Black level for GR pixels */ "
+" float level_r; /* Black level for R pixels */ "
+" float level_b; /* Black level for B pixels */ "
+" float level_gb; /* Black level for GB pixels */ "
+"}BLCConfig; "
+" "
+"__kernel void kernel_blc (__read_only image2d_t input, "
+" __write_only image2d_t output, "
+" BLCConfig blc_config) "
+"{ "
+" int x0 = 2 * get_global_id (0); "
+" int y0 = 2 * get_global_id (1); "
+" int x1 = x0 + 1; "
+" int y1 = y0 + 1; "
+" sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE | CLK_FILTER_NEAREST; "
+" "
+" int2 pos_r = (int2)(x0, y0); "
+" int2 pos_gr = (int2)(x1, y0); "
+" int2 pos_gb = (int2)(x0, y1); "
+" int2 pos_b = (int2)(x1, y1); "
+" "
+" uint4 pixel; "
+" pixel = read_imageui(input, sampler, pos_r); "
+" pixel.x = pixel.x << 6; "
+" pixel.x = pixel.x - blc_config.level_r * 65536; "
+" write_imageui(output, pos_r, pixel); "
+" "
+" pixel = read_imageui(input, sampler, pos_gr); "
+" pixel.x = pixel.x << 6; "
+" pixel.x = pixel.x - blc_config.level_gr * 65536; "
+" write_imageui(output, pos_gr, pixel); "
+" "
+" pixel = read_imageui(input, sampler, pos_gb); "
+" pixel.x = pixel.x << 6; "
+" pixel.x = pixel.x - blc_config.level_gb * 65536; "
+" write_imageui(output, pos_gb, pixel); "
+" "
+" pixel = read_imageui(input, sampler, pos_b); "
+" pixel.x = pixel.x << 6; "
+" pixel.x = pixel.x - blc_config.level_b * 65536; "
+" write_imageui(output, pos_b, pixel); "
+"} "
diff --git a/tests/test-cl-image.cpp b/tests/test-cl-image.cpp
index a5a471b..1479d41 100644
--- a/tests/test-cl-image.cpp
+++ b/tests/test-cl-image.cpp
@@ -23,6 +23,7 @@
#include "cl_context.h"
#include "cl_demo_handler.h"
#include "cl_hdr_handler.h"
+#include "cl_blc_handler.h"
#include "drm_bo_buffer.h"
#include "cl_demosaic_handler.h"
#include "cl_csc_handler.h"
@@ -202,6 +203,7 @@
image_handler = create_cl_demo_image_handler (context);
break;
case TestHandlerBlackLevel:
+ image_handler = create_cl_blc_image_handler (context);
break;
case TestHandlerDefect:
break;
diff --git a/xcore/Makefile.am b/xcore/Makefile.am
index d655d0d..8aa06f3 100644
--- a/xcore/Makefile.am
+++ b/xcore/Makefile.am
@@ -100,6 +100,7 @@
cl_image_handler.cpp \
cl_image_processor.cpp \
cl_demo_handler.cpp \
+ cl_blc_handler.cpp \
drm_bo_buffer.cpp \
cl_hdr_handler.cpp \
cl_demosaic_handler.cpp \
diff --git a/xcore/cl_blc_handler.cpp b/xcore/cl_blc_handler.cpp
new file mode 100644
index 0000000..b4e992a
--- /dev/null
+++ b/xcore/cl_blc_handler.cpp
@@ -0,0 +1,111 @@
+/*
+ * cl_blc_handler.cpp - CL black level correction handler
+ *
+ * Copyright (c) 2015 Intel Corporation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ *
+ * Author: Shincy Tu <shincy.tu@intel.com>
+ */
+#include "xcam_utils.h"
+#include "cl_blc_handler.h"
+
+namespace XCam {
+
+CLBlcImageKernel::CLBlcImageKernel (SmartPtr<CLContext> &context)
+ : CLImageKernel (context, "kernel_blc")
+{
+}
+
+XCamReturn
+CLBlcImageKernel::prepare_arguments (
+ SmartPtr<DrmBoBuffer> &input, SmartPtr<DrmBoBuffer> &output,
+ CLArgument args[], uint32_t &arg_count,
+ CLWorkSize &work_size)
+{
+ SmartPtr<CLContext> context = get_context ();
+ const VideoBufferInfo & video_info = input->get_video_info ();
+ cl_libva_image image_info;
+ uint32_t channel_bits = XCAM_ALIGN_UP (video_info.color_bits, 8);
+
+ xcam_mem_clear (&image_info);
+ image_info.fmt.image_channel_order = CL_R;
+ if (channel_bits == 8)
+ image_info.fmt.image_channel_data_type = CL_UNSIGNED_INT8;
+ else if (channel_bits == 16)
+ image_info.fmt.image_channel_data_type = CL_UNSIGNED_INT16;
+ image_info.offset = 0;
+ image_info.width = video_info.width;
+ image_info.height = (video_info.size / video_info.strides[0]) / 4 * 4;
+ image_info.row_pitch = video_info.strides[0];
+
+ _image_in = new CLVaImage (context, input, &image_info);
+ _image_out = new CLVaImage (context, output, &image_info);
+
+ XCAM_ASSERT (_image_in->is_valid () && _image_out->is_valid ());
+ XCAM_FAIL_RETURN (
+ WARNING,
+ _image_in->is_valid () && _image_out->is_valid (),
+ XCAM_RETURN_ERROR_MEM,
+ "cl image kernel(%s) in/out memory not available", get_kernel_name ());
+
+ _blc_config.level_b = (cl_float)XCAM_CL_BLACK_LEVEL / XCAM_CL_10BIT_NOR;
+ _blc_config.level_gr = (cl_float)XCAM_CL_BLACK_LEVEL / XCAM_CL_10BIT_NOR;
+ _blc_config.level_gb = (cl_float)XCAM_CL_BLACK_LEVEL / XCAM_CL_10BIT_NOR;
+ _blc_config.level_r = (cl_float)XCAM_CL_BLACK_LEVEL / XCAM_CL_10BIT_NOR;
+
+ //set args;
+ args[0].arg_adress = &_image_in->get_mem_id ();
+ args[0].arg_size = sizeof (cl_mem);
+ args[1].arg_adress = &_image_out->get_mem_id ();
+ args[1].arg_size = sizeof (cl_mem);
+ args[2].arg_adress = &_blc_config;
+ args[2].arg_size = sizeof (BLCConfig);
+ arg_count = 3;
+
+ work_size.dim = XCAM_DEFAULT_IMAGE_DIM;
+ work_size.global[0] = image_info.width / 2;
+ work_size.global[1] = image_info.height / 2;
+ work_size.local[0] = 8;
+ work_size.local[1] = 4;
+
+ return XCAM_RETURN_NO_ERROR;
+}
+
+SmartPtr<CLImageHandler>
+create_cl_blc_image_handler (SmartPtr<CLContext> &context)
+{
+ SmartPtr<CLImageHandler> blc_handler;
+ SmartPtr<CLImageKernel> blc_kernel;
+ XCamReturn ret = XCAM_RETURN_NO_ERROR;
+
+ blc_kernel = new CLBlcImageKernel (context);
+ {
+ XCAM_CL_KERNEL_FUNC_SOURCE_BEGIN(kernel_blc)
+#include "kernel_blc.cl"
+ XCAM_CL_KERNEL_FUNC_END;
+ ret = blc_kernel->load_from_source (kernel_blc_body, strlen (kernel_blc_body));
+ XCAM_FAIL_RETURN (
+ WARNING,
+ ret == XCAM_RETURN_NO_ERROR,
+ NULL,
+ "CL image handler(%s) load source failed", blc_kernel->get_kernel_name());
+ }
+ XCAM_ASSERT (blc_kernel->is_valid ());
+ blc_handler = new CLImageHandler ("cl_handler_blc");
+ blc_handler->add_kernel (blc_kernel);
+
+ return blc_handler;
+}
+
+}
diff --git a/xcore/cl_blc_handler.h b/xcore/cl_blc_handler.h
new file mode 100644
index 0000000..afbbcf3
--- /dev/null
+++ b/xcore/cl_blc_handler.h
@@ -0,0 +1,67 @@
+/*
+ * cl_blc_handler.h - CL black level correction handler
+ *
+ * Copyright (c) 2015 Intel Corporation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ *
+ * Author: Shincy Tu <shincy.tu@intel.com>
+ */
+
+#ifndef XCAM_CL_BLC_HANLDER_H
+#define XCAM_CL_BLC_HANLDER_H
+
+#include "xcam_utils.h"
+#include "cl_image_handler.h"
+
+#define XCAM_CL_BLACK_LEVEL 0x3c
+#define XCAM_CL_10BIT_NOR 0x400 /* Normalization for 10bit data */
+
+namespace XCam {
+
+class CLBlcImageKernel
+ : public CLImageKernel
+{
+
+public:
+ /* Black level correction configuration
+ *
+ */
+ typedef struct
+ {
+ cl_float level_gr; /* Black level for GR pixels */
+ cl_float level_r; /* Black level for R pixels */
+ cl_float level_b; /* Black level for B pixels */
+ cl_float level_gb; /* Black level for GB pixels */
+ } BLCConfig;
+
+public:
+ explicit CLBlcImageKernel (SmartPtr<CLContext> &context);
+
+protected:
+ virtual XCamReturn prepare_arguments (
+ SmartPtr<DrmBoBuffer> &input, SmartPtr<DrmBoBuffer> &output,
+ CLArgument args[], uint32_t &arg_count,
+ CLWorkSize &work_size);
+ BLCConfig _blc_config;
+
+private:
+ XCAM_DEAD_COPY (CLBlcImageKernel);
+};
+
+SmartPtr<CLImageHandler>
+create_cl_blc_image_handler (SmartPtr<CLContext> &context);
+
+};
+
+#endif //XCAM_CL_BLC_HANLDER_H
diff --git a/xcore/cl_image_processor.cpp b/xcore/cl_image_processor.cpp
index bc37387..eac91f8 100644
--- a/xcore/cl_image_processor.cpp
+++ b/xcore/cl_image_processor.cpp
@@ -23,6 +23,8 @@
#include "cl_image_handler.h"
#include "drm_display.h"
#include "cl_demo_handler.h"
+#include "cl_blc_handler.h"
+
namespace XCam {
@@ -115,6 +117,15 @@
"CLImageProcessor create demo handler failed");
add_handler (demo_handler);
+ SmartPtr<CLImageHandler> blc_handler;
+ blc_handler = create_cl_blc_image_handler (_context);
+ XCAM_FAIL_RETURN (
+ WARNING,
+ blc_handler.ptr (),
+ XCAM_RETURN_ERROR_CL,
+ "CLImageProcessor create blc handler failed");
+ add_handler (blc_handler);
+
return XCAM_RETURN_NO_ERROR;
}