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; }