cl: add calculator for 3a statistics
* need calculate 3a stats before whitebalance
* 3a stats buffer read out from cl buffer and copy to
x3astats pool
diff --git a/cl_kernel/kernel_3a_stats.cl b/cl_kernel/kernel_3a_stats.cl
new file mode 100644
index 0000000..4c5cf81
--- /dev/null
+++ b/cl_kernel/kernel_3a_stats.cl
@@ -0,0 +1,154 @@
+/*
+ * function: kernel_3a_stats
+ * input: image2d_t as read only
+ * output: XCamGridStat, stats results
+ */
+
+"typedef struct"
+"{"
+" unsigned int avg_y;"
+
+" unsigned int avg_r;"
+" unsigned int avg_gr;"
+" unsigned int avg_gb;"
+" unsigned int avg_b;"
+" unsigned int valid_wb_count;"
+
+" unsigned int f_value1;"
+" unsigned int f_value2;"
+"} XCamGridStat;\n"
+
+"__kernel void kernel_3a_stats (__read_only image2d_t input, __global XCamGridStat *output) "
+"{ "
+" int x = get_global_id (0); "
+" int y = get_global_id (1); "
+" int w = get_global_size (0); "
+" sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE | CLK_FILTER_NEAREST; "
+" "
+" int x0 = 16 * x; "
+" int y0 = 16 * y; "
+" float sum_gr = 0.0f, sum_r = 0.0f, sum_b = 0.0f, sum_gb=0.0f; "
+" float avg_gr = 0.0f, avg_r = 0.0f, avg_b = 0.0f, avg_gb = 0.0f; "
+" int i = 0, j = 0; "
+" float count = (16.0 / 2) * (16.0 / 2); "
+" float4 p[4]; "
+
+"\n#pragma unroll\n"
+" for (j = 0; j < 16; j += 2) { "
+
+// grid (0, 0)
+" i = 0; "
+" p[0] = read_imagef (input, sampler, (int2)(x0 + i, y0 + j)); "
+" p[1] = read_imagef (input, sampler, (int2)(x0 + i, y0 + j + 1)); "
+" ++i; "
+" p[2] = read_imagef (input, sampler, (int2)(x0 + i, y0 + j)); "
+" p[3] = read_imagef (input, sampler, (int2)(x0 + i, y0 + j + 1)); "
+" sum_gr += p[0].x; "
+" sum_b += p[1].x; "
+" sum_r += p[2].x; "
+" sum_gb += p[3].x; "
+
+// grid (1, 0)
+" ++i; "
+" p[0] = read_imagef (input, sampler, (int2)(x0 + i, y0 + j)); "
+" p[1] = read_imagef (input, sampler, (int2)(x0 + i, y0 + j + 1)); "
+" ++i; "
+" p[2] = read_imagef (input, sampler, (int2)(x0 + i, y0 + j)); "
+" p[3] = read_imagef (input, sampler, (int2)(x0 + i, y0 + j + 1)); "
+" sum_gr += p[0].x; "
+" sum_b += p[1].x; "
+" sum_r += p[2].x; "
+" sum_gb += p[3].x; "
+
+// grid (2, 0)
+" ++i; "
+" p[0] = read_imagef (input, sampler, (int2)(x0 + i, y0 + j)); "
+" p[1] = read_imagef (input, sampler, (int2)(x0 + i, y0 + j + 1)); "
+" ++i; "
+" p[2] = read_imagef (input, sampler, (int2)(x0 + i, y0 + j)); "
+" p[3] = read_imagef (input, sampler, (int2)(x0 + i, y0 + j + 1)); "
+" sum_gr += p[0].x; "
+" sum_b += p[1].x; "
+" sum_r += p[2].x; "
+" sum_gb += p[3].x; "
+
+// grid (3, 0)
+" ++i; "
+" p[0] = read_imagef (input, sampler, (int2)(x0 + i, y0 + j)); "
+" p[1] = read_imagef (input, sampler, (int2)(x0 + i, y0 + j + 1)); "
+" ++i; "
+" p[2] = read_imagef (input, sampler, (int2)(x0 + i, y0 + j)); "
+" p[3] = read_imagef (input, sampler, (int2)(x0 + i, y0 + j + 1)); "
+" sum_gr += p[0].x; "
+" sum_b += p[1].x; "
+" sum_r += p[2].x; "
+" sum_gb += p[3].x; "
+
+// grid (4, 0)
+" ++i; "
+" p[0] = read_imagef (input, sampler, (int2)(x0 + i, y0 + j)); "
+" p[1] = read_imagef (input, sampler, (int2)(x0 + i, y0 + j + 1)); "
+" ++i; "
+" p[2] = read_imagef (input, sampler, (int2)(x0 + i, y0 + j)); "
+" p[3] = read_imagef (input, sampler, (int2)(x0 + i, y0 + j + 1)); "
+" sum_gr += p[0].x; "
+" sum_b += p[1].x; "
+" sum_r += p[2].x; "
+" sum_gb += p[3].x; "
+
+// grid (5, 0)
+" ++i; "
+" p[0] = read_imagef (input, sampler, (int2)(x0 + i, y0 + j)); "
+" p[1] = read_imagef (input, sampler, (int2)(x0 + i, y0 + j + 1)); "
+" ++i; "
+" p[2] = read_imagef (input, sampler, (int2)(x0 + i, y0 + j)); "
+" p[3] = read_imagef (input, sampler, (int2)(x0 + i, y0 + j + 1)); "
+" sum_gr += p[0].x; "
+" sum_b += p[1].x; "
+" sum_r += p[2].x; "
+" sum_gb += p[3].x; "
+
+// grid (6, 0)
+" ++i; "
+" p[0] = read_imagef (input, sampler, (int2)(x0 + i, y0 + j)); "
+" p[1] = read_imagef (input, sampler, (int2)(x0 + i, y0 + j + 1)); "
+" ++i; "
+" p[2] = read_imagef (input, sampler, (int2)(x0 + i, y0 + j)); "
+" p[3] = read_imagef (input, sampler, (int2)(x0 + i, y0 + j + 1)); "
+" sum_gr += p[0].x; "
+" sum_b += p[1].x; "
+" sum_r += p[2].x; "
+" sum_gb += p[3].x; "
+
+
+// grid (7, 0)
+" ++i; "
+" p[0] = read_imagef (input, sampler, (int2)(x0 + i, y0 + j)); "
+" p[1] = read_imagef (input, sampler, (int2)(x0 + i, y0 + j + 1)); "
+" ++i; "
+" p[2] = read_imagef (input, sampler, (int2)(x0 + i, y0 + j)); "
+" p[3] = read_imagef (input, sampler, (int2)(x0 + i, y0 + j + 1)); "
+" sum_gr += p[0].x; "
+" sum_b += p[1].x; "
+" sum_r += p[2].x; "
+" sum_gb += p[3].x; "
+
+//end for loop
+" } "
+
+" avg_gr = sum_gr/count; "
+" avg_r = sum_r/count; "
+" avg_b = sum_b/count; "
+" avg_gb = sum_gb/count; "
+
+" output[y * w + x].avg_gr = convert_uint(avg_gr * 256.0); "
+" output[y * w + x].avg_r = convert_uint(avg_r * 256.0); "
+" output[y * w + x].avg_b = convert_uint(avg_b * 256.0); "
+" output[y * w + x].avg_gb = convert_uint(avg_gb * 256.0); "
+" output[y * w + x].valid_wb_count = 255; "
+" output[y * w + x].avg_y = convert_uint(((avg_gr + avg_gb)/2.0f)*256.0); "
+" output[y * w + x].f_value1 = 0; "
+" output[y * w + x].f_value2 = 0; "
+
+"} "
+
diff --git a/xcore/Makefile.am b/xcore/Makefile.am
index 1807ffb..0108cb0 100644
--- a/xcore/Makefile.am
+++ b/xcore/Makefile.am
@@ -66,6 +66,7 @@
x3a_analyzer_manager.cpp \
x3a_analyzer_simple.cpp \
x3a_image_process_center.cpp \
+ x3a_stats_pool.cpp \
x3a_isp_config.cpp \
x3a_result.cpp \
x3a_result_factory.cpp \
@@ -102,15 +103,16 @@
cl_image_bo_buffer.cpp \
cl_image_handler.cpp \
cl_image_processor.cpp \
- cl_3a_image_processor.cpp \
+ cl_3a_image_processor.cpp \
+ cl_3a_stats_calculator.cpp \
cl_demo_handler.cpp \
cl_blc_handler.cpp \
drm_bo_buffer.cpp \
cl_hdr_handler.cpp \
cl_demosaic_handler.cpp \
cl_csc_handler.cpp \
- cl_wb_handler.cpp \
- cl_denoise_handler.cpp \
+ cl_wb_handler.cpp \
+ cl_denoise_handler.cpp \
cl_gamma_handler.cpp \
$(NULL)
endif
diff --git a/xcore/cl_3a_stats_calculator.cpp b/xcore/cl_3a_stats_calculator.cpp
new file mode 100644
index 0000000..79038c8
--- /dev/null
+++ b/xcore/cl_3a_stats_calculator.cpp
@@ -0,0 +1,175 @@
+/*
+ * cl_3a_stats_calculator.cpp - CL 3a calculator
+ *
+ * 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: Wind Yuan <feng.yuan@intel.com>
+ */
+
+#include "xcam_utils.h"
+#include "cl_3a_stats_calculator.h"
+
+namespace XCam {
+
+CL3AStatsCalculatorKernel::CL3AStatsCalculatorKernel (SmartPtr<CLContext> &context)
+ : CLImageKernel (context, "kernel_3a_stats")
+ , _data_allocated (false)
+{
+ xcam_mem_clear (&_stats_info);
+}
+
+XCamReturn
+CL3AStatsCalculatorKernel::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 ();
+
+ XCAM_UNUSED (output);
+
+ if (!_data_allocated && !allocate_data (video_info)) {
+ XCAM_LOG_WARNING ("CL3AStatsCalculatorKernel allocate data failed");
+ return XCAM_RETURN_ERROR_MEM;
+ }
+
+ _image_in = new CLVaImage (context, input);
+
+ //set args;
+ args[0].arg_adress = &_image_in->get_mem_id ();
+ args[0].arg_size = sizeof (cl_mem);
+ args[1].arg_adress = &_stats_cl_buffer->get_mem_id ();
+ args[1].arg_size = sizeof (cl_mem);
+ arg_count = 2;
+
+ work_size.dim = XCAM_DEFAULT_IMAGE_DIM;
+ // grid_size default 16
+ work_size.global[0] = _stats_info.aligned_width;
+ work_size.global[1] = _stats_info.aligned_height;
+ work_size.local[0] = 8;
+ work_size.local[1] = 1;
+
+ return XCAM_RETURN_NO_ERROR;
+}
+
+XCamReturn
+CL3AStatsCalculatorKernel::post_execute ()
+{
+ SmartPtr<CLContext> context = get_context ();
+ SmartPtr<BufferProxy> buffer;
+ SmartPtr<X3aStats> stats;
+ SmartPtr<CLEvent> event = new CLEvent;
+ XCam3AStats *stats_ptr = NULL;
+ XCamReturn ret = XCAM_RETURN_NO_ERROR;
+
+ context->finish ();
+ _image_in.release ();
+ //copy out and post 3a stats
+ buffer = _stats_pool->get_buffer (_stats_pool);
+ stats = buffer.dynamic_cast_ptr<X3aStats> ();
+ XCAM_ASSERT (stats.ptr ());
+ stats_ptr = stats->get_stats ();
+ ret = _stats_cl_buffer->enqueue_read (
+ stats_ptr->stats,
+ 0, _stats_info.aligned_width * _stats_info.aligned_height * sizeof (stats_ptr->stats[0]),
+ CLEvent::EmptyList, event);
+
+ XCAM_FAIL_RETURN (WARNING, ret == XCAM_RETURN_NO_ERROR, ret, "3a stats enqueue read buffer failed.");
+ XCAM_ASSERT (event->get_event_id ());
+
+ ret = event->wait ();
+ XCAM_FAIL_RETURN (WARNING, ret == XCAM_RETURN_NO_ERROR, ret, "3a stats buffer enqueue event wait failed");
+ event.release ();
+
+ //post stats out
+ return post_stats (stats);
+}
+
+XCamReturn
+CL3AStatsCalculatorKernel::post_stats (const SmartPtr<X3aStats> &stats)
+{
+ return XCAM_RETURN_NO_ERROR;
+}
+
+bool
+CL3AStatsCalculatorKernel::allocate_data (const VideoBufferInfo &buffer_info)
+{
+ SmartPtr<CLContext> context = get_context ();
+
+ _stats_pool = new X3aStatsPool ();
+ _stats_pool->set_video_info (buffer_info);
+
+ XCAM_FAIL_RETURN (
+ WARNING,
+ _stats_pool->reserve (6),
+ false,
+ "reserve cl stats buffer failed");
+
+ _stats_info = _stats_pool->get_stats_info ();
+ _stats_cl_buffer = new CLBuffer (
+ context,
+ _stats_info.aligned_width * _stats_info.aligned_height * sizeof (XCamGridStat));
+
+ XCAM_FAIL_RETURN (
+ WARNING,
+ _stats_cl_buffer->is_valid (),
+ false,
+ "allocate cl stats buffer failed");
+ _data_allocated = true;
+
+ return true;
+}
+
+CL3AStatsCalculator::CL3AStatsCalculator ()
+ : CLImageHandler ("CL3AStatsCalculator")
+{
+}
+
+XCamReturn
+CL3AStatsCalculator::prepare_output_buf (SmartPtr<DrmBoBuffer> &input, SmartPtr<DrmBoBuffer> &output)
+{
+ output = input;
+ return XCAM_RETURN_NO_ERROR;
+}
+
+SmartPtr<CLImageHandler>
+create_cl_3a_stats_image_handler (SmartPtr<CLContext> &context)
+{
+ SmartPtr<CLImageHandler> x3a_stats_handler;
+ SmartPtr<CLImageKernel> x3a_stats_kernel;
+ XCamReturn ret = XCAM_RETURN_NO_ERROR;
+
+ x3a_stats_kernel = new CL3AStatsCalculatorKernel (context);
+ {
+ XCAM_CL_KERNEL_FUNC_SOURCE_BEGIN(kernel_3a_stats)
+#include "kernel_3a_stats.cl"
+ XCAM_CL_KERNEL_FUNC_END;
+ ret = x3a_stats_kernel->load_from_source (kernel_3a_stats_body, strlen (kernel_3a_stats_body));
+ XCAM_FAIL_RETURN (
+ WARNING,
+ ret == XCAM_RETURN_NO_ERROR,
+ NULL,
+ "CL image handler(%s) load source failed", x3a_stats_kernel->get_kernel_name());
+ }
+ XCAM_ASSERT (x3a_stats_kernel->is_valid ());
+ x3a_stats_handler = new CL3AStatsCalculator ();
+ x3a_stats_handler->add_kernel (x3a_stats_kernel);
+
+ return x3a_stats_handler;
+}
+
+
+};
diff --git a/xcore/cl_3a_stats_calculator.h b/xcore/cl_3a_stats_calculator.h
new file mode 100644
index 0000000..b6bd4d4
--- /dev/null
+++ b/xcore/cl_3a_stats_calculator.h
@@ -0,0 +1,76 @@
+/*
+ * cl_3a_stats_calculator.h - CL 3a calculator
+ *
+ * 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: Wind Yuan <feng.yuan@intel.com>
+ */
+
+#ifndef XCAM_CL_3A_CALCULATOR_H
+#define XCAM_CL_3A_CALCULATOR_H
+
+#include "xcam_utils.h"
+#include "cl_image_handler.h"
+#include "cl_memory.h"
+#include "x3a_stats_pool.h"
+
+namespace XCam {
+
+class CL3AStatsCalculatorKernel
+ : public CLImageKernel
+{
+public:
+ explicit CL3AStatsCalculatorKernel (SmartPtr<CLContext> &context);
+
+public:
+ virtual XCamReturn prepare_arguments (
+ SmartPtr<DrmBoBuffer> &input, SmartPtr<DrmBoBuffer> &output,
+ CLArgument args[], uint32_t &arg_count,
+ CLWorkSize &work_size);
+
+ virtual XCamReturn post_execute ();
+
+private:
+ bool allocate_data (const VideoBufferInfo &buffer_info);
+ XCamReturn post_stats (const SmartPtr<X3aStats> &stats);
+
+ XCAM_DEAD_COPY (CL3AStatsCalculatorKernel);
+
+private:
+ SmartPtr<X3aStatsPool> _stats_pool;
+ SmartPtr<CLBuffer> _stats_cl_buffer;
+ XCam3AStatsInfo _stats_info;
+ bool _data_allocated;
+};
+
+class CL3AStatsCalculator
+ : public CLImageHandler
+{
+public:
+ explicit CL3AStatsCalculator ();
+
+public:
+ virtual XCamReturn prepare_output_buf (SmartPtr<DrmBoBuffer> &input, SmartPtr<DrmBoBuffer> &output);
+
+private:
+ XCAM_DEAD_COPY (CL3AStatsCalculator);
+};
+
+SmartPtr<CLImageHandler>
+create_cl_3a_stats_image_handler (SmartPtr<CLContext> &context);
+
+};
+
+#endif // XCAM_CL_3A_CALCULATOR_H
diff --git a/xcore/cl_image_handler.h b/xcore/cl_image_handler.h
index 00ea5cf..e6088ed 100644
--- a/xcore/cl_image_handler.h
+++ b/xcore/cl_image_handler.h
@@ -87,7 +87,9 @@
virtual XCamReturn prepare_buffer_pool_video_info (
const VideoBufferInfo &input,
VideoBufferInfo &output);
- XCamReturn prepare_output_buf (SmartPtr<DrmBoBuffer> &input, SmartPtr<DrmBoBuffer> &output);
+
+ // if derive prepare_output_buf, then prepare_buffer_pool_video_info is not involked
+ virtual XCamReturn prepare_output_buf (SmartPtr<DrmBoBuffer> &input, SmartPtr<DrmBoBuffer> &output);
XCamReturn create_buffer_pool (const VideoBufferInfo &video_info);
SmartPtr<BufferPool> &get_buffer_pool () {
return _buf_pool;
diff --git a/xcore/x3a_stats_pool.cpp b/xcore/x3a_stats_pool.cpp
new file mode 100644
index 0000000..04c9b97
--- /dev/null
+++ b/xcore/x3a_stats_pool.cpp
@@ -0,0 +1,115 @@
+/*
+ * x3a_stats_pool.cpp - 3a stats pool
+ *
+ * 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: Wind Yuan <feng.yuan@intel.com>
+ */
+
+#include "x3a_stats_pool.h"
+
+namespace XCam {
+
+X3aStatsData::X3aStatsData (XCam3AStats *data)
+ : _data (data)
+{
+ XCAM_ASSERT (_data);
+}
+
+X3aStatsData::~X3aStatsData ()
+{
+ if (_data)
+ xcam_free (_data);
+}
+
+uint8_t *
+X3aStatsData::map ()
+{
+ return (uint8_t*)(intptr_t)(_data);
+}
+
+bool
+X3aStatsData::unmap ()
+{
+ return true;
+}
+
+X3aStats::X3aStats (const SmartPtr<X3aStatsData> &data)
+ : BufferProxy (SmartPtr<BufferData>(data))
+{
+}
+
+
+XCam3AStats *
+X3aStats::get_stats ()
+{
+ SmartPtr<BufferData> data = get_buffer_data ();
+ SmartPtr<X3aStatsData> stats = data.dynamic_cast_ptr<X3aStatsData> ();
+
+ XCAM_FAIL_RETURN(
+ WARNING,
+ stats.ptr(),
+ NULL,
+ "X3aStats get_stats failed with NULL");
+ return stats->get_stats ();
+}
+
+X3aStatsPool::X3aStatsPool ()
+{
+}
+
+void
+X3aStatsPool::set_stats_info (const XCam3AStatsInfo &info)
+{
+ _stats_info = info;
+}
+
+bool
+X3aStatsPool::fixate_video_info (VideoBufferInfo &info)
+{
+ const uint32_t grid = 16;
+
+ _stats_info.aligned_width = (info.width + grid - 1) / grid;
+ _stats_info.aligned_height = (info.height + grid - 1) / grid;
+
+ _stats_info.width = info.width / grid;
+ _stats_info.height = info.height / grid;
+ _stats_info.grid_pixel_size = grid;
+ _stats_info.bit_depth = 8;
+ return true;
+}
+
+SmartPtr<BufferData>
+X3aStatsPool::allocate_data (const VideoBufferInfo &buffer_info)
+{
+ XCam3AStats *stats = NULL;
+ stats =
+ (XCam3AStats *) xcam_malloc0 (
+ sizeof (XCam3AStats) +
+ sizeof (XCamGridStat) * _stats_info.aligned_width * _stats_info.aligned_height);
+ stats->info = _stats_info;
+ return new X3aStatsData (stats);
+}
+
+SmartPtr<BufferProxy>
+X3aStatsPool::create_buffer_from_data (SmartPtr<BufferData> &data)
+{
+ SmartPtr<X3aStatsData> stats_data = data.dynamic_cast_ptr<X3aStatsData> ();
+ XCAM_ASSERT (stats_data.ptr ());
+
+ return new X3aStats (stats_data);
+}
+
+};
diff --git a/xcore/x3a_stats_pool.h b/xcore/x3a_stats_pool.h
new file mode 100644
index 0000000..d979446
--- /dev/null
+++ b/xcore/x3a_stats_pool.h
@@ -0,0 +1,87 @@
+/*
+ * x3a_stats_pool.h - 3a stats pool
+ *
+ * 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: Wind Yuan <feng.yuan@intel.com>
+ */
+
+#ifndef XCAM_3A_STATS_POOL_H
+#define XCAM_3A_STATS_POOL_H
+
+#include "xcam_utils.h"
+#include "buffer_pool.h"
+#include <base/xcam_3a_stats.h>
+
+namespace XCam {
+
+class X3aStatsData
+ : public BufferData
+{
+public:
+ explicit X3aStatsData (XCam3AStats *data);
+ ~X3aStatsData ();
+ XCam3AStats *get_stats () {
+ return _data;
+ }
+
+ virtual uint8_t *map ();
+ virtual bool unmap ();
+
+private:
+ XCAM_DEAD_COPY (X3aStatsData);
+private:
+ XCam3AStats *_data;
+};
+
+class X3aStats
+ : public BufferProxy
+{
+ friend class X3aStatsPool;
+public:
+ XCam3AStats *get_stats ();
+
+protected:
+ explicit X3aStats (const SmartPtr<X3aStatsData> &data);
+ XCAM_DEAD_COPY (X3aStats);
+
+};
+
+class X3aStatsPool
+ : public BufferPool
+{
+public:
+ explicit X3aStatsPool ();
+ XCam3AStatsInfo &get_stats_info () {
+ return _stats_info;
+ }
+ void set_stats_info (const XCam3AStatsInfo &info);
+
+protected:
+ virtual bool fixate_video_info (VideoBufferInfo &info);
+ virtual SmartPtr<BufferData> allocate_data (const VideoBufferInfo &buffer_info);
+ virtual SmartPtr<BufferProxy> create_buffer_from_data (SmartPtr<BufferData> &data);
+
+private:
+ XCAM_DEAD_COPY (X3aStatsPool);
+
+private:
+ XCam3AStatsInfo _stats_info;
+};
+
+};
+
+#endif //XCAM_3A_STATS_POOL_H
+