int init_sin ()
{
cl_int error;
if (sin_init)
return 1;
context = get_cl_context();
device = get_cl_device();
const char *src=sin_cl;
size_t srcsize=strlen(sin_cl);
const char *srcptr[]={src};
// build CL program
error = buildcl (srcptr, &srcsize, &prog, "", NUM_GPUS);
// create kernel
k_sin = clCreateKernel(prog, "sin_cl", &error);
// get the shared CQ
cq = get_command_queue();
if (!error)
sin_init = 1;
return error;
}
int init_rot13 ()
{
cl_int error;
if (rot13_init)
return 1;
context = get_cl_context();
device = get_cl_device();
const char *src=rot13_cl;
size_t srcsize=strlen(rot13_cl);
const char *srcptr[]={src};
// build CL program
error = buildcl (srcptr, &srcsize, &prog, "", NUM_GPUS);
// create kernel
k_rot13 = clCreateKernel(prog, "rot13", &error);
// get the shared CQ
cq = get_command_queue();
// we are initialised
if (!error)
rot13_init = 1;
return error;
}
XCamReturn
CL3aImageProcessor::create_handlers ()
{
SmartPtr<CLImageHandler> image_handler;
SmartPtr<CLContext> context = get_cl_context ();
XCAM_ASSERT (context.ptr ());
/* black leve as first */
image_handler = create_cl_blc_image_handler (context);
_black_level = image_handler;
XCAM_FAIL_RETURN (
WARNING,
image_handler.ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create blc handler failed");
add_handler (image_handler);
/* hdr */
if (_enable_hdr) {
image_handler = create_cl_hdr_image_handler (context, CL_HDR_TYPE_RGB);
_hdr = image_handler;
XCAM_FAIL_RETURN (
WARNING,
_hdr.ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create hdr handler failed");
add_handler (image_handler);
}
/* demosaic */
image_handler = create_cl_demosaic_image_handler (context);
_demosaic = image_handler.dynamic_cast_ptr<CLBayer2RGBImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
_demosaic.ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create demosaic handler failed");
add_handler (image_handler);
/* color space conversion */
if (_out_smaple_type == OutSampleYuv) {
image_handler = create_cl_csc_image_handler (context, CL_CSC_TYPE_RGBATONV12);
_csc = image_handler.dynamic_cast_ptr<CLCscImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
_csc .ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create csc handler failed");
add_handler (image_handler);
} else if (_out_smaple_type == OutSampleRGB) {
_demosaic->set_output_format (_output_fourcc);
}
return XCAM_RETURN_NO_ERROR;
}
int init_mandelbrot ()
{
int i;
cl_int error;
if (mandelbrot_init)
return 1;
context = get_cl_context();
device = get_cl_device();
FILE *fp;
fp = fopen(CLKERNELDEFS, "r");
if (!fp) {
fprintf(stderr, "Failed to load kernel.\n");
return(1);
}
char *src = (char*) malloc (MAX_SOURCE_SIZE);
size_t srcsize = fread (src, 1, MAX_SOURCE_SIZE, fp);
fclose (fp);
const char *srcptr[]={src};
// get the number of GPUS/DEVICES available
numdevices = getNumDevices();
#if CAN_USE_DOUBLE
// build CL program with a USE_DOUBLE define if we found the correct extension
if (getCorrectDevice("cl_khr_fp64") == CL_SUCCESS) {
error = buildcl (srcptr, &srcsize, &prog[0], "-D USE_DOUBLE -cl-fast-relaxed-math -cl-mad-enable", numdevices);
}
else {
#endif
mandelbrot_cl_float = 1;
error = buildcl (srcptr, &srcsize, &prog[0], "-D USE_FLOAT -cl-fast-relaxed-math -cl-mad-enable", numdevices);
#if CAN_USE_DOUBLE
}
#endif
// create kernel
for (i=0; i<numdevices; i++) {
k_mandelbrot[i] = clCreateKernel(prog[i], "mandelbrot", &error);
}
// get the shared CQ
cq = get_command_queue();
if (!error)
mandelbrot_init = 1;
return error;
}
int mandelbrotvis (cl_int *data, cl_fract *job)
{
cl_int error;
int i;
#if MULTI_GPUS
int currentdevice = nextDevice();
#else
int currentdevice = 0;
#endif
cl_command_queue *cqm = get_command_queue();
cl_context *cm = get_cl_context();
// Allocate memory for the kernel to work with
cl_mem mem1, mem2;
mem1 = clCreateBuffer(*cm, CL_MEM_WRITE_ONLY, sizeof(cl_int)*(visheight*viswidth*framesperworker), 0, &error);
if (mandelbrot_cl_float) {
cl_float jobfloat[framesperworker*JOBS_PER_FRAME];
for (i=0; i<framesperworker*JOBS_PER_FRAME; i++)
jobfloat[i] = (cl_float) job[i];
mem2 = clCreateBuffer(*cm, CL_MEM_COPY_HOST_PTR, sizeof(cl_float)*framesperworker*JOBS_PER_FRAME, jobfloat, &error);
} else {
mem2 = clCreateBuffer(*cm, CL_MEM_COPY_HOST_PTR, sizeof(cl_fract)*framesperworker*JOBS_PER_FRAME, job, &error);
}
// get a handle and map parameters for the kernel
error = clSetKernelArg(k_mandelbrotvis[currentdevice], 0, sizeof(mem1), &mem1);
error = clSetKernelArg(k_mandelbrotvis[currentdevice], 1, sizeof(mem2), &mem2);
size_t worksize[3] = {visheight, viswidth, framesperworker};
error = clEnqueueNDRangeKernel(*cqm, k_mandelbrotvis[currentdevice], 3, NULL, &worksize[0], 0, 0, 0, 0);
// Read the result back into data
error = clEnqueueReadBuffer(*cqm, mem1, CL_TRUE, 0, sizeof(cl_int)*(visheight*viswidth*framesperworker), data, 0, 0, 0);
// cleanup - don't perform a flush as the queue is now shared between all executions. The
// blocking clEnqueueReadBuffer should be enough
clReleaseMemObject(mem1);
clReleaseMemObject(mem2);
if (error) {
fprintf (stderr, "ERROR! : %s\n", errorMessageCL(error));
exit(10);
}
return error;
}
XCamReturn
CLCscImageProcessor::create_handlers ()
{
SmartPtr<CLImageHandler> image_handler;
SmartPtr<CLContext> context = get_cl_context ();
XCAM_ASSERT (context.ptr ());
/* color space conversion */
image_handler = create_cl_csc_image_handler (context, CL_CSC_TYPE_YUYVTORGBA);
_csc = image_handler.dynamic_cast_ptr<CLCscImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
_csc .ptr (),
XCAM_RETURN_ERROR_CL,
"CLCscImageProcessor create csc handler failed");
image_handler->set_pool_type (CLImageHandler::CLVideoPoolType);
add_handler (image_handler);
return XCAM_RETURN_NO_ERROR;
}
XCamReturn
CL3aImageProcessor::create_handlers ()
{
SmartPtr<CLImageHandler> image_handler;
SmartPtr<CLContext> context = get_cl_context ();
XCAM_ASSERT (context.ptr ());
#if 1
/* bayer pipeline */
image_handler = create_cl_bayer_pipe_image_handler (context);
_bayer_pipe = image_handler.dynamic_cast_ptr<CLBayerPipeImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
image_handler.ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create bayer pipe handler failed");
_bayer_pipe->set_stats_callback (_stats_callback);
#if 0
if (get_profile () >= AdvancedPipelineProfile) {
_bayer_pipe->set_output_format (V4L2_PIX_FMT_ABGR32);
}
#endif
_bayer_pipe->enable_denoise (XCAM_DENOISE_TYPE_BNR & _snr_mode);
_bayer_pipe->enable_gamma (_enable_gamma);
image_handler->set_pool_size (XCAM_CL_3A_IMAGE_MAX_POOL_SIZE * 2);
add_handler (image_handler);
if(_capture_stage == BasicbayerStage)
return XCAM_RETURN_NO_ERROR;
#else
/* black leve as first */
image_handler = create_cl_blc_image_handler (context);
_black_level = image_handler.dynamic_cast_ptr<CLBlcImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
image_handler.ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create blc handler failed");
add_handler (image_handler);
image_handler = create_cl_dpc_image_handler (context);
_dpc = image_handler.dynamic_cast_ptr<CLDpcImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
image_handler.ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create dpc handler failed");
_dpc->set_kernels_enable(_enable_dpc);
add_handler (image_handler);
image_handler = create_cl_bnr_image_handler (context);
_bnr = image_handler.dynamic_cast_ptr<CLBnrImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
_bnr.ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create bnr handler failed");
_bnr->set_kernels_enable (XCAM_DENOISE_TYPE_BNR & _snr_mode);
add_handler (image_handler);
image_handler = create_cl_3a_stats_image_handler (context);
_x3a_stats_calculator = image_handler.dynamic_cast_ptr<CL3AStatsCalculator> ();
XCAM_FAIL_RETURN (
WARNING,
_x3a_stats_calculator.ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create 3a stats calculator failed");
_x3a_stats_calculator->set_stats_callback (_stats_callback);
add_handler (image_handler);
image_handler = create_cl_wb_image_handler (context);
_wb = image_handler.dynamic_cast_ptr<CLWbImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
_wb.ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create whitebalance handler failed");
add_handler (image_handler);
/* gamma */
image_handler = create_cl_gamma_image_handler (context);
_gamma = image_handler.dynamic_cast_ptr<CLGammaImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
_gamma.ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create gamma handler failed");
_gamma->set_kernels_enable (_enable_gamma);
add_handler (image_handler);
/* hdr */
image_handler = create_cl_hdr_image_handler (context, CL_HDR_DISABLE);
_hdr = image_handler.dynamic_cast_ptr<CLHdrImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
_hdr.ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create hdr handler failed");
if(_hdr_mode == CL_HDR_TYPE_RGB)
_hdr->set_mode (_hdr_mode);
add_handler (image_handler);
/* demosaic */
image_handler = create_cl_demosaic_image_handler (context);
_demosaic = image_handler.dynamic_cast_ptr<CLBayer2RGBImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
_demosaic.ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create demosaic handler failed");
image_handler->set_pool_size (XCAM_CL_3A_IMAGE_MAX_POOL_SIZE);
add_handler (image_handler);
#endif
/* hdr-lab*/
image_handler = create_cl_hdr_image_handler (context, CL_HDR_DISABLE);
_hdr = image_handler.dynamic_cast_ptr<CLHdrImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
_hdr.ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create hdr handler failed");
if(_hdr_mode == CL_HDR_TYPE_LAB)
_hdr->set_mode (_hdr_mode);
image_handler->set_pool_size (XCAM_CL_3A_IMAGE_MAX_POOL_SIZE);
add_handler (image_handler);
/* bilateral noise reduction */
image_handler = create_cl_denoise_image_handler (context);
_binr = image_handler.dynamic_cast_ptr<CLDenoiseImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
_binr.ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create denoise handler failed");
_binr->set_kernels_enable (XCAM_DENOISE_TYPE_BILATERAL & _snr_mode);
image_handler->set_pool_size (XCAM_CL_3A_IMAGE_MAX_POOL_SIZE);
add_handler (image_handler);
#if 0
image_handler = create_cl_rgb_pipe_image_handler (context);
_rgb_pipe = image_handler.dynamic_cast_ptr<CLRgbPipeImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
_rgb_pipe.ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create rgb pipe handler failed");
_rgb_pipe->set_kernels_enable (get_profile () >= AdvancedPipelineProfile);
add_handler (image_handler);
/* Temporal Noise Reduction (RGB domain) */
image_handler = create_cl_tnr_image_handler(context, CL_TNR_TYPE_RGB);
_tnr_rgb = image_handler.dynamic_cast_ptr<CLTnrImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
_tnr_rgb.ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create tnr handler failed");
_tnr_rgb->set_mode (CL_TNR_TYPE_RGB & _tnr_mode);
add_handler (image_handler);
#else
/* simple noise reduction */
image_handler = create_cl_snr_image_handler (context);
_snr = image_handler.dynamic_cast_ptr<CLSnrImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
_snr.ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create snr handler failed");
_snr->set_kernels_enable (XCAM_DENOISE_TYPE_SIMPLE & _snr_mode);
image_handler->set_pool_size (XCAM_CL_3A_IMAGE_MAX_POOL_SIZE);
add_handler (image_handler);
#endif
/* tone mapping*/
image_handler = create_cl_tonemapping_image_handler (context);
_tonemapping = image_handler.dynamic_cast_ptr<CLTonemappingImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
_tonemapping.ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create tonemapping handler failed");
_tonemapping->set_kernels_enable (_enable_tonemapping);
image_handler->set_pool_size (XCAM_CL_3A_IMAGE_MAX_POOL_SIZE);
add_handler (image_handler);
#if 1
image_handler = create_cl_yuv_pipe_image_handler (context);
_yuv_pipe = image_handler.dynamic_cast_ptr<CLYuvPipeImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
_yuv_pipe.ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create macc handler failed");
_yuv_pipe->set_tnr_enable (_tnr_mode & CL_TNR_TYPE_RGB, _tnr_mode & CL_TNR_TYPE_YUV);
image_handler->set_pool_size (XCAM_CL_3A_IMAGE_MAX_POOL_SIZE);
add_handler (image_handler);
#else
/* macc */
image_handler = create_cl_macc_image_handler (context);
_macc = image_handler.dynamic_cast_ptr<CLMaccImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
_macc.ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create macc handler failed");
_macc->set_kernels_enable (_enable_macc);
add_handler (image_handler);
/* color space conversion */
image_handler = create_cl_csc_image_handler (context, CL_CSC_TYPE_RGBATONV12);
_csc = image_handler.dynamic_cast_ptr<CLCscImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
_csc .ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create csc handler failed");
image_handler->set_pool_type (CLImageHandler::DrmBoPoolType);
add_handler (image_handler);
/* Temporal Noise Reduction (YUV domain) */
image_handler = create_cl_tnr_image_handler(context, CL_TNR_TYPE_YUV);
_tnr_yuv = image_handler.dynamic_cast_ptr<CLTnrImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
_tnr_yuv.ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create tnr handler failed");
_tnr_yuv->set_mode (CL_TNR_TYPE_YUV & _tnr_mode);
image_handler->set_pool_type (CLImageHandler::DrmBoPoolType);
add_handler (image_handler);
#endif
/* ee */
image_handler = create_cl_ee_image_handler (context);
_ee = image_handler.dynamic_cast_ptr<CLEeImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
_ee.ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create ee handler failed");
_ee->set_kernels_enable (XCAM_DENOISE_TYPE_EE & _snr_mode);
image_handler->set_pool_type (CLImageHandler::DrmBoPoolType);
image_handler->set_pool_size (XCAM_CL_3A_IMAGE_MAX_POOL_SIZE);
add_handler (image_handler);
/* biyuv */
image_handler = create_cl_biyuv_image_handler (context);
_biyuv = image_handler.dynamic_cast_ptr<CLBiyuvImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
_biyuv.ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create biyuv handler failed");
_biyuv->set_kernels_enable (XCAM_DENOISE_TYPE_BIYUV & _snr_mode);
image_handler->set_pool_type (CLImageHandler::DrmBoPoolType);
image_handler->set_pool_size (XCAM_CL_3A_IMAGE_MAX_POOL_SIZE);
add_handler (image_handler);
/* image scaler */
image_handler = create_cl_image_scaler_handler (context, V4L2_PIX_FMT_NV12);
_scaler = image_handler.dynamic_cast_ptr<CLImageScaler> ();
XCAM_FAIL_RETURN (
WARNING,
_scaler.ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create scaler handler failed");
_scaler->set_scaler_factor (XCAM_CL_3A_IMAGE_SCALER_FACTOR);
_scaler->set_buffer_callback (_stats_callback);
image_handler->set_pool_type (CLImageHandler::DrmBoPoolType);
image_handler->set_kernels_enable (false);
add_handler (image_handler);
if (_out_smaple_type == OutSampleRGB) {
image_handler = create_cl_csc_image_handler (context, CL_CSC_TYPE_NV12TORGBA);
_csc = image_handler.dynamic_cast_ptr<CLCscImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
_csc .ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create csc handler failed");
image_handler->set_pool_type (CLImageHandler::DrmBoPoolType);
image_handler->set_pool_size (XCAM_CL_3A_IMAGE_MAX_POOL_SIZE);
add_handler (image_handler);
}
return XCAM_RETURN_NO_ERROR;
}
int init_mandelbrotvis ()
{
int i;
cl_int error;
if (mandelbrotvis_init)
return 1;
context = get_cl_context();
device = get_cl_device();
FILE *fp;
fp = fopen(CLVISKERNELDEFS, "r");
if (!fp) {
fprintf(stderr, "Failed to load kernel.\n");
return(1);
}
char *src = (char*) malloc (MAX_SOURCE_SIZE);
size_t srcsize = fread (src, 1, MAX_SOURCE_SIZE, fp);
fclose (fp);
const char *srcptr[]={src};
// get the number of GPUS/DEVICES available
numdevices = getNumDevices();
// build CL program with a USE_DOUBLE define if we found the correct extension
char *precision = " ";
#if CAN_USE_DOUBLE
if (getCorrectDevice("cl_khr_fp64") == CL_SUCCESS) {
precision = "-D USE_DOUBLE";
}
else {
#endif
mandelbrot_cl_float = 1;
precision = "-D USE_FLOAT";
#if CAN_USE_DOUBLE
}
#endif
char options[MAX_BUILD_LINE_LENGTH];
// following options seem to speed things up a little
char *compile_opt = "-cl-fast-relaxed-math -cl-mad-enable";
snprintf(options, MAX_BUILD_LINE_LENGTH,
"%s -D IMAGEWIDTHVIS=%d -D IMAGEHEIGHTVIS=%d %s", precision, viswidth, visheight, compile_opt);
error = buildcl (srcptr, &srcsize, &progvis[0], options, numdevices);
// printf("%s\n\n\n", options);
// create kernel
for (i=0; i<numdevices; i++) {
k_mandelbrotvis[i] = clCreateKernel(progvis[i], "mandelbrot_vis", &error);
}
// get the shared CQ
cq = get_command_queue();
// initialise the jobs array
initialiseJobs();
if (!error)
mandelbrotvis_init = 1;
return error;
}
int mandelbrot (cl_char (*data)[200], cl_fract *job)
{
cl_int error;
int i;
#if ERROR_CHECK
if (prog == NULL) {
init_mandelbrot();
}
#endif
#if MULTI_GPUS
// move to new context/cq
int currentdevice = nextDevice();
cl_command_queue *cq = get_command_queue();
cl_context *context = get_cl_context();
#else
int currentdevice = 0;
#endif
// Allocate memory for the kernel to work with
cl_mem mem1, mem2;
mem1 = clCreateBuffer(*context, CL_MEM_WRITE_ONLY, sizeof(cl_char)*(IMAGEHEIGHT*IMAGEWIDTH*2), 0, &error);
if (mandelbrot_cl_float) {
cl_float jobfloat[4];
for (i=0; i<4; i++)
jobfloat[i] = (cl_float) job[i];
mem2 = clCreateBuffer(*context, CL_MEM_COPY_HOST_PTR, sizeof(cl_float)*4, jobfloat, &error);
} else {
mem2 = clCreateBuffer(*context, CL_MEM_COPY_HOST_PTR, sizeof(cl_fract)*4, job, &error);
}
// get a handle and map parameters for the kernel
error = clSetKernelArg(k_mandelbrot[currentdevice], 0, sizeof(cl_mem), &mem1);
error = clSetKernelArg(k_mandelbrot[currentdevice], 1, sizeof(cl_mem), &mem2);
// Perform the operation (width is 100 in this example)
size_t worksize[3] = {IMAGEHEIGHT, IMAGEWIDTH, 0};
error = clEnqueueNDRangeKernel(*cq, k_mandelbrot[currentdevice], 2, NULL, &worksize[0], 0, 0, 0, 0);
// Read the result back into data
error = clEnqueueReadBuffer(*cq, mem1, CL_TRUE, 0, sizeof(cl_char)*(IMAGEHEIGHT*IMAGEWIDTH*2), data, 0, 0, 0);
// cleanup - don't perform a flush as the queue is now shared between all executions. The
// blocking clEnqueueReadBuffer should be enough
clReleaseMemObject(mem1);
clReleaseMemObject(mem2);
if (error) {
fprintf (stderr, "ERROR! : %s\n", errorMessageCL(error));
exit(10);
}
#if C_PRINT
// this will print a frame coming out of the CL kernel in a dirty but functional manner
int j;
int colour = -1;
for (i=0; i < IMAGEHEIGHT; i++) {
for (j=0; j < IMAGEWIDTH*2; j++) {
if (colour != data[i][j]) {
colour = data[i][j];
textcolour(colour);
}
j++;
fprintf (stdout, "%c", data[i][j]);
}
fprintf(stdout, "\n");
}
#endif
return error;
}
XCamReturn
CL3aImageProcessor::create_handlers ()
{
SmartPtr<CLImageHandler> image_handler;
SmartPtr<CLContext> context = get_cl_context ();
XCAM_ASSERT (context.ptr ());
/* bayer pipeline */
image_handler = create_cl_bayer_basic_image_handler (context, _enable_gamma, _3a_stats_bits);
_bayer_basic_pipe = image_handler.dynamic_cast_ptr<CLBayerBasicImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
_bayer_basic_pipe.ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create bayer basic pipe handler failed");
image_handler->set_pool_size (XCAM_CL_3A_IMAGE_MAX_POOL_SIZE);
_bayer_basic_pipe->set_stats_callback (_stats_callback);
add_handler (image_handler);
/* tone mapping */
switch(_wdr_mode) {
case Gaussian: {
image_handler = create_cl_tonemapping_image_handler (context);
_tonemapping = image_handler.dynamic_cast_ptr<CLTonemappingImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
_tonemapping.ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create tonemapping handler failed");
_tonemapping->set_kernels_enable (true);
image_handler->set_pool_size (XCAM_CL_3A_IMAGE_MAX_POOL_SIZE);
add_handler (image_handler);
break;
}
case Haleq: {
image_handler = create_cl_newtonemapping_image_handler (context);
_newtonemapping = image_handler.dynamic_cast_ptr<CLNewTonemappingImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
_newtonemapping.ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create tonemapping handler failed");
_newtonemapping->set_kernels_enable (true);
image_handler->set_pool_size (XCAM_CL_3A_IMAGE_MAX_POOL_SIZE);
add_handler (image_handler);
break;
}
default:
XCAM_LOG_DEBUG ("WDR disabled");
break;
}
/* bayer pipe */
image_handler = create_cl_bayer_pipe_image_handler (context);
_bayer_pipe = image_handler.dynamic_cast_ptr<CLBayerPipeImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
image_handler.ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create bayer pipe handler failed");
_bayer_pipe->enable_denoise (XCAM_DENOISE_TYPE_BNR & _snr_mode);
image_handler->set_pool_size (XCAM_CL_3A_IMAGE_MAX_POOL_SIZE * 2);
//image_handler->set_pool_type (CLImageHandler::DrmBoPoolType);
add_handler (image_handler);
if(_capture_stage == BasicbayerStage)
return XCAM_RETURN_NO_ERROR;
image_handler = create_cl_yuv_pipe_image_handler (context);
_yuv_pipe = image_handler.dynamic_cast_ptr<CLYuvPipeImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
_yuv_pipe.ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create yuv pipe handler failed");
_yuv_pipe->set_tnr_enable (_tnr_mode & CL_TNR_TYPE_YUV);
image_handler->set_pool_size (XCAM_CL_3A_IMAGE_MAX_POOL_SIZE * 2);
add_handler (image_handler);
#if ENABLE_YEENR_HANDLER
/* ee */
image_handler = create_cl_ee_image_handler (context);
_ee = image_handler.dynamic_cast_ptr<CLEeImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
_ee.ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create ee handler failed");
_ee->set_kernels_enable (XCAM_DENOISE_TYPE_EE & _snr_mode);
image_handler->set_pool_type (CLImageHandler::DrmBoPoolType);
image_handler->set_pool_size (XCAM_CL_3A_IMAGE_MAX_POOL_SIZE);
add_handler (image_handler);
#endif
/* wavelet denoise */
switch (_wavelet_basis) {
case CL_WAVELET_HAT: {
image_handler = create_cl_wavelet_denoise_image_handler (context, _wavelet_channel);
_wavelet = image_handler.dynamic_cast_ptr<CLWaveletDenoiseImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
_wavelet.ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create wavelet denoise handler failed");
_wavelet->set_kernels_enable (true);
image_handler->set_pool_type (CLImageHandler::DrmBoPoolType);
image_handler->set_pool_size (XCAM_CL_3A_IMAGE_MAX_POOL_SIZE);
add_handler (image_handler);
break;
}
case CL_WAVELET_HAAR: {
image_handler = create_cl_newwavelet_denoise_image_handler (context, _wavelet_channel, _wavelet_bayes_shrink);
_newwavelet = image_handler.dynamic_cast_ptr<CLNewWaveletDenoiseImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
_newwavelet.ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create new wavelet denoise handler failed");
_newwavelet->set_kernels_enable (true);
image_handler->set_pool_type (CLImageHandler::DrmBoPoolType);
image_handler->set_pool_size (XCAM_CL_3A_IMAGE_MAX_POOL_SIZE);
add_handler (image_handler);
break;
}
case CL_WAVELET_DISABLED:
default :
XCAM_LOG_DEBUG ("unknown or disable wavelet (%d)", _wavelet_basis);
break;
}
/* image scaler */
image_handler = create_cl_image_scaler_handler (context, V4L2_PIX_FMT_NV12);
_scaler = image_handler.dynamic_cast_ptr<CLImageScaler> ();
XCAM_FAIL_RETURN (
WARNING,
_scaler.ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create scaler handler failed");
_scaler->set_scaler_factor (_scaler_factor);
_scaler->set_buffer_callback (_stats_callback);
image_handler->set_pool_type (CLImageHandler::DrmBoPoolType);
image_handler->set_kernels_enable (_enable_scaler);
add_handler (image_handler);
/* wire frame */
image_handler = create_cl_wire_frame_image_handler (context);
_wire_frame = image_handler.dynamic_cast_ptr<CLWireFrameImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
_wire_frame.ptr (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor create wire frame handler failed");
_wire_frame->set_kernels_enable (_enable_wireframe);
image_handler->set_pool_type (CLImageHandler::DrmBoPoolType);
image_handler->set_pool_size (XCAM_CL_3A_IMAGE_MAX_POOL_SIZE);
add_handler (image_handler);
XCAM_FAIL_RETURN (
WARNING,
post_config (),
XCAM_RETURN_ERROR_CL,
"CL3aImageProcessor post_config failed");
return XCAM_RETURN_NO_ERROR;
}
int main(int argc, char **argv)
{
cl_int err = 0;
cl_context context = 0;
cl_device_id * devices = NULL;
cl_command_queue queue = 0;
cl_program program = 0;
cl_mem cl_a = 0, cl_b = 0, cl_res = 0;
cl_kernel adder = 0;
cl_event event;
// The iteration variable
int i;
// Define our data set
cl_float a[DATA_SIZE], b[DATA_SIZE], res[DATA_SIZE];
// Initialize array
srand(time(0));
for (i = 0; i < DATA_SIZE; i++) {
a[i] = (rand() % 100) / 100.0;
b[i] = (rand() % 100) / 100.0;
res[i] = 0;
}
check_release(get_cl_context(&context, &devices, 0) == false,
"Fail to create context");
// Specify the queue to be profile-able
queue = clCreateCommandQueue(context, devices[0], CL_QUEUE_PROFILING_ENABLE, 0);
check_release(queue == NULL, "Can't create command queue");
program = load_program(context, devices[0], "shader.cl");
check_release(program == NULL, "Fail to build program");
cl_a =
clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(cl_float) * DATA_SIZE, NULL, NULL);
cl_b =
clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(cl_float) * DATA_SIZE, NULL, NULL);
cl_res = clCreateBuffer(
context, CL_MEM_WRITE_ONLY, sizeof(cl_float) * DATA_SIZE, NULL, NULL);
if (cl_a == 0 || cl_b == 0 || cl_res == 0) {
printf("Can't create OpenCL buffer\n");
goto release;
}
check_release(clEnqueueWriteBuffer(
queue, cl_a, CL_TRUE, 0, sizeof(cl_float) * DATA_SIZE, a, 0, 0, 0),
"Write Buffer 1");
check_release(clEnqueueWriteBuffer(
queue, cl_b, CL_TRUE, 0, sizeof(cl_float) * DATA_SIZE, b, 0, 0, 0),
"Write Buffer 2");
adder = clCreateKernel(program, "adder", &err);
if (err == CL_INVALID_KERNEL_NAME) printf("CL_INVALID_KERNEL_NAME\n");
check_release(adder == NULL, "Can't load kernel");
clSetKernelArg(adder, 0, sizeof(cl_mem), &cl_a);
clSetKernelArg(adder, 1, sizeof(cl_mem), &cl_b);
clSetKernelArg(adder, 2, sizeof(cl_mem), &cl_res);
size_t work_size = DATA_SIZE;
check_release(clEnqueueNDRangeKernel(queue, adder, 1, 0, &work_size, 0, 0, 0, &event),
"Can't enqueue kernel");
check_release(
clEnqueueReadBuffer(
queue, cl_res, CL_TRUE, 0, sizeof(cl_float) * DATA_SIZE, res, 0, 0, 0),
"Can't enqueue read buffer");
clWaitForEvents(1, &event);
printf("Execution Time: %.04lf ms\n\n", get_event_exec_time(event));
// Make sure everything is done before we do anything
clFinish(queue);
err = 0;
for (i = 0; i < DATA_SIZE; i++) {
if (res[i] != a[i] + b[i]) {
printf("%f + %f = %f(answer %f)\n", a[i], b[i], res[i], a[i] + b[i]);
err++;
}
}
if (err == 0)
printf("Validation passed\n");
else
printf("Validation failed\n");
printf("------\n");
//--------------------------------
// Second test
for (i = 0; i < DATA_SIZE; i++) {
a[i] = i;
b[i] = i;
res[i] = 0;
}
check_err(clEnqueueWriteBuffer(
queue, cl_a, CL_TRUE, 0, sizeof(cl_float) * DATA_SIZE, a, 0, 0, 0),
"Write Buffer 1");
check_err(clEnqueueWriteBuffer(
queue, cl_b, CL_TRUE, 0, sizeof(cl_float) * DATA_SIZE, b, 0, 0, 0),
"Write Buffer 2");
check_err(clEnqueueNDRangeKernel(queue, adder, 1, 0, &work_size, 0, 0, 0, &event),
"Can't enqueue kernel");
check_err(clEnqueueReadBuffer(
queue, cl_res, CL_TRUE, 0, sizeof(cl_float) * DATA_SIZE, res, 0, 0, 0),
"Can't enqueue read buffer");
clWaitForEvents(1, &event);
printf("Execution Time: %.04lf ms\n\n", get_event_exec_time(event));
// Make sure everything is done before we do anything
clFinish(queue);
err = 0;
for (i = 0; i < DATA_SIZE; i++) {
if (res[i] != a[i] + b[i]) {
printf("%f + %f = %f(answer %f)\n", a[i], b[i], res[i], a[i] + b[i]);
err++;
}
}
if (err == 0)
printf("Validation passed\n");
else
printf("Validation failed\n");
release:
clReleaseKernel(adder);
clReleaseProgram(program);
clReleaseMemObject(cl_a);
clReleaseMemObject(cl_b);
clReleaseMemObject(cl_res);
clReleaseCommandQueue(queue);
clReleaseContext(context);
return 0;
}
XCamReturn
CLPostImageProcessor::create_handlers ()
{
SmartPtr<CLImageHandler> image_handler;
SmartPtr<CLContext> context = get_cl_context ();
XCAM_ASSERT (context.ptr ());
/* defog: retinex */
image_handler = create_cl_retinex_image_handler (context);
_retinex = image_handler.dynamic_cast_ptr<CLRetinexImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
_retinex.ptr (),
XCAM_RETURN_ERROR_CL,
"CLPostImageProcessor create retinex handler failed");
_retinex->set_kernels_enable (_defog_mode == CLPostImageProcessor::DefogRetinex);
image_handler->set_pool_type (CLImageHandler::DrmBoPoolType);
image_handler->set_pool_size (XCAM_CL_POST_IMAGE_MAX_POOL_SIZE);
add_handler (image_handler);
/* defog: dark channel prior */
image_handler = create_cl_defog_dcp_image_handler (context);
_defog_dcp = image_handler.dynamic_cast_ptr<CLDefogDcpImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
_retinex.ptr (),
XCAM_RETURN_ERROR_CL,
"CLPostImageProcessor create retinex handler failed");
_defog_dcp->set_kernels_enable (_defog_mode == CLPostImageProcessor::DefogDarkChannelPrior);
image_handler->set_pool_type (CLImageHandler::DrmBoPoolType);
image_handler->set_pool_size (XCAM_CL_POST_IMAGE_MAX_POOL_SIZE);
add_handler (image_handler);
/* Temporal Noise Reduction */
if (_defog_mode != CLPostImageProcessor::DefogDisabled) {
switch (_tnr_mode) {
case TnrYuv: {
image_handler = create_cl_tnr_image_handler (context, CL_TNR_TYPE_YUV);
_tnr = image_handler.dynamic_cast_ptr<CLTnrImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
_tnr.ptr (),
XCAM_RETURN_ERROR_CL,
"CLPostImageProcessor create tnr handler failed");
_tnr->set_mode (CL_TNR_TYPE_YUV);
image_handler->set_pool_type (CLImageHandler::DrmBoPoolType);
image_handler->set_pool_size (XCAM_CL_POST_IMAGE_DEFAULT_POOL_SIZE);
add_handler (image_handler);
break;
}
case TnrDisable:
XCAM_LOG_DEBUG ("CLPostImageProcessor disable tnr");
break;
default:
XCAM_LOG_WARNING ("CLPostImageProcessor unknown tnr mode (%d)", _tnr_mode);
break;
}
}
/* csc (nv12torgba) */
image_handler = create_cl_csc_image_handler (context, CL_CSC_TYPE_NV12TORGBA);
_csc = image_handler.dynamic_cast_ptr<CLCscImageHandler> ();
XCAM_FAIL_RETURN (
WARNING,
_csc .ptr (),
XCAM_RETURN_ERROR_CL,
"CLPostImageProcessor create csc handler failed");
_csc->set_kernels_enable (_out_sample_type == OutSampleRGB);
_csc->set_output_format (_output_fourcc);
image_handler->set_pool_type (CLImageHandler::DrmBoPoolType);
image_handler->set_pool_size (XCAM_CL_POST_IMAGE_DEFAULT_POOL_SIZE);
add_handler (image_handler);
return XCAM_RETURN_NO_ERROR;
}