XCamReturn
X3aAnalyzerAiq::internal_init (uint32_t width, uint32_t height, double framerate)
{
XCAM_ASSERT (_cpf_path);
CpfReader reader (_cpf_path);
ia_binary_data binary;
XCAM_ASSERT (_aiq_compositor.ptr ());
_aiq_compositor->set_framerate (framerate);
xcam_mem_clear (binary);
XCAM_FAIL_RETURN (
ERROR,
reader.read(binary),
XCAM_RETURN_ERROR_AIQ,
"read cpf file(%s) failed", _cpf_path);
_aiq_compositor->set_size (width, height);
XCAM_FAIL_RETURN (
ERROR,
_aiq_compositor->open (binary),
XCAM_RETURN_ERROR_AIQ,
"AIQ open failed");
return XCAM_RETURN_NO_ERROR;
}
SmartPtr<CLImageHandler>
create_cl_bayer_basic_image_handler (const SmartPtr<CLContext> &context, bool enable_gamma, uint32_t stats_bits)
{
SmartPtr<CLBayerBasicImageHandler> bayer_planar_handler;
SmartPtr<CLBayerBasicImageKernel> basic_kernel;
char build_options[1024];
bayer_planar_handler = new CLBayerBasicImageHandler (context, "cl_handler_bayer_basic");
bayer_planar_handler->set_stats_bits (stats_bits);
basic_kernel = new CLBayerBasicImageKernel (context);
XCAM_ASSERT (basic_kernel.ptr ());
xcam_mem_clear (build_options);
snprintf (build_options, sizeof (build_options),
" -DENABLE_GAMMA=%d "
" -DENABLE_IMAGE_2D_INPUT=%d "
" -DSTATS_BITS=%d ",
(enable_gamma ? 1 : 0),
ENABLE_IMAGE_2D_INPUT,
stats_bits);
XCAM_FAIL_RETURN (
ERROR, basic_kernel->build_kernel (kernel_bayer_basic_info, build_options) == XCAM_RETURN_NO_ERROR, NULL,
"build bayer-basic kernel(%s) failed", kernel_bayer_basic_info.kernel_name);
XCAM_ASSERT (basic_kernel->is_valid ());
bayer_planar_handler->set_bayer_kernel (basic_kernel);
return bayer_planar_handler;
}
CLBayerBasicImageHandler::CLBayerBasicImageHandler (
const SmartPtr<CLContext> &context, const char *name)
: CLImageHandler (context, name)
, _is_first_buf (true)
{
_blc_config.level_gr = XCAM_CL_BLC_DEFAULT_LEVEL;
_blc_config.level_r = XCAM_CL_BLC_DEFAULT_LEVEL;
_blc_config.level_b = XCAM_CL_BLC_DEFAULT_LEVEL;
_blc_config.level_gb = XCAM_CL_BLC_DEFAULT_LEVEL;
_blc_config.color_bits = 10;
_wb_config.r_gain = 1.0;
_wb_config.gr_gain = 1.0;
_wb_config.gb_gain = 1.0;
_wb_config.b_gain = 1.0;
for(int i = 0; i < XCAM_GAMMA_TABLE_SIZE; i++)
_gamma_table[i] = (float)i / 256.0f;
_gamma_table[XCAM_GAMMA_TABLE_SIZE] = 0.9999f;
SmartPtr<CL3AStatsCalculatorContext> stats_context = new CL3AStatsCalculatorContext (context);
XCAM_ASSERT (stats_context.ptr ());
_3a_stats_context = stats_context;
SmartPtr<CLBayer3AStatsThread> stats_thread = new CLBayer3AStatsThread (this);
XCAM_ASSERT (stats_thread.ptr ());
_3a_stats_thread = stats_thread;
XCAM_OBJ_PROFILING_INIT;
}
bool
CLBayer3AStatsThread::loop ()
{
XCamReturn ret = XCAM_RETURN_NO_ERROR;
SmartPtr<BayerPostData> data;
data = _stats_process_list.pop ();
if (!data.ptr ()) {
XCAM_LOG_INFO ("cl bayer 3a-stats thread is going to stop, processing data empty");
return false;
}
XCAM_ASSERT (data->image_buffer.ptr ());
XCAM_ASSERT (data->stats_cl_buf.ptr ());
XCAM_ASSERT (_handler);
ret = _handler->process_stats_buffer (data->image_buffer, data->stats_cl_buf);
XCAM_FAIL_RETURN (
WARNING,
ret == XCAM_RETURN_NO_ERROR,
false,
"cl bayer 3a-stats thread has error buffer on kernel post processing");
XCAM_FAIL_RETURN (
ERROR,
_buffer_done_list.push (data->image_buffer),
false,
"cl bayer 3a-stats thread failed to queue done-buffers");
return true;
}
XCamReturn
CLKernel::load_from_binary (const uint8_t *binary, size_t length)
{
cl_kernel new_kernel_id = NULL;
XCAM_ASSERT (binary);
if (!binary || !length) {
XCAM_LOG_WARNING ("kernel:%s binary empty", XCAM_STR (_name));
return XCAM_RETURN_ERROR_PARAM;
}
if (_kernel_id) {
XCAM_LOG_WARNING ("kernel:%s already build yet", XCAM_STR (_name));
return XCAM_RETURN_ERROR_PARAM;
}
XCAM_ASSERT (_context.ptr ());
new_kernel_id =
_context->generate_kernel_id (
this,
binary, length,
CLContext::KERNEL_BUILD_BINARY);
XCAM_FAIL_RETURN(
WARNING,
new_kernel_id != NULL,
XCAM_RETURN_ERROR_CL,
"cl kernel(%s) load from binary failed", XCAM_STR (_name));
_kernel_id = new_kernel_id;
return XCAM_RETURN_NO_ERROR;
}
SmartPtr<CLImageHandler>
create_cl_wavelet_denoise_image_handler (SmartPtr<CLContext> &context)
{
SmartPtr<CLWaveletDenoiseImageHandler> wavelet_handler;
SmartPtr<CLWaveletDenoiseImageKernel> wavelet_kernel;
XCamReturn ret = XCAM_RETURN_NO_ERROR;
XCAM_CL_KERNEL_FUNC_SOURCE_BEGIN(kernel_wavelet_denoise)
#include "kernel_wavelet_denoise.clx"
XCAM_CL_KERNEL_FUNC_END;
wavelet_handler = new CLWaveletDenoiseImageHandler ("cl_handler_wavelet_denoise");
XCAM_ASSERT (wavelet_handler.ptr ());
for (int layer = 1; layer <= WAVELET_DECOMPOSITION_LEVELS; layer++) {
wavelet_kernel = new CLWaveletDenoiseImageKernel (context, "kernel_wavelet_denoise", wavelet_handler, layer);
ret = wavelet_kernel->load_from_source (
kernel_wavelet_denoise_body, strlen (kernel_wavelet_denoise_body),
NULL, NULL,
WAVELET_DENOISE_UV ? "-DWAVELET_DENOISE_UV=1" : "-DWAVELET_DENOISE_UV=0");
XCAM_FAIL_RETURN (
WARNING,
ret == XCAM_RETURN_NO_ERROR,
NULL,
"CL image handler(%s) load source failed", wavelet_kernel->get_kernel_name());
XCAM_ASSERT (wavelet_kernel->is_valid ());
SmartPtr<CLImageKernel> image_kernel = wavelet_kernel;
wavelet_handler->add_kernel (image_kernel);
}
return wavelet_handler;
}
/*
* Default kernel arguments
* arg0:
* input, __read_only image2d_t
* arg1:
* output, __write_only image2d_t
* suppose cl can get width/height pixels from
* get_image_width/get_image_height
*/
XCamReturn
CLImageKernel::pre_execute (SmartPtr<DrmBoBuffer> &input, SmartPtr<DrmBoBuffer> &output)
{
XCamReturn ret = XCAM_RETURN_NO_ERROR;
SmartPtr<CLContext> context = get_context ();
#define XCAM_CL_MAX_ARGS 256
CLArgument args[XCAM_CL_MAX_ARGS];
uint32_t arg_count = XCAM_CL_MAX_ARGS;
CLWorkSize work_size;
ret = prepare_arguments (input, output, args, arg_count, work_size);
XCAM_ASSERT (arg_count);
for (uint32_t i = 0; i < arg_count; ++i) {
ret = set_argument (i, args[i].arg_adress, args[i].arg_size);
XCAM_FAIL_RETURN (
WARNING,
ret == XCAM_RETURN_NO_ERROR,
ret,
"cl image kernel(%s) set argc(%d) failed", get_kernel_name (), i);
}
XCAM_ASSERT (work_size.global[0]);
ret = set_work_size (work_size.dim, work_size.global, work_size.local);
XCAM_FAIL_RETURN (
WARNING,
ret == XCAM_RETURN_NO_ERROR,
ret,
"cl image kernel(%s) set work size failed", get_kernel_name ());
return XCAM_RETURN_NO_ERROR;
}
void
VKDevice::destroy_cmd_pool (VkCommandPool pool)
{
XCAM_ASSERT (XCAM_IS_VALID_VK_ID (_dev_id));
XCAM_ASSERT (XCAM_IS_VALID_VK_ID (pool));
vkDestroyCommandPool (_dev_id, pool, _allocator.ptr ());
}
XCamReturn
cl_events_wait (CLEventList &event_list)
{
#define XCAM_MAX_CL_EVENT_COUNT 256
cl_event event_ids [XCAM_MAX_CL_EVENT_COUNT];
uint32_t event_count = 0;
cl_int error_code = CL_SUCCESS;
if (event_list.empty ())
return XCAM_RETURN_NO_ERROR;
xcam_mem_clear (&event_ids);
for (CLEventList::iterator iter = event_list.begin ();
iter != event_list.end (); ++iter) {
SmartPtr<CLEvent> &event = *iter;
XCAM_ASSERT (event->get_event_id ());
event_ids[event_count++] = event->get_event_id ();
if (event_count >= XCAM_MAX_CL_EVENT_COUNT)
break;
}
XCAM_ASSERT (event_count > 0);
error_code = clWaitForEvents (event_count, event_ids);
XCAM_FAIL_RETURN (
WARNING,
error_code == CL_SUCCESS,
XCAM_RETURN_ERROR_CL,
"cl events wait failed with error cod:%d", error_code);
return XCAM_RETURN_NO_ERROR;
}
VkCommandPool
VKDevice::create_cmd_pool (VkFlags queue_flag)
{
XCAM_ASSERT (XCAM_IS_VALID_VK_ID (_dev_id));
XCAM_ASSERT (_instance.ptr ());
VkCommandPool pool_id = VK_NULL_HANDLE;
VkCommandPoolCreateInfo create_pool_info = {};
create_pool_info.sType = VK_STRUCTURE_TYPE_COMMAND_POOL_CREATE_INFO;
create_pool_info.flags = VK_COMMAND_POOL_CREATE_RESET_COMMAND_BUFFER_BIT;
if (queue_flag == VK_QUEUE_COMPUTE_BIT)
create_pool_info.queueFamilyIndex = _instance->get_compute_queue_family_idx ();
else if (queue_flag == VK_QUEUE_GRAPHICS_BIT)
create_pool_info.queueFamilyIndex = _instance->get_graphics_queue_family_idx ();
else {
XCAM_LOG_WARNING ("VKDevice create command pool failed, queue_flag(%d) not supported.", queue_flag);
return VK_NULL_HANDLE;
}
XCAM_VK_CHECK_RETURN (
ERROR, vkCreateCommandPool (_dev_id, &create_pool_info, _allocator.ptr (), &pool_id),
VK_NULL_HANDLE, "VKDevice create command pool failed.");
return pool_id;
}
XCamReturn
CLKernel::load_from_source (const char *source, size_t length)
{
cl_kernel new_kernel_id = NULL;
XCAM_ASSERT (source);
if (!source) {
XCAM_LOG_WARNING ("kernel:%s source empty", XCAM_STR (_name));
return XCAM_RETURN_ERROR_PARAM;
}
if (_kernel_id) {
XCAM_LOG_WARNING ("kernel:%s already build yet", XCAM_STR (_name));
return XCAM_RETURN_ERROR_PARAM;
}
XCAM_ASSERT (_context.ptr ());
if (length == 0)
length = strlen (source);
new_kernel_id =
_context->generate_kernel_id (
this,
(const uint8_t *)source, length,
CLContext::KERNEL_BUILD_SOURCE);
XCAM_FAIL_RETURN(
WARNING,
new_kernel_id != NULL,
XCAM_RETURN_ERROR_CL,
"cl kernel(%s) load from source failed", XCAM_STR (_name));
_kernel_id = new_kernel_id;
return XCAM_RETURN_NO_ERROR;
}
void
VKDevice::destroy_desc_pool (VkDescriptorPool pool)
{
XCAM_ASSERT (XCAM_IS_VALID_VK_ID (_dev_id));
XCAM_ASSERT (XCAM_IS_VALID_VK_ID (pool));
vkDestroyDescriptorPool (_dev_id, pool, _allocator.ptr ());
}
bool
translate_3a_stats (XCam3AStats *from, struct atomisp_3a_statistics *to)
{
XCAM_ASSERT (from);
XCAM_ASSERT (to);
struct atomisp_grid_info &to_info = to->grid_info;
XCam3AStatsInfo &from_info = from->info;
uint32_t color_count = (from_info.grid_pixel_size / 2) * (from_info.grid_pixel_size / 2);
XCAM_ASSERT (to_info.bqs_per_grid_cell == 8);
for (uint32_t i = 0; i < from_info.height; ++i)
for (uint32_t j = 0; j < from_info.width; ++j) {
to->data [i * to_info.aligned_width + j].ae_y =
from->stats [i * from_info.aligned_width + j].avg_y * color_count;
to->data [i * to_info.aligned_width + j].awb_gr =
from->stats [i * from_info.aligned_width + j].avg_gr * color_count;
to->data [i * to_info.aligned_width + j].awb_r =
from->stats [i * from_info.aligned_width + j].avg_r * color_count;
to->data [i * to_info.aligned_width + j].awb_b =
from->stats [i * from_info.aligned_width + j].avg_b * color_count;
to->data [i * to_info.aligned_width + j].awb_gb =
from->stats [i * from_info.aligned_width + j].avg_gb * color_count;
to->data [i * to_info.aligned_width + j].awb_cnt =
from->stats [i * from_info.aligned_width + j].valid_wb_count;
to->data [i * to_info.aligned_width + j].af_hpf1 =
from->stats [i * from_info.aligned_width + j].f_value1;
to->data [i * to_info.aligned_width + j].af_hpf2 =
from->stats [i * from_info.aligned_width + j].f_value2;
}
return true;
}
bool
CLVaImage::init_va_image (
SmartPtr<CLContext> &context, SmartPtr<DrmBoBuffer> &bo,
const CLImageDesc &cl_desc, uint32_t offset)
{
uint32_t bo_name = 0;
cl_mem mem_id = 0;
bool need_create = true;
cl_libva_image va_image_info;
cl_import_image_info_intel import_image_info;
xcam_mem_clear (va_image_info);
xcam_mem_clear (import_image_info);
import_image_info.offset = va_image_info.offset = offset;
import_image_info.width = va_image_info.width = cl_desc.width;
import_image_info.height = va_image_info.height = cl_desc.height;
import_image_info.fmt = va_image_info.fmt = cl_desc.format;
import_image_info.row_pitch = va_image_info.row_pitch = cl_desc.row_pitch;
import_image_info.size = cl_desc.size;
import_image_info.type = CL_MEM_OBJECT_IMAGE2D;
XCAM_ASSERT (bo.ptr ());
SmartPtr<CLImageBoBuffer> cl_image_buffer = bo.dynamic_cast_ptr<CLImageBoBuffer> ();
if (cl_image_buffer.ptr ()) {
SmartPtr<CLImage> cl_image_data = cl_image_buffer->get_cl_image ();
XCAM_ASSERT (cl_image_data.ptr ());
CLImageDesc old_desc = cl_image_data->get_image_desc ();
if (cl_desc == old_desc) {
need_create = false;
mem_id = cl_image_data->get_mem_id ();
}
}
if (need_create) {
import_image_info.fd = bo->get_fd();
if (import_image_info.fd != -1)
mem_id = context->import_dma_image (import_image_info);
if (mem_id == NULL) {
if (drm_intel_bo_flink (bo->get_bo (), &bo_name) == 0) {
va_image_info.bo_name = bo_name;
mem_id = context->create_va_image (va_image_info);
}
if (mem_id == NULL) {
XCAM_LOG_WARNING ("create va image failed");
return false;
}
}
} else {
va_image_info.bo_name = uint32_t(-1);
}
set_mem_id (mem_id, need_create);
init_desc_by_image ();
_va_image_info = va_image_info;
return true;
}
void
VKDevice::free_mem_id (VkDeviceMemory mem)
{
XCAM_ASSERT (XCAM_IS_VALID_VK_ID (_dev_id));
XCAM_ASSERT (XCAM_IS_VALID_VK_ID (mem));
vkFreeMemory (_dev_id, mem, _allocator.ptr ());
}
VKDevice::VKDevice (VkDevice id, const SmartPtr<VKInstance> &instance)
: _dev_id (id)
, _instance (instance)
{
XCAM_ASSERT (instance.ptr ());
XCAM_ASSERT (XCAM_IS_VALID_VK_ID (id));
_allocator = instance->get_allocator ();
}
CLCommandQueue::CLCommandQueue (SmartPtr<CLContext> &context, cl_command_queue id)
: _context (context)
, _cmd_queue_id (id)
{
XCAM_ASSERT (context.ptr ());
XCAM_ASSERT (id);
XCAM_LOG_DEBUG ("CLCommandQueue constructed");
}
void
VKDevice::destroy_fence (VkFence fence)
{
XCAM_ASSERT (XCAM_IS_VALID_VK_ID (_dev_id));
XCAM_ASSERT (XCAM_IS_VALID_VK_ID (fence));
vkDestroyFence (_dev_id, fence, _allocator.ptr ());
}
void
VKDevice::free_cmd_buffer (VkCommandPool pool, VkCommandBuffer buf)
{
XCAM_ASSERT (XCAM_IS_VALID_VK_ID (_dev_id));
XCAM_ASSERT (XCAM_IS_VALID_VK_ID (pool));
XCAM_ASSERT (XCAM_IS_VALID_VK_ID (buf));
vkFreeCommandBuffers (_dev_id, pool, 1, &buf);
}
XCamReturn
VKDevice::reset_fence (VkFence fence)
{
XCAM_ASSERT (XCAM_IS_VALID_VK_ID (_dev_id));
XCAM_ASSERT (XCAM_IS_VALID_VK_ID (fence));
XCAM_VK_CHECK_RETURN (
ERROR, vkResetFences (_dev_id, 1, &fence),
XCAM_RETURN_ERROR_VULKAN, "VKDevice reset fence failed.");
return XCAM_RETURN_NO_ERROR;
}
bool
CLContext::init_cmd_queue (SmartPtr<CLContext> &self)
{
XCAM_ASSERT (_cmd_queue_list.empty ());
XCAM_ASSERT (self.ptr() == this);
SmartPtr<CLCommandQueue> cmd_queue = create_cmd_queue (self);
if (!cmd_queue.ptr ())
return false;
_cmd_queue_list.push_back (cmd_queue);
return true;
}
XCamReturn
CLContext::execute_kernel (
CLKernel *kernel,
CLCommandQueue *queue,
CLEventList &events_wait,
SmartPtr<CLEvent> &event_out)
{
cl_int error_code = CL_SUCCESS;
cl_command_queue cmd_queue_id = NULL;
cl_kernel kernel_id = kernel->get_kernel_id ();
uint32_t work_dims = kernel->get_work_dims ();
const size_t *global_sizes = kernel->get_work_global_size ();
const size_t *local_sizes = kernel->get_work_local_size ();
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 = 1;
XCAM_ASSERT (kernel);
if (queue == NULL) {
SmartPtr<CLCommandQueue> cmd_queue = get_default_cmd_queue ();
queue = cmd_queue.ptr ();
}
XCAM_ASSERT (queue);
cmd_queue_id = queue->get_cmd_queue_id ();
num_of_events_wait = event_list_2_id_array (events_wait, events_id_wait, XCAM_CL_MAX_EVENT_SIZE);
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,
work_dims, NULL, global_sizes, local_sizes,
num_of_events_wait, (num_of_events_wait ? events_id_wait : NULL),
event_out_id);
XCAM_FAIL_RETURN(
WARNING,
error_code == CL_SUCCESS,
XCAM_RETURN_ERROR_CL,
"execute kernel(%s) failed with error_code:%d",
kernel->get_kernel_name (), error_code);
return XCAM_RETURN_NO_ERROR;
}
static SmartPtr<CLImageKernel>
create_scale_kernel (
const SmartPtr<CLContext> &context, SmartPtr<CLImageScaler> &handler, CLImageScalerMemoryLayout layout)
{
SmartPtr<CLImageKernel> kernel;
kernel = new CLImageScalerKernel (context, layout, handler);
XCAM_ASSERT (kernel.ptr ());
XCAM_FAIL_RETURN (
ERROR, kernel->build_kernel (kernel_scale_info, NULL) == XCAM_RETURN_NO_ERROR, NULL,
"build scaler kernel(%s) failed", kernel_scale_info.kernel_name);
XCAM_ASSERT (kernel->is_valid ());
return kernel;
}
XCamReturn
CLKernel::execute (
const SmartPtr<CLKernel> self,
bool block,
CLEventList &events,
SmartPtr<CLEvent> &event_out)
{
XCAM_ASSERT (self.ptr () == this);
XCAM_ASSERT (_context.ptr ());
SmartPtr<CLEvent> kernel_event = event_out;
if (!block && !kernel_event.ptr ()) {
kernel_event = new CLEvent;
}
#if ENABLE_DEBUG_KERNEL
XCAM_OBJ_PROFILING_START;
#endif
XCamReturn ret = _context->execute_kernel (self, NULL, events, kernel_event);
XCAM_FAIL_RETURN (
ERROR,
ret == XCAM_RETURN_NO_ERROR,
ret,
"kernel(%s) execute failed", XCAM_STR(_name));
if (block) {
_context->finish ();
} else {
XCAM_ASSERT (kernel_event.ptr () && kernel_event->get_event_id ());
KernelUserData *user_data = new KernelUserData (self, kernel_event);
user_data->arg_list.swap (_arg_list);
ret = _context->set_event_callback (kernel_event, CL_COMPLETE, event_notify, user_data);
if (ret != XCAM_RETURN_NO_ERROR) {
XCAM_LOG_WARNING ("kernel(%s) set event callback failed", XCAM_STR (_name));
_context->finish ();
delete user_data;
}
}
_arg_list.clear ();
#if ENABLE_DEBUG_KERNEL
_context->finish ();
char name[1024];
snprintf (name, 1024, "%s-%p", XCAM_STR (_name), this);
XCAM_OBJ_PROFILING_END (name, XCAM_OBJ_DUR_FRAME_NUM);
#endif
return ret;
}
XCamReturn
VKDevice::free_desc_set (VkDescriptorSet set, VkDescriptorPool pool)
{
XCAM_ASSERT (XCAM_IS_VALID_VK_ID (_dev_id));
XCAM_ASSERT (XCAM_IS_VALID_VK_ID (set));
XCAM_ASSERT (XCAM_IS_VALID_VK_ID (pool));
XCAM_VK_CHECK_RETURN (
ERROR,
vkFreeDescriptorSets (_dev_id, pool, 1, &set),
XCAM_RETURN_ERROR_VULKAN,
"vkdevice free desriptor set from pool failed");
return XCAM_RETURN_NO_ERROR;
}
CLKernel::CLKernel(SmartPtr<CLContext> &context, const char *name)
: _name (NULL)
, _kernel_id (NULL)
, _context (context)
, _work_dim (0)
{
XCAM_ASSERT (context.ptr ());
XCAM_ASSERT (name);
if (name)
_name = strdup (name);
set_default_work_size ();
}
SmartPtr<GLBuffer>
GLBuffer::create_buffer (
GLenum target,
const GLvoid *data, uint32_t size,
GLenum usage)
{
XCAM_ASSERT (size > 0);
GLuint buf_id = 0;
glGenBuffers (1, &buf_id);
GLenum error = gl_error ();
XCAM_FAIL_RETURN (
ERROR, buf_id && (error == GL_NO_ERROR), NULL,
"GL buffer creation failed, error flag: %s", gl_error_string (error));
glBindBuffer (target, buf_id);
XCAM_FAIL_RETURN (
ERROR, (error = gl_error ()) == GL_NO_ERROR, NULL,
"GL buffer creation failed when bind buffer:%d, error flag: %s",
buf_id, gl_error_string (error));
glBufferData (target, size, data, usage);
XCAM_FAIL_RETURN (
ERROR, (error = gl_error ()) == GL_NO_ERROR, NULL,
"GL buffer creation failed in glBufferData, id:%d, error flag: %s",
buf_id, gl_error_string (error));
SmartPtr<GLBuffer> buf_obj =
new GLBuffer (buf_id, target, usage, size);
return buf_obj;
}
SmartPtr<CLImageHandler>
create_cl_tnr_image_handler (SmartPtr<CLContext> &context, CLTnrType type)
{
SmartPtr<CLTnrImageHandler> tnr_handler;
SmartPtr<CLTnrImageKernel> tnr_kernel;
XCamReturn ret = XCAM_RETURN_NO_ERROR;
XCAM_CL_KERNEL_FUNC_SOURCE_BEGIN(kernel_tnr_yuv)
#include "kernel_tnr_yuv.clx"
XCAM_CL_KERNEL_FUNC_END;
XCAM_CL_KERNEL_FUNC_SOURCE_BEGIN(kernel_tnr_rgb)
#include "kernel_tnr_rgb.clx"
XCAM_CL_KERNEL_FUNC_END;
if (CL_TNR_TYPE_YUV == type) {
tnr_kernel = new CLTnrImageKernel (context, "kernel_tnr_yuv", CL_TNR_TYPE_YUV);
ret = tnr_kernel->load_from_source (kernel_tnr_yuv_body, strlen (kernel_tnr_yuv_body));
} else if (CL_TNR_TYPE_RGB == type) {
tnr_kernel = new CLTnrImageKernel (context, "kernel_tnr_rgb", CL_TNR_TYPE_RGB);
ret = tnr_kernel->load_from_source (kernel_tnr_rgb_body, strlen (kernel_tnr_rgb_body));
}
XCAM_FAIL_RETURN (
WARNING,
ret == XCAM_RETURN_NO_ERROR,
NULL,
"CL image handler(%s) load source failed", tnr_kernel->get_kernel_name());
tnr_handler = new CLTnrImageHandler ("cl_handler_tnr");
XCAM_ASSERT (tnr_kernel->is_valid ());
tnr_handler->set_tnr_kernel (tnr_kernel);
return tnr_handler;
}
XCamReturn
CLDemosaicImageKernel::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 = output->get_video_info ();
_image_in = new CLVaImage (context, input);
_image_out = new CLVaImage (context, output);
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 ());
//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);
arg_count = 2;
work_size.dim = XCAM_DEFAULT_IMAGE_DIM;
work_size.global[0] = video_info.width / 2;
work_size.global[1] = video_info.height / 2;
work_size.local[0] = 4;
work_size.local[1] = 4;
return XCAM_RETURN_NO_ERROR;
}
SmartPtr<CLImageHandler>
create_cl_demosaic_image_handler (SmartPtr<CLContext> &context)
{
SmartPtr<CLImageHandler> demosaic_handler;
SmartPtr<CLImageKernel> demosaic_kernel;
XCamReturn ret = XCAM_RETURN_NO_ERROR;
demosaic_kernel = new CLDemosaicImageKernel (context);
{
XCAM_CL_KERNEL_FUNC_SOURCE_BEGIN(kernel_demosaic)
#include "kernel_demosaic.cl"
XCAM_CL_KERNEL_FUNC_END;
ret = demosaic_kernel->load_from_source (kernel_demosaic_body, strlen (kernel_demosaic_body));
XCAM_FAIL_RETURN (
WARNING,
ret == XCAM_RETURN_NO_ERROR,
NULL,
"CL image handler(%s) load source failed", demosaic_kernel->get_kernel_name());
}
XCAM_ASSERT (demosaic_kernel->is_valid ());
demosaic_handler = new CLBayer2RGBImageHandler ("cl_handler_demosaic");
demosaic_handler->add_kernel (demosaic_kernel);
return demosaic_handler;
}
