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