generic-library/opencv
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

ocl: file-based ProgramCache refactoring

Browse Source
This commit is contained in:
Alexander Alekhin 2013-09-25 19:07:14 +04:00
parent b00f79ac5f
commit dd9ff587dc
15 changed files with 529 additions and 287 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3
cmake/OpenCVModule.cmake
View File

@ -445,6 +445,8 @@ macro(ocv_glob_module_sources)
source_group("Src\\Cuda" FILES ${lib_cuda_srcs} ${lib_cuda_hdrs}) source_group("Src\\Cuda" FILES ${lib_cuda_srcs} ${lib_cuda_hdrs})
endif() endif()
source_group("Src" FILES ${lib_srcs} ${lib_int_hdrs})
file(GLOB cl_kernels "src/opencl/*.cl") file(GLOB cl_kernels "src/opencl/*.cl")
if(HAVE_OPENCL AND cl_kernels) if(HAVE_OPENCL AND cl_kernels)
@ -457,7 +459,6 @@ macro(ocv_glob_module_sources)
list(APPEND lib_srcs ${cl_kernels} "${CMAKE_CURRENT_BINARY_DIR}/opencl_kernels.cpp" "${CMAKE_CURRENT_BINARY_DIR}/opencl_kernels.hpp") list(APPEND lib_srcs ${cl_kernels} "${CMAKE_CURRENT_BINARY_DIR}/opencl_kernels.cpp" "${CMAKE_CURRENT_BINARY_DIR}/opencl_kernels.hpp")
endif() endif()
source_group("Src" FILES ${lib_srcs} ${lib_int_hdrs})
source_group("Include" FILES ${lib_hdrs}) source_group("Include" FILES ${lib_hdrs})
source_group("Include\\detail" FILES ${lib_hdrs_detail}) source_group("Include\\detail" FILES ${lib_hdrs_detail})

19
cmake/cl2cpp.cmake
View File

@ -20,6 +20,7 @@ namespace cv
{ {
namespace ocl namespace ocl
{ {
") ")
foreach(cl ${cl_list}) foreach(cl ${cl_list})
@ -43,12 +44,22 @@ foreach(cl ${cl_list})
string(REGEX REPLACE "\"$" "" lines "${lines}") # unneeded " at the eof string(REGEX REPLACE "\"$" "" lines "${lines}") # unneeded " at the eof
set(STR_CPP "${STR_CPP}const char* ${cl_filename}=\"${lines};\n") string(MD5 hash "${lines}")
set(STR_HPP "${STR_HPP}extern const char* ${cl_filename};\n")
set(STR_CPP "${STR_CPP}const struct ProgramEntry ${cl_filename}={\"${cl_filename}\",\n\"${lines}, \"${hash}\"};\n")
set(STR_HPP "${STR_HPP}extern const struct ProgramEntry ${cl_filename};\n")
endforeach() endforeach()
set(STR_CPP "${STR_CPP}}\n}\n") set(STR_CPP "${STR_CPP}}\n}\n")
set(STR_HPP "${STR_HPP}}\n}\n") set(STR_HPP "${STR_HPP}}\n}\n")
file(WRITE ${OUTPUT} "${STR_CPP}") file(WRITE "${OUTPUT}" "${STR_CPP}")
file(WRITE ${OUTPUT_HPP} "${STR_HPP}")
if(EXISTS "${OUTPUT_HPP}")
file(READ "${OUTPUT_HPP}" hpp_lines)
endif()
if("${hpp_lines}" STREQUAL "${STR_HPP}")
message(STATUS "${OUTPUT_HPP} contains same content")
else()
file(WRITE "${OUTPUT_HPP}" "${STR_HPP}")
endif()

4
modules/nonfree/src/surf.ocl.cpp
View File

@ -55,11 +55,11 @@ namespace cv
{ {
namespace ocl namespace ocl
{ {
const char noImage2dOption [] = "-D DISABLE_IMAGE2D"; static const char noImage2dOption[] = "-D DISABLE_IMAGE2D";
static bool use_image2d = false; static bool use_image2d = false;
static void openCLExecuteKernelSURF(Context *clCxt , const char **source, string kernelName, size_t globalThreads[3], static void openCLExecuteKernelSURF(Context *clCxt, const cv::ocl::ProgramEntry* source, string kernelName, size_t globalThreads[3],
size_t localThreads[3], std::vector< std::pair<size_t, const void *> > &args, int channels, int depth) size_t localThreads[3], std::vector< std::pair<size_t, const void *> > &args, int channels, int depth)
{ {
char optBuf [100] = {0}; char optBuf [100] = {0};

21
modules/ocl/include/opencv2/ocl/ocl.hpp
View File

@ -199,24 +199,6 @@ namespace cv
void CV_EXPORTS finish(); void CV_EXPORTS finish();
//! Calls a kernel, by string. Pass globalThreads = NULL, and cleanUp = true, to finally clean-up without executing.
CV_EXPORTS double openCLExecuteKernelInterop(Context *clCxt ,
const char **source, string kernelName,
size_t globalThreads[3], size_t localThreads[3],
std::vector< std::pair<size_t, const void *> > &args,
int channels, int depth, const char *build_options,
bool finish = true, bool measureKernelTime = false,
bool cleanUp = true);
//! Calls a kernel, by file. Pass globalThreads = NULL, and cleanUp = true, to finally clean-up without executing.
CV_EXPORTS double openCLExecuteKernelInterop(Context *clCxt ,
const char **fileName, const int numFiles, string kernelName,
size_t globalThreads[3], size_t localThreads[3],
std::vector< std::pair<size_t, const void *> > &args,
int channels, int depth, const char *build_options,
bool finish = true, bool measureKernelTime = false,
bool cleanUp = true);
//! Enable or disable OpenCL program binary caching onto local disk //! Enable or disable OpenCL program binary caching onto local disk
// After a program (*.cl files in opencl/ folder) is built at runtime, we allow the // After a program (*.cl files in opencl/ folder) is built at runtime, we allow the
// compiled OpenCL program to be cached to the path automatically as "path/*.clb" // compiled OpenCL program to be cached to the path automatically as "path/*.clb"
@ -233,12 +215,11 @@ namespace cv
CACHE_DEBUG = 0x1 << 0, // cache OpenCL binary when built in debug mode (only work with MSVC) CACHE_DEBUG = 0x1 << 0, // cache OpenCL binary when built in debug mode (only work with MSVC)
CACHE_RELEASE = 0x1 << 1, // default behavior, only cache when built in release mode (only work with MSVC) CACHE_RELEASE = 0x1 << 1, // default behavior, only cache when built in release mode (only work with MSVC)
CACHE_ALL = CACHE_DEBUG | CACHE_RELEASE, // always cache opencl binary CACHE_ALL = CACHE_DEBUG | CACHE_RELEASE, // always cache opencl binary
CACHE_UPDATE = 0x1 << 2 // if the binary cache file with the same name is already on the disk, it will be updated.
}; };
CV_EXPORTS void setBinaryDiskCache(int mode = CACHE_RELEASE, cv::String path = "./"); CV_EXPORTS void setBinaryDiskCache(int mode = CACHE_RELEASE, cv::String path = "./");
//! set where binary cache to be saved to //! set where binary cache to be saved to
CV_EXPORTS void setBinpath(const char *path); CV_EXPORTS void setBinaryPath(const char *path);
class CV_EXPORTS oclMatExpr; class CV_EXPORTS oclMatExpr;
//////////////////////////////// oclMat //////////////////////////////// //////////////////////////////// oclMat ////////////////////////////////

76
modules/ocl/include/opencv2/ocl/private/util.hpp
View File

@ -55,6 +55,13 @@ namespace cv
namespace ocl namespace ocl
{ {
struct ProgramEntry
{
const char* name;
const char* programStr;
const char* programHash;
};
inline cl_device_id getClDeviceID(const Context *ctx) inline cl_device_id getClDeviceID(const Context *ctx)
{ {
return *(cl_device_id*)(ctx->getOpenCLDeviceIDPtr()); return *(cl_device_id*)(ctx->getOpenCLDeviceIDPtr());
@ -78,41 +85,39 @@ enum openCLMemcpyKind
}; };
///////////////////////////OpenCL call wrappers//////////////////////////// ///////////////////////////OpenCL call wrappers////////////////////////////
void CV_EXPORTS openCLMallocPitch(Context *clCxt, void **dev_ptr, size_t *pitch, void CV_EXPORTS openCLMallocPitch(Context *clCxt, void **dev_ptr, size_t *pitch,
size_t widthInBytes, size_t height); size_t widthInBytes, size_t height);
void CV_EXPORTS openCLMallocPitchEx(Context *clCxt, void **dev_ptr, size_t *pitch, void CV_EXPORTS openCLMallocPitchEx(Context *clCxt, void **dev_ptr, size_t *pitch,
size_t widthInBytes, size_t height, DevMemRW rw_type, DevMemType mem_type); size_t widthInBytes, size_t height, DevMemRW rw_type, DevMemType mem_type);
void CV_EXPORTS openCLMemcpy2D(Context *clCxt, void *dst, size_t dpitch, void CV_EXPORTS openCLMemcpy2D(Context *clCxt, void *dst, size_t dpitch,
const void *src, size_t spitch, const void *src, size_t spitch,
size_t width, size_t height, openCLMemcpyKind kind, int channels = -1); size_t width, size_t height, openCLMemcpyKind kind, int channels = -1);
void CV_EXPORTS openCLCopyBuffer2D(Context *clCxt, void *dst, size_t dpitch, int dst_offset, void CV_EXPORTS openCLCopyBuffer2D(Context *clCxt, void *dst, size_t dpitch, int dst_offset,
const void *src, size_t spitch, const void *src, size_t spitch,
size_t width, size_t height, int src_offset); size_t width, size_t height, int src_offset);
void CV_EXPORTS openCLFree(void *devPtr); void CV_EXPORTS openCLFree(void *devPtr);
cl_mem CV_EXPORTS openCLCreateBuffer(Context *clCxt, size_t flag, size_t size); cl_mem CV_EXPORTS openCLCreateBuffer(Context *clCxt, size_t flag, size_t size);
void CV_EXPORTS openCLReadBuffer(Context *clCxt, cl_mem dst_buffer, void *host_buffer, size_t size); void CV_EXPORTS openCLReadBuffer(Context *clCxt, cl_mem dst_buffer, void *host_buffer, size_t size);
cl_kernel CV_EXPORTS openCLGetKernelFromSource(const Context *clCxt, cl_kernel CV_EXPORTS openCLGetKernelFromSource(const Context *clCxt,
const char **source, std::string kernelName); const cv::ocl::ProgramEntry* source, std::string kernelName);
cl_kernel CV_EXPORTS openCLGetKernelFromSource(const Context *clCxt, cl_kernel CV_EXPORTS openCLGetKernelFromSource(const Context *clCxt,
const char **source, std::string kernelName, const char *build_options); const cv::ocl::ProgramEntry* source, std::string kernelName, const char *build_options);
void CV_EXPORTS openCLVerifyKernel(const Context *clCxt, cl_kernel kernel, size_t *localThreads); void CV_EXPORTS openCLVerifyKernel(const Context *clCxt, cl_kernel kernel, size_t *localThreads);
void CV_EXPORTS openCLExecuteKernel(Context *clCxt , const char **source, string kernelName, std::vector< std::pair<size_t, const void *> > &args, void CV_EXPORTS openCLExecuteKernel(Context *clCxt , const cv::ocl::ProgramEntry* source, string kernelName, std::vector< std::pair<size_t, const void *> > &args,
int globalcols , int globalrows, size_t blockSize = 16, int kernel_expand_depth = -1, int kernel_expand_channel = -1); int globalcols , int globalrows, size_t blockSize = 16, int kernel_expand_depth = -1, int kernel_expand_channel = -1);
void CV_EXPORTS openCLExecuteKernel_(Context *clCxt , const char **source, std::string kernelName, void CV_EXPORTS openCLExecuteKernel_(Context *clCxt, const cv::ocl::ProgramEntry* source, std::string kernelName,
size_t globalThreads[3], size_t localThreads[3], size_t globalThreads[3], size_t localThreads[3],
std::vector< std::pair<size_t, const void *> > &args, int channels, int depth, const char *build_options); std::vector< std::pair<size_t, const void *> > &args, int channels, int depth, const char *build_options);
void CV_EXPORTS openCLExecuteKernel(Context *clCxt , const char **source, std::string kernelName, size_t globalThreads[3], void CV_EXPORTS openCLExecuteKernel(Context *clCxt, const cv::ocl::ProgramEntry* source, std::string kernelName, size_t globalThreads[3],
size_t localThreads[3], std::vector< std::pair<size_t, const void *> > &args, int channels, int depth); size_t localThreads[3], std::vector< std::pair<size_t, const void *> > &args, int channels, int depth);
void CV_EXPORTS openCLExecuteKernel(Context *clCxt , const char **source, std::string kernelName, size_t globalThreads[3], void CV_EXPORTS openCLExecuteKernel(Context *clCxt, const cv::ocl::ProgramEntry* source, std::string kernelName, size_t globalThreads[3],
size_t localThreads[3], std::vector< std::pair<size_t, const void *> > &args, int channels, size_t localThreads[3], std::vector< std::pair<size_t, const void *> > &args, int channels,
int depth, const char *build_options); int depth, const char *build_options);
cl_mem CV_EXPORTS load_constant(cl_context context, cl_command_queue command_queue, const void *value, cl_mem CV_EXPORTS load_constant(cl_context context, cl_command_queue command_queue, const void *value,
const size_t size); const size_t size);
cl_mem CV_EXPORTS openCLMalloc(cl_context clCxt, size_t size, cl_mem_flags flags, void *host_ptr); cl_mem CV_EXPORTS openCLMalloc(cl_context clCxt, size_t size, cl_mem_flags flags, void *host_ptr);
int CV_EXPORTS savetofile(const Context *clcxt, cl_program &program, const char *fileName);
enum FLUSH_MODE enum FLUSH_MODE
{ {
CLFINISH = 0, CLFINISH = 0,
@ -120,11 +125,12 @@ enum FLUSH_MODE
DISABLE DISABLE
}; };
void CV_EXPORTS openCLExecuteKernel2(Context *clCxt , const char **source, std::string kernelName, size_t globalThreads[3], void CV_EXPORTS openCLExecuteKernel2(Context *clCxt, const cv::ocl::ProgramEntry* source, std::string kernelName, size_t globalThreads[3],
size_t localThreads[3], std::vector< std::pair<size_t, const void *> > &args, int channels, int depth, FLUSH_MODE finish_mode = DISABLE); size_t localThreads[3], std::vector< std::pair<size_t, const void *> > &args, int channels, int depth, FLUSH_MODE finish_mode = DISABLE);
void CV_EXPORTS openCLExecuteKernel2(Context *clCxt , const char **source, std::string kernelName, size_t globalThreads[3], void CV_EXPORTS openCLExecuteKernel2(Context *clCxt, const cv::ocl::ProgramEntry* source, std::string kernelName, size_t globalThreads[3],
size_t localThreads[3], std::vector< std::pair<size_t, const void *> > &args, int channels, size_t localThreads[3], std::vector< std::pair<size_t, const void *> > &args, int channels,
int depth, char *build_options, FLUSH_MODE finish_mode = DISABLE); int depth, char *build_options, FLUSH_MODE finish_mode = DISABLE);
// bind oclMat to OpenCL image textures // bind oclMat to OpenCL image textures
// note: // note:
// 1. there is no memory management. User need to explicitly release the resource // 1. there is no memory management. User need to explicitly release the resource
@ -183,6 +189,24 @@ inline size_t roundUp(size_t sz, size_t n)
return result; return result;
} }
//! Calls a kernel, by string. Pass globalThreads = NULL, and cleanUp = true, to finally clean-up without executing.
CV_EXPORTS double openCLExecuteKernelInterop(Context *clCxt,
const cv::ocl::ProgramEntry* source, string kernelName,
size_t globalThreads[3], size_t localThreads[3],
std::vector< std::pair<size_t, const void *> > &args,
int channels, int depth, const char *build_options,
bool finish = true, bool measureKernelTime = false,
bool cleanUp = true);
//! Calls a kernel, by file. Pass globalThreads = NULL, and cleanUp = true, to finally clean-up without executing.
CV_EXPORTS double openCLExecuteKernelInterop(Context *clCxt,
const cv::ocl::ProgramEntry* source, const int numFiles, string kernelName,
size_t globalThreads[3], size_t localThreads[3],
std::vector< std::pair<size_t, const void *> > &args,
int channels, int depth, const char *build_options,
bool finish = true, bool measureKernelTime = false,
bool cleanUp = true);
}//namespace ocl }//namespace ocl
}//namespace cv }//namespace cv

1
modules/ocl/perf/main.cpp
View File

@ -91,7 +91,6 @@ int main(int argc, char ** argv)
} }
cv::ocl::setDevice(devicesInfo[device]); cv::ocl::setDevice(devicesInfo[device]);
cv::ocl::setBinaryDiskCache(cv::ocl::CACHE_UPDATE);
cout << "Device type:" << type << endl cout << "Device type:" << type << endl
<< "Platform name:" << devicesInfo[device]->platform->platformName << endl << "Platform name:" << devicesInfo[device]->platform->platformName << endl

4
modules/ocl/src/brute_force_matcher.cpp
View File

@ -45,10 +45,14 @@
//M*/ //M*/
#include "precomp.hpp" #include "precomp.hpp"
#include <functional>
#include <iterator>
#include <vector>
#include "opencl_kernels.hpp" #include "opencl_kernels.hpp"
using namespace cv; using namespace cv;
using namespace cv::ocl; using namespace cv::ocl;
using namespace std;
static const int OPT_SIZE = 100; static const int OPT_SIZE = 100;

57
modules/ocl/src/cl_context.cpp
View File

@ -48,15 +48,16 @@
#include "precomp.hpp" #include "precomp.hpp"
#include <iomanip> #include <iomanip>
#include <fstream> #include <fstream>
#include "binarycaching.hpp" #include "cl_programcache.hpp"
#if defined _MSC_VER && _MSC_VER >= 1200
# pragma warning( disable: 4100 4101 4127 4244 4267 4510 4512 4610)
#endif
#undef __CL_ENABLE_EXCEPTIONS #undef __CL_ENABLE_EXCEPTIONS
#include <CL/cl.hpp> #include <CL/cl.hpp>
namespace cv { namespace ocl { namespace cv {
namespace ocl {
extern void fft_teardown();
extern void clBlasTeardown();
struct PlatformInfoImpl struct PlatformInfoImpl
{ {
@ -174,7 +175,7 @@ static int initializeOpenCLDevices()
deviceInfo.info.platform = &platformInfo.info; deviceInfo.info.platform = &platformInfo.info;
platformInfo.deviceIDs[j] = deviceInfo.info._id; platformInfo.deviceIDs[j] = deviceInfo.info._id;
cl_device_type type = -1; cl_device_type type = cl_device_type(-1);
openCLSafeCall(device.getInfo(CL_DEVICE_TYPE, &type)); openCLSafeCall(device.getInfo(CL_DEVICE_TYPE, &type));
deviceInfo.info.deviceType = DeviceType(type); deviceInfo.info.deviceType = DeviceType(type);
@ -182,7 +183,7 @@ static int initializeOpenCLDevices()
openCLSafeCall(device.getInfo(CL_DEVICE_VERSION, &deviceInfo.info.deviceVersion)); openCLSafeCall(device.getInfo(CL_DEVICE_VERSION, &deviceInfo.info.deviceVersion));
openCLSafeCall(device.getInfo(CL_DEVICE_NAME, &deviceInfo.info.deviceName)); openCLSafeCall(device.getInfo(CL_DEVICE_NAME, &deviceInfo.info.deviceName));
openCLSafeCall(device.getInfo(CL_DEVICE_VENDOR, &deviceInfo.info.deviceVendor)); openCLSafeCall(device.getInfo(CL_DEVICE_VENDOR, &deviceInfo.info.deviceVendor));
cl_uint vendorID = -1; cl_uint vendorID = 0;
openCLSafeCall(device.getInfo(CL_DEVICE_VENDOR_ID, &vendorID)); openCLSafeCall(device.getInfo(CL_DEVICE_VENDOR_ID, &vendorID));
deviceInfo.info.deviceVendorId = vendorID; deviceInfo.info.deviceVendorId = vendorID;
openCLSafeCall(device.getInfo(CL_DRIVER_VERSION, &deviceInfo.info.deviceDriverVersion)); openCLSafeCall(device.getInfo(CL_DRIVER_VERSION, &deviceInfo.info.deviceDriverVersion));
@ -347,9 +348,6 @@ static bool __termination = false;
ContextImpl::~ContextImpl() ContextImpl::~ContextImpl()
{ {
fft_teardown();
clBlasTeardown();
#ifdef WIN32 #ifdef WIN32
// if process is on termination stage (ExitProcess was called and other threads were terminated) // if process is on termination stage (ExitProcess was called and other threads were terminated)
// then disable command queue release because it may cause program hang // then disable command queue release because it may cause program hang
@ -370,8 +368,14 @@ ContextImpl::~ContextImpl()
clContext = NULL; clContext = NULL;
} }
void fft_teardown();
void clBlasTeardown();
void ContextImpl::cleanupContext(void) void ContextImpl::cleanupContext(void)
{ {
fft_teardown();
clBlasTeardown();
cv::AutoLock lock(currentContextMutex); cv::AutoLock lock(currentContextMutex);
if (currentContext) if (currentContext)
delete currentContext; delete currentContext;
@ -382,6 +386,15 @@ void ContextImpl::setContext(const DeviceInfo* deviceInfo)
{ {
CV_Assert(deviceInfo->_id >= 0 && deviceInfo->_id < (int)global_devices.size()); CV_Assert(deviceInfo->_id >= 0 && deviceInfo->_id < (int)global_devices.size());
{
cv::AutoLock lock(currentContextMutex);
if (currentContext)
{
if (currentContext->deviceInfo._id == deviceInfo->_id)
return;
}
}
DeviceInfoImpl& infoImpl = global_devices[deviceInfo->_id]; DeviceInfoImpl& infoImpl = global_devices[deviceInfo->_id];
CV_Assert(deviceInfo == &infoImpl.info); CV_Assert(deviceInfo == &infoImpl.info);
@ -466,6 +479,30 @@ int getOpenCLDevices(std::vector<const DeviceInfo*> &devices, int deviceType, co
} }
} }
if (currentContext == NULL)
{
// select default device
const DeviceInfo* selectedDevice = NULL;
for (size_t i = 0; i < devices.size(); i++)
{
const DeviceInfo* dev = devices[i];
if (dev->deviceType == CL_DEVICE_TYPE_GPU)
{
selectedDevice = dev;
break;
}
else if (dev->deviceType == CL_DEVICE_TYPE_CPU && (selectedDevice == NULL))
{
selectedDevice = dev;
}
}
if (selectedDevice)
{
setDevice(selectedDevice);
}
}
return (int)devices.size(); return (int)devices.size();
} }

47
modules/ocl/src/cl_operations.cpp
View File

@ -48,10 +48,7 @@
#include "precomp.hpp" #include "precomp.hpp"
#include <iomanip> #include <iomanip>
#include <fstream> #include <fstream>
#include "binarycaching.hpp" #include "cl_programcache.hpp"
#undef __CL_ENABLE_EXCEPTIONS
#include <CL/cl.hpp>
//#define PRINT_KERNEL_RUN_TIME //#define PRINT_KERNEL_RUN_TIME
#define RUN_TIMES 100 #define RUN_TIMES 100
@ -60,7 +57,8 @@
#endif #endif
//#define AMD_DOUBLE_DIFFER //#define AMD_DOUBLE_DIFFER
namespace cv { namespace ocl { namespace cv {
namespace ocl {
DevMemType gDeviceMemType = DEVICE_MEM_DEFAULT; DevMemType gDeviceMemType = DEVICE_MEM_DEFAULT;
DevMemRW gDeviceMemRW = DEVICE_MEM_R_W; DevMemRW gDeviceMemRW = DEVICE_MEM_R_W;
@ -179,21 +177,22 @@ void openCLFree(void *devPtr)
openCLSafeCall(clReleaseMemObject((cl_mem)devPtr)); openCLSafeCall(clReleaseMemObject((cl_mem)devPtr));
} }
cl_kernel openCLGetKernelFromSource(const Context *ctx, const char **source, string kernelName) cl_kernel openCLGetKernelFromSource(const Context *ctx, const cv::ocl::ProgramEntry* source, string kernelName)
{ {
return openCLGetKernelFromSource(ctx, source, kernelName, NULL); return openCLGetKernelFromSource(ctx, source, kernelName, NULL);
} }
cl_kernel openCLGetKernelFromSource(const Context *ctx, const char **source, string kernelName, cl_kernel openCLGetKernelFromSource(const Context *ctx, const cv::ocl::ProgramEntry* source, string kernelName,
const char *build_options) const char *build_options)
{ {
cl_kernel kernel; cl_kernel kernel;
cl_int status = 0; cl_int status = 0;
CV_Assert(ProgramCache::getProgramCache() != NULL); CV_Assert(ProgramCache::getProgramCache() != NULL);
cl_program program = ProgramCache::getProgramCache()->getProgram(ctx, source, kernelName, build_options); cl_program program = ProgramCache::getProgramCache()->getProgram(ctx, source, build_options);
CV_Assert(program != NULL); CV_Assert(program != NULL);
kernel = clCreateKernel(program, kernelName.c_str(), &status); kernel = clCreateKernel(program, kernelName.c_str(), &status);
openCLVerifyCall(status); openCLVerifyCall(status);
openCLVerifyCall(clReleaseProgram(program));
return kernel; return kernel;
} }
@ -213,7 +212,7 @@ void openCLVerifyKernel(const Context *ctx, cl_kernel kernel, size_t *localThrea
static double total_execute_time = 0; static double total_execute_time = 0;
static double total_kernel_time = 0; static double total_kernel_time = 0;
#endif #endif
void openCLExecuteKernel_(Context *ctx , const char **source, string kernelName, size_t globalThreads[3], void openCLExecuteKernel_(Context *ctx, const cv::ocl::ProgramEntry* source, string kernelName, size_t globalThreads[3],
size_t localThreads[3], vector< pair<size_t, const void *> > &args, int channels, size_t localThreads[3], vector< pair<size_t, const void *> > &args, int channels,
int depth, const char *build_options) int depth, const char *build_options)
{ {
@ -275,14 +274,14 @@ void openCLExecuteKernel_(Context *ctx , const char **source, string kernelName,
openCLSafeCall(clReleaseKernel(kernel)); openCLSafeCall(clReleaseKernel(kernel));
} }
void openCLExecuteKernel(Context *ctx , const char **source, string kernelName, void openCLExecuteKernel(Context *ctx, const cv::ocl::ProgramEntry* source, string kernelName,
size_t globalThreads[3], size_t localThreads[3], size_t globalThreads[3], size_t localThreads[3],
vector< pair<size_t, const void *> > &args, int channels, int depth) vector< pair<size_t, const void *> > &args, int channels, int depth)
{ {
openCLExecuteKernel(ctx, source, kernelName, globalThreads, localThreads, args, openCLExecuteKernel(ctx, source, kernelName, globalThreads, localThreads, args,
channels, depth, NULL); channels, depth, NULL);
} }
void openCLExecuteKernel(Context *ctx , const char **source, string kernelName, void openCLExecuteKernel(Context *ctx, const cv::ocl::ProgramEntry* source, string kernelName,
size_t globalThreads[3], size_t localThreads[3], size_t globalThreads[3], size_t localThreads[3],
vector< pair<size_t, const void *> > &args, int channels, int depth, const char *build_options) vector< pair<size_t, const void *> > &args, int channels, int depth, const char *build_options)
@ -316,7 +315,7 @@ void openCLExecuteKernel(Context *ctx , const char **source, string kernelName,
#endif #endif
} }
double openCLExecuteKernelInterop(Context *ctx , const char **source, string kernelName, double openCLExecuteKernelInterop(Context *ctx, const cv::ocl::ProgramEntry* source, string kernelName,
size_t globalThreads[3], size_t localThreads[3], size_t globalThreads[3], size_t localThreads[3],
vector< pair<size_t, const void *> > &args, int channels, int depth, const char *build_options, vector< pair<size_t, const void *> > &args, int channels, int depth, const char *build_options,
bool finish, bool measureKernelTime, bool cleanUp) bool finish, bool measureKernelTime, bool cleanUp)
@ -391,29 +390,6 @@ double openCLExecuteKernelInterop(Context *ctx , const char **source, string ker
return kernelTime; return kernelTime;
} }
//double openCLExecuteKernelInterop(Context *ctx , const char **fileName, const int numFiles, string kernelName,
// size_t globalThreads[3], size_t localThreads[3],
// vector< pair<size_t, const void *> > &args, int channels, int depth, const char *build_options,
// bool finish, bool measureKernelTime, bool cleanUp)
//
//{
// std::vector<std::string> fsource;
// for (int i = 0 ; i < numFiles ; i++)
// {
// std::string str;
// if (convertToString(fileName[i], str) >= 0)
// fsource.push_back(str);
// }
// const char **source = new const char *[numFiles];
// for (int i = 0 ; i < numFiles ; i++)
// source[i] = fsource[i].c_str();
// double kernelTime = openCLExecuteKernelInterop(ctx ,source, kernelName, globalThreads, localThreads,
// args, channels, depth, build_options, finish, measureKernelTime, cleanUp);
// fsource.clear();
// delete []source;
// return kernelTime;
//}
cl_mem load_constant(cl_context context, cl_command_queue command_queue, const void *value, cl_mem load_constant(cl_context context, cl_command_queue command_queue, const void *value,
const size_t size) const size_t size)
{ {
@ -427,7 +403,6 @@ cl_mem load_constant(cl_context context, cl_command_queue command_queue, const v
value, 0, 0, 0)); value, 0, 0, 0));
return con_struct; return con_struct;
} }
}//namespace ocl }//namespace ocl

512
modules/ocl/src/cl_programcache.cpp
View File

@ -48,85 +48,93 @@
#include "precomp.hpp" #include "precomp.hpp"
#include <iomanip> #include <iomanip>
#include <fstream> #include <fstream>
#include "binarycaching.hpp" #include "cl_programcache.hpp"
#if defined _MSC_VER && _MSC_VER >= 1200
# pragma warning( disable: 4100 4244 4267 4510 4512 4610)
#endif
#undef __CL_ENABLE_EXCEPTIONS #undef __CL_ENABLE_EXCEPTIONS
#include <CL/cl.hpp> #include <CL/cl.hpp>
namespace cv { namespace ocl { namespace cv { namespace ocl {
#define MAX_PROG_CACHE_SIZE 1024
/* /*
* The binary caching system to eliminate redundant program source compilation. * The binary caching system to eliminate redundant program source compilation.
* Strictly, this is not a cache because we do not implement evictions right now. * Strictly, this is not a cache because we do not implement evictions right now.
* We shall add such features to trade-off memory consumption and performance when necessary. * We shall add such features to trade-off memory consumption and performance when necessary.
*/ */
cv::Mutex ProgramCache::mutexFiles;
cv::Mutex ProgramCache::mutexCache;
std::auto_ptr<ProgramCache> _programCache; std::auto_ptr<ProgramCache> _programCache;
ProgramCache* ProgramCache::getProgramCache() ProgramCache* ProgramCache::getProgramCache()
{ {
if (NULL == _programCache.get()) if (NULL == _programCache.get())
_programCache.reset(new ProgramCache()); _programCache.reset(new ProgramCache());
return _programCache.get(); return _programCache.get();
} }
ProgramCache::ProgramCache() ProgramCache::ProgramCache()
{ {
codeCache.clear(); codeCache.clear();
cacheSize = 0; cacheSize = 0;
} }
ProgramCache::~ProgramCache() ProgramCache::~ProgramCache()
{ {
releaseProgram(); releaseProgram();
} }
cl_program ProgramCache::progLookup(string srcsign) cl_program ProgramCache::progLookup(const string& srcsign)
{ {
map<string, cl_program>::iterator iter; map<string, cl_program>::iterator iter;
iter = codeCache.find(srcsign); iter = codeCache.find(srcsign);
if(iter != codeCache.end()) if(iter != codeCache.end())
return iter->second; return iter->second;
else else
return NULL; return NULL;
} }
void ProgramCache::addProgram(string srcsign , cl_program program) void ProgramCache::addProgram(const string& srcsign, cl_program program)
{ {
if(!progLookup(srcsign)) if (!progLookup(srcsign))
{ {
codeCache.insert(map<string, cl_program>::value_type(srcsign, program)); clRetainProgram(program);
} codeCache.insert(map<string, cl_program>::value_type(srcsign, program));
}
} }
void ProgramCache::releaseProgram() void ProgramCache::releaseProgram()
{ {
map<string, cl_program>::iterator iter; map<string, cl_program>::iterator iter;
for(iter = codeCache.begin(); iter != codeCache.end(); iter++) for(iter = codeCache.begin(); iter != codeCache.end(); iter++)
{ {
openCLSafeCall(clReleaseProgram(iter->second)); openCLSafeCall(clReleaseProgram(iter->second));
} }
codeCache.clear(); codeCache.clear();
cacheSize = 0; cacheSize = 0;
} }
static int enable_disk_cache = static int enable_disk_cache = true ||
#ifdef _DEBUG #ifdef _DEBUG
false; false;
#else #else
true; true;
#endif #endif
static int update_disk_cache = false;
static String binpath = ""; static String binpath = "";
void setBinaryDiskCache(int mode, String path) void setBinaryDiskCache(int mode, String path)
{ {
enable_disk_cache = 0;
binpath = "";
if(mode == CACHE_NONE) if(mode == CACHE_NONE)
{ {
update_disk_cache = 0;
enable_disk_cache = 0;
return; return;
} }
update_disk_cache |= (mode & CACHE_UPDATE) == CACHE_UPDATE; enable_disk_cache =
enable_disk_cache |=
#ifdef _DEBUG #ifdef _DEBUG
(mode & CACHE_DEBUG) == CACHE_DEBUG; (mode & CACHE_DEBUG) == CACHE_DEBUG;
#else #else
@ -138,108 +146,286 @@ void setBinaryDiskCache(int mode, String path)
} }
} }
void setBinpath(const char *path) void setBinaryPath(const char *path)
{ {
binpath = path; binpath = path;
} }
int savetofile(const Context*, cl_program &program, const char *fileName) static const int MAX_ENTRIES = 64;
struct ProgramFileCache
{ {
size_t binarySize; struct CV_DECL_ALIGNED(1) ProgramFileHeader
openCLSafeCall(clGetProgramInfo(program,
CL_PROGRAM_BINARY_SIZES,
sizeof(size_t),
&binarySize, NULL));
char* binary = (char*)malloc(binarySize);
if(binary == NULL)
{ {
CV_Error(CV_StsNoMem, "Failed to allocate host memory."); int hashLength;
} //char hash[];
openCLSafeCall(clGetProgramInfo(program, };
CL_PROGRAM_BINARIES,
sizeof(char *),
&binary,
NULL));
FILE *fp = fopen(fileName, "wb+"); struct CV_DECL_ALIGNED(1) ProgramFileTable
if(fp != NULL)
{ {
fwrite(binary, binarySize, 1, fp); int numberOfEntries;
free(binary); //int firstEntryOffset[];
fclose(fp); };
}
return 1;
}
cl_program ProgramCache::getProgram(const Context *ctx, const char **source, string kernelName, struct CV_DECL_ALIGNED(1) ProgramFileConfigurationEntry
const char *build_options)
{
cl_program program;
cl_int status = 0;
stringstream src_sign;
string srcsign;
string filename;
if (NULL != build_options)
{ {
src_sign << (int64)(*source) << getClContext(ctx) << "_" << build_options; int nextEntry;
} int dataSize;
else int optionsLength;
{ //char options[];
src_sign << (int64)(*source) << getClContext(ctx); // char data[];
} };
srcsign = src_sign.str();
program = NULL; string fileName_;
program = ProgramCache::getProgramCache()->progLookup(srcsign); const char* hash_;
std::fstream f;
if (!program) ProgramFileCache(const string& fileName, const char* hash)
: fileName_(fileName), hash_(hash)
{ {
//config build programs if (hash_ != NULL)
std::string all_build_options;
if (!ctx->getDeviceInfo().compilationExtraOptions.empty())
all_build_options += ctx->getDeviceInfo().compilationExtraOptions;
if (build_options != NULL)
{ {
all_build_options += " "; f.open(fileName_.c_str(), ios::in|ios::out|ios::binary);
all_build_options += build_options; if(f.is_open())
{
int hashLength = 0;
f.read((char*)&hashLength, sizeof(int));
std::vector<char> fhash(hashLength + 1);
f.read(&fhash[0], hashLength);
if (f.eof() || strncmp(hash_, &fhash[0], hashLength) != 0)
{
f.close();
remove(fileName_.c_str());
return;
}
}
} }
filename = binpath + kernelName + "_" + ctx->getDeviceInfo().deviceName + all_build_options + ".clb"; }
FILE *fp = enable_disk_cache ? fopen(filename.c_str(), "rb") : NULL; int getHash(const string& options)
if(fp == NULL || update_disk_cache) {
int hash = 0;
for (size_t i = 0; i < options.length(); i++)
{ {
if(fp != NULL) hash = (hash << 2) ^ (hash >> 17) ^ options[i];
fclose(fp); }
return (hash + (hash >> 16)) & (MAX_ENTRIES - 1);
}
program = clCreateProgramWithSource( bool readConfigurationFromFile(const string& options, std::vector<char>& buf)
getClContext(ctx), 1, source, NULL, &status); {
openCLVerifyCall(status); if (hash_ == NULL)
cl_device_id device = getClDeviceID(ctx); return false;
status = clBuildProgram(program, 1, &device, all_build_options.c_str(), NULL, NULL);
if(status == CL_SUCCESS && enable_disk_cache) if (!f.is_open())
savetofile(ctx, program, filename.c_str()); return false;
f.seekg(0, std::fstream::end);
size_t fileSize = (size_t)f.tellg();
if (fileSize == 0)
{
std::cerr << "Invalid file (empty): " << fileName_ << std::endl;
f.close();
remove(fileName_.c_str());
return false;
}
f.seekg(0, std::fstream::beg);
int hashLength = 0;
f.read((char*)&hashLength, sizeof(int));
CV_Assert(hashLength > 0);
f.seekg(sizeof(hashLength) + hashLength, std::fstream::beg);
int numberOfEntries = 0;
f.read((char*)&numberOfEntries, sizeof(int));
CV_Assert(numberOfEntries > 0);
if (numberOfEntries != MAX_ENTRIES)
{
std::cerr << "Invalid file: " << fileName_ << std::endl;
f.close();
remove(fileName_.c_str());
return false;
}
std::vector<int> firstEntryOffset(numberOfEntries);
f.read((char*)&firstEntryOffset[0], sizeof(int)*numberOfEntries);
int entryNum = getHash(options);
int entryOffset = firstEntryOffset[entryNum];
ProgramFileConfigurationEntry entry;
while (entryOffset > 0)
{
f.seekg(entryOffset, std::fstream::beg);
assert(sizeof(entry) == sizeof(int)*3);
f.read((char*)&entry, sizeof(entry));
std::vector<char> foptions(entry.optionsLength);
if ((int)options.length() == entry.optionsLength)
{
if (entry.optionsLength > 0)
f.read(&foptions[0], entry.optionsLength);
if (memcmp(&foptions[0], options.c_str(), entry.optionsLength) == 0)
{
buf.resize(entry.dataSize);
f.read(&buf[0], entry.dataSize);
f.seekg(0, std::fstream::beg);
return true;
}
}
if (entry.nextEntry <= 0)
break;
entryOffset = entry.nextEntry;
}
return false;
}
bool writeConfigurationToFile(const string& options, std::vector<char>& buf)
{
if (hash_ == NULL)
return true; // don't save dynamic kernels
if (!f.is_open())
{
f.open(fileName_.c_str(), ios::in|ios::out|ios::binary);
if (!f.is_open())
{
f.open(fileName_.c_str(), ios::out|ios::binary);
if (!f.is_open())
return false;
}
}
f.seekg(0, std::fstream::end);
size_t fileSize = (size_t)f.tellg();
if (fileSize == 0)
{
f.seekp(0, std::fstream::beg);
int hashLength = strlen(hash_);
f.write((char*)&hashLength, sizeof(int));
f.write(hash_, hashLength);
int numberOfEntries = MAX_ENTRIES;
f.write((char*)&numberOfEntries, sizeof(int));
std::vector<int> firstEntryOffset(MAX_ENTRIES, 0);
f.write((char*)&firstEntryOffset[0], sizeof(int)*numberOfEntries);
f.close();
f.open(fileName_.c_str(), ios::in|ios::out|ios::binary);
CV_Assert(f.is_open());
f.seekg(0, std::fstream::end);
fileSize = (size_t)f.tellg();
}
f.seekg(0, std::fstream::beg);
int hashLength = 0;
f.read((char*)&hashLength, sizeof(int));
CV_Assert(hashLength > 0);
f.seekg(sizeof(hashLength) + hashLength, std::fstream::beg);
int numberOfEntries = 0;
f.read((char*)&numberOfEntries, sizeof(int));
CV_Assert(numberOfEntries > 0);
if (numberOfEntries != MAX_ENTRIES)
{
std::cerr << "Invalid file: " << fileName_ << std::endl;
f.close();
remove(fileName_.c_str());
return false;
}
size_t tableEntriesOffset = (size_t)f.tellg();
std::vector<int> firstEntryOffset(numberOfEntries);
f.read((char*)&firstEntryOffset[0], sizeof(int)*numberOfEntries);
int entryNum = getHash(options);
int entryOffset = firstEntryOffset[entryNum];
ProgramFileConfigurationEntry entry;
while (entryOffset > 0)
{
f.seekg(entryOffset, std::fstream::beg);
assert(sizeof(entry) == sizeof(int)*3);
f.read((char*)&entry, sizeof(entry));
std::vector<char> foptions(entry.optionsLength);
if ((int)options.length() == entry.optionsLength)
{
if (entry.optionsLength > 0)
f.read(&foptions[0], entry.optionsLength);
CV_Assert(memcmp(&foptions, options.c_str(), entry.optionsLength) != 0);
}
if (entry.nextEntry <= 0)
break;
entryOffset = entry.nextEntry;
}
if (entryOffset > 0)
{
f.seekp(entryOffset, std::fstream::beg);
entry.nextEntry = fileSize;
f.write((char*)&entry, sizeof(entry));
} }
else else
{ {
fseek(fp, 0, SEEK_END); firstEntryOffset[entryNum] = fileSize;
size_t binarySize = ftell(fp); f.seekp(tableEntriesOffset, std::fstream::beg);
fseek(fp, 0, SEEK_SET); f.write((char*)&firstEntryOffset[0], sizeof(int)*numberOfEntries);
char *binary = new char[binarySize]; }
CV_Assert(1 == fread(binary, binarySize, 1, fp)); f.seekp(fileSize, std::fstream::beg);
fclose(fp); entry.nextEntry = 0;
cl_int status = 0; entry.dataSize = buf.size();
cl_device_id device = getClDeviceID(ctx); entry.optionsLength = options.length();
program = clCreateProgramWithBinary(getClContext(ctx), f.write((char*)&entry, sizeof(entry));
1, f.write(options.c_str(), entry.optionsLength);
&device, f.write(&buf[0], entry.dataSize);
(const size_t *)&binarySize, return true;
(const unsigned char **)&binary, }
NULL,
&status); cl_program getOrBuildProgram(const Context* ctx, const cv::ocl::ProgramEntry* source, const string& options)
{
cl_int status = 0;
cl_program program = NULL;
std::vector<char> binary;
if (!enable_disk_cache || !readConfigurationFromFile(options, binary))
{
program = clCreateProgramWithSource(getClContext(ctx), 1, (const char**)&source->programStr, NULL, &status);
openCLVerifyCall(status); openCLVerifyCall(status);
status = clBuildProgram(program, 1, &device, all_build_options.c_str(), NULL, NULL); cl_device_id device = getClDeviceID(ctx);
delete[] binary; status = clBuildProgram(program, 1, &device, options.c_str(), NULL, NULL);
if(status == CL_SUCCESS)
{
if (enable_disk_cache)
{
size_t binarySize;
openCLSafeCall(clGetProgramInfo(program,
CL_PROGRAM_BINARY_SIZES,
sizeof(size_t),
&binarySize, NULL));
std::vector<char> binary(binarySize);
char* ptr = &binary[0];
openCLSafeCall(clGetProgramInfo(program,
CL_PROGRAM_BINARIES,
sizeof(char*),
&ptr,
NULL));
if (!writeConfigurationToFile(options, binary))
{
std::cerr << "Can't write data to file: " << fileName_ << std::endl;
}
}
}
}
else
{
cl_device_id device = getClDeviceID(ctx);
size_t size = binary.size();
const char* ptr = &binary[0];
program = clCreateProgramWithBinary(getClContext(ctx),
1, &device,
(const size_t *)&size, (const unsigned char **)&ptr,
NULL, &status);
openCLVerifyCall(status);
status = clBuildProgram(program, 1, &device, options.c_str(), NULL, NULL);
} }
if(status != CL_SUCCESS) if(status != CL_SUCCESS)
@ -259,53 +445,77 @@ cl_program ProgramCache::getProgram(const Context *ctx, const char **source, str
memset(buildLog, 0, buildLogSize); memset(buildLog, 0, buildLogSize);
openCLSafeCall(clGetProgramBuildInfo(program, getClDeviceID(ctx), openCLSafeCall(clGetProgramBuildInfo(program, getClDeviceID(ctx),
CL_PROGRAM_BUILD_LOG, buildLogSize, buildLog, NULL)); CL_PROGRAM_BUILD_LOG, buildLogSize, buildLog, NULL));
std::cout << "\n\t\t\tBUILD LOG\n"; std::cout << "\nBUILD LOG: " << options << "\n";
std::cout << buildLog << endl; std::cout << buildLog << endl;
delete [] buildLog; delete [] buildLog;
} }
openCLVerifyCall(status); openCLVerifyCall(status);
} }
//Cache the binary for future use if build_options is null return program;
if( (this->cacheSize += 1) < MAX_PROG_CACHE_SIZE) }
this->addProgram(srcsign, program); };
else
cout << "Warning: code cache has been full.\n"; cl_program ProgramCache::getProgram(const Context *ctx, const cv::ocl::ProgramEntry* source,
const char *build_options)
{
stringstream src_sign;
src_sign << (int64)(source->programStr);
src_sign << getClContext(ctx);
if (NULL != build_options)
{
src_sign << "_" << build_options;
}
{
cv::AutoLock lockCache(mutexCache);
cl_program program = ProgramCache::getProgramCache()->progLookup(src_sign.str());
if (!!program)
{
clRetainProgram(program);
return program;
}
}
cv::AutoLock lockCache(mutexFiles);
// second check
{
cv::AutoLock lockCache(mutexCache);
cl_program program = ProgramCache::getProgramCache()->progLookup(src_sign.str());
if (!!program)
{
clRetainProgram(program);
return program;
}
}
string all_build_options;
if (!ctx->getDeviceInfo().compilationExtraOptions.empty())
all_build_options += ctx->getDeviceInfo().compilationExtraOptions;
if (build_options != NULL)
{
all_build_options += " ";
all_build_options += build_options;
}
const DeviceInfo& devInfo = ctx->getDeviceInfo();
string filename = binpath + (source->name ? source->name : "NULL") + "_" + devInfo.platform->platformName + "_" + devInfo.deviceName + ".clb";
ProgramFileCache programFileCache(filename, source->programHash);
cl_program program = programFileCache.getOrBuildProgram(ctx, source, all_build_options);
//Cache the binary for future use if build_options is null
if( (this->cacheSize += 1) < MAX_PROG_CACHE_SIZE)
{
cv::AutoLock lockCache(mutexCache);
this->addProgram(src_sign.str(), program);
}
else
{
cout << "Warning: code cache has been full.\n";
} }
return program; return program;
} }
//// Converts the contents of a file into a string
//static int convertToString(const char *filename, std::string& s)
//{
// size_t size;
// char* str;
//
// std::fstream f(filename, (std::fstream::in | std::fstream::binary));
// if(f.is_open())
// {
// size_t fileSize;
// f.seekg(0, std::fstream::end);
// size = fileSize = (size_t)f.tellg();
// f.seekg(0, std::fstream::beg);
//
// str = new char[size+1];
// if(!str)
// {
// f.close();
// return -1;
// }
//
// f.read(str, fileSize);
// f.close();
// str[size] = '\0';
//
// s = str;
// delete[] str;
// return 0;
// }
// printf("Error: Failed to open file %s\n", filename);
// return -1;
//}
} // namespace ocl } // namespace ocl
} // namespace cv } // namespace cv

47
modules/ocl/src/binarycaching.hpp → modules/ocl/src/cl_programcache.hpp
View File

@ -44,41 +44,42 @@
#include "precomp.hpp" #include "precomp.hpp"
using namespace cv; namespace cv {
using namespace cv::ocl; namespace ocl {
using namespace std;
using std::cout;
using std::endl;
namespace cv { namespace ocl {
class ProgramCache class ProgramCache
{ {
protected: protected:
ProgramCache(); ProgramCache();
~ProgramCache(); ~ProgramCache();
friend class std::auto_ptr<ProgramCache>; friend class std::auto_ptr<ProgramCache>;
public: public:
static ProgramCache *getProgramCache(); static ProgramCache *getProgramCache();
cl_program getProgram(const Context *ctx, const char **source, string kernelName, cl_program getProgram(const Context *ctx, const cv::ocl::ProgramEntry* source,
const char *build_options); const char *build_options);
void releaseProgram(); void releaseProgram();
protected: protected:
//lookup the binary given the file name //lookup the binary given the file name
cl_program progLookup(string srcsign); // (with acquired mutexCache)
cl_program progLookup(const string& srcsign);
//add program to the cache //add program to the cache
void addProgram(string srcsign, cl_program program); // (with acquired mutexCache)
void addProgram(const string& srcsign, cl_program program);
map <string, cl_program> codeCache; map <string, cl_program> codeCache;
unsigned int cacheSize; unsigned int cacheSize;
//The presumed watermark for the cache volume (256MB). Is it enough? //The presumed watermark for the cache volume (256MB). Is it enough?
//We may need more delicate algorithms when necessary later. //We may need more delicate algorithms when necessary later.
//Right now, let's just leave it along. //Right now, let's just leave it along.
static const unsigned MAX_PROG_CACHE_SIZE = 1024; static const unsigned MAX_PROG_CACHE_SIZE = 1024;
// acquire both mutexes in this order: 1) mutexFiles 2) mutexCache
static cv::Mutex mutexFiles;
static cv::Mutex mutexCache;
}; };
}//namespace ocl }//namespace ocl

12
modules/ocl/src/imgproc.cpp
View File

@ -1108,7 +1108,7 @@ namespace cv
CV_Assert(Dx.offset == 0 && Dy.offset == 0); CV_Assert(Dx.offset == 0 && Dy.offset == 0);
} }
static void corner_ocl(const char *src_str, string kernelName, int block_size, float k, oclMat &Dx, oclMat &Dy, static void corner_ocl(const cv::ocl::ProgramEntry* source, string kernelName, int block_size, float k, oclMat &Dx, oclMat &Dy,
oclMat &dst, int border_type) oclMat &dst, int border_type)
{ {
char borderType[30]; char borderType[30];
@ -1160,7 +1160,7 @@ namespace cv
args.push_back( make_pair(sizeof(cl_int), (void *)&dst.cols)); args.push_back( make_pair(sizeof(cl_int), (void *)&dst.cols));
args.push_back( make_pair(sizeof(cl_int), (void *)&dst.step)); args.push_back( make_pair(sizeof(cl_int), (void *)&dst.step));
args.push_back( make_pair( sizeof(cl_float) , (void *)&k)); args.push_back( make_pair( sizeof(cl_float) , (void *)&k));
openCLExecuteKernel(dst.clCxt, &src_str, kernelName, gt, lt, args, -1, -1, build_options); openCLExecuteKernel(dst.clCxt, source, kernelName, gt, lt, args, -1, -1, build_options);
} }
void cornerHarris(const oclMat &src, oclMat &dst, int blockSize, int ksize, void cornerHarris(const oclMat &src, oclMat &dst, int blockSize, int ksize,
@ -1181,7 +1181,7 @@ namespace cv
CV_Assert(borderType == cv::BORDER_CONSTANT || borderType == cv::BORDER_REFLECT101 || borderType == cv::BORDER_REPLICATE || borderType == cv::BORDER_REFLECT); CV_Assert(borderType == cv::BORDER_CONSTANT || borderType == cv::BORDER_REFLECT101 || borderType == cv::BORDER_REPLICATE || borderType == cv::BORDER_REFLECT);
extractCovData(src, dx, dy, blockSize, ksize, borderType); extractCovData(src, dx, dy, blockSize, ksize, borderType);
dst.create(src.size(), CV_32F); dst.create(src.size(), CV_32F);
corner_ocl(imgproc_calcHarris, "calcHarris", blockSize, static_cast<float>(k), dx, dy, dst, borderType); corner_ocl(&imgproc_calcHarris, "calcHarris", blockSize, static_cast<float>(k), dx, dy, dst, borderType);
} }
void cornerMinEigenVal(const oclMat &src, oclMat &dst, int blockSize, int ksize, int borderType) void cornerMinEigenVal(const oclMat &src, oclMat &dst, int blockSize, int ksize, int borderType)
@ -1200,7 +1200,7 @@ namespace cv
CV_Assert(borderType == cv::BORDER_CONSTANT || borderType == cv::BORDER_REFLECT101 || borderType == cv::BORDER_REPLICATE || borderType == cv::BORDER_REFLECT); CV_Assert(borderType == cv::BORDER_CONSTANT || borderType == cv::BORDER_REFLECT101 || borderType == cv::BORDER_REPLICATE || borderType == cv::BORDER_REFLECT);
extractCovData(src, dx, dy, blockSize, ksize, borderType); extractCovData(src, dx, dy, blockSize, ksize, borderType);
dst.create(src.size(), CV_32F); dst.create(src.size(), CV_32F);
corner_ocl(imgproc_calcMinEigenVal, "calcMinEigenVal", blockSize, 0, dx, dy, dst, borderType); corner_ocl(&imgproc_calcMinEigenVal, "calcMinEigenVal", blockSize, 0, dx, dy, dst, borderType);
} }
/////////////////////////////////// MeanShiftfiltering /////////////////////////////////////////////// /////////////////////////////////// MeanShiftfiltering ///////////////////////////////////////////////
static void meanShiftFiltering_gpu(const oclMat &src, oclMat dst, int sp, int sr, int maxIter, float eps) static void meanShiftFiltering_gpu(const oclMat &src, oclMat dst, int sp, int sr, int maxIter, float eps)
@ -1749,7 +1749,7 @@ namespace cv
} }
//////////////////////////////////convolve//////////////////////////////////////////////////// //////////////////////////////////convolve////////////////////////////////////////////////////
static void convolve_run(const oclMat &src, const oclMat &temp1, oclMat &dst, string kernelName, const char **kernelString) static void convolve_run(const oclMat &src, const oclMat &temp1, oclMat &dst, string kernelName, const cv::ocl::ProgramEntry* source)
{ {
CV_Assert(src.depth() == CV_32FC1); CV_Assert(src.depth() == CV_32FC1);
CV_Assert(temp1.depth() == CV_32F); CV_Assert(temp1.depth() == CV_32F);
@ -1784,7 +1784,7 @@ static void convolve_run(const oclMat &src, const oclMat &temp1, oclMat &dst, st
args.push_back( make_pair( sizeof(cl_int), (void *)&temp1.rows )); args.push_back( make_pair( sizeof(cl_int), (void *)&temp1.rows ));
args.push_back( make_pair( sizeof(cl_int), (void *)&temp1.cols )); args.push_back( make_pair( sizeof(cl_int), (void *)&temp1.cols ));
openCLExecuteKernel(clCxt, kernelString, kernelName, globalThreads, localThreads, args, -1, depth); openCLExecuteKernel(clCxt, source, kernelName, globalThreads, localThreads, args, -1, depth);
} }
void cv::ocl::convolve(const oclMat &x, const oclMat &t, oclMat &y) void cv::ocl::convolve(const oclMat &x, const oclMat &t, oclMat &y)
{ {

10
modules/ocl/src/mcwutil.cpp
View File

@ -72,7 +72,7 @@ namespace cv
namespace ocl namespace ocl
{ {
// provide additional methods for the user to interact with the command queue after a task is fired // provide additional methods for the user to interact with the command queue after a task is fired
static void openCLExecuteKernel_2(Context *clCxt , const char **source, string kernelName, size_t globalThreads[3], static void openCLExecuteKernel_2(Context *clCxt, const cv::ocl::ProgramEntry* source, string kernelName, size_t globalThreads[3],
size_t localThreads[3], vector< pair<size_t, const void *> > &args, int channels, size_t localThreads[3], vector< pair<size_t, const void *> > &args, int channels,
int depth, char *build_options, FLUSH_MODE finish_mode) int depth, char *build_options, FLUSH_MODE finish_mode)
{ {
@ -118,14 +118,14 @@ namespace cv
openCLSafeCall(clReleaseKernel(kernel)); openCLSafeCall(clReleaseKernel(kernel));
} }
void openCLExecuteKernel2(Context *clCxt , const char **source, string kernelName, void openCLExecuteKernel2(Context *clCxt, const cv::ocl::ProgramEntry* source, string kernelName,
size_t globalThreads[3], size_t localThreads[3], size_t globalThreads[3], size_t localThreads[3],
vector< pair<size_t, const void *> > &args, int channels, int depth, FLUSH_MODE finish_mode) vector< pair<size_t, const void *> > &args, int channels, int depth, FLUSH_MODE finish_mode)
{ {
openCLExecuteKernel2(clCxt, source, kernelName, globalThreads, localThreads, args, openCLExecuteKernel2(clCxt, source, kernelName, globalThreads, localThreads, args,
channels, depth, NULL, finish_mode); channels, depth, NULL, finish_mode);
} }
void openCLExecuteKernel2(Context *clCxt , const char **source, string kernelName, void openCLExecuteKernel2(Context *clCxt, const cv::ocl::ProgramEntry* source, string kernelName,
size_t globalThreads[3], size_t localThreads[3], size_t globalThreads[3], size_t localThreads[3],
vector< pair<size_t, const void *> > &args, int channels, int depth, char *build_options, FLUSH_MODE finish_mode) vector< pair<size_t, const void *> > &args, int channels, int depth, char *build_options, FLUSH_MODE finish_mode)
@ -249,7 +249,7 @@ namespace cv
bool support_image2d(Context *clCxt) bool support_image2d(Context *clCxt)
{ {
static const char * _kernel_string = "__kernel void test_func(image2d_t img) {}"; const cv::ocl::ProgramEntry _kernel = {NULL, "__kernel void test_func(image2d_t img) {}", NULL};
static bool _isTested = false; static bool _isTested = false;
static bool _support = false; static bool _support = false;
if(_isTested) if(_isTested)
@ -258,7 +258,7 @@ namespace cv
} }
try try
{ {
cv::ocl::openCLGetKernelFromSource(clCxt, &_kernel_string, "test_func"); cv::ocl::openCLGetKernelFromSource(clCxt, &_kernel, "test_func");
cv::ocl::finish(); cv::ocl::finish();
_support = true; _support = true;
} }

2
modules/ocl/src/moments.cpp
View File

@ -229,7 +229,7 @@ static void ocl_cvMoments( const void* array, CvMoments* mom, int binary )
CV_Error( CV_StsBadArg, "The passed sequence is not a valid contour" ); CV_Error( CV_StsBadArg, "The passed sequence is not a valid contour" );
} }
if( !moments ) if( !mom )
CV_Error( CV_StsNullPtr, "" ); CV_Error( CV_StsNullPtr, "" );
memset( mom, 0, sizeof(*mom)); memset( mom, 0, sizeof(*mom));

1
modules/ocl/test/main.cpp
View File

@ -118,7 +118,6 @@ int main(int argc, char **argv)
} }
cv::ocl::setDevice(devicesInfo[device]); cv::ocl::setDevice(devicesInfo[device]);
setBinaryDiskCache(CACHE_UPDATE);
cout << "Device type: " << type << endl cout << "Device type: " << type << endl
<< "Platform name: " << devicesInfo[device]->platform->platformName << endl << "Platform name: " << devicesInfo[device]->platform->platformName << endl