From d55f462c9c57b9fe5c3c146e126b50ed72090a3d Mon Sep 17 00:00:00 2001 From: Ray Smith Date: Tue, 6 Dec 2016 13:45:49 -0800 Subject: [PATCH] More clang-tidy from previous commits --- api/baseapi.h | 4 +- ccmain/thresholder.cpp | 8 +- ccutil/strngs.h | 4 +- lstm/lstmtrainer.cpp | 3 +- lstm/lstmtrainer.h | 3 +- opencl/opencl_device_selection.h | 65 ++- opencl/openclwrapper.cpp | 840 ++++++++++++++----------------- opencl/openclwrapper.h | 11 +- textord/blkocc.h | 5 +- 9 files changed, 412 insertions(+), 531 deletions(-) diff --git a/api/baseapi.h b/api/baseapi.h index 0f489e0ba5..2e1bb8f30c 100644 --- a/api/baseapi.h +++ b/api/baseapi.h @@ -847,9 +847,7 @@ class TESS_API TessBaseAPI { int** y1, PAGE_RES* page_res); - TESS_LOCAL const PAGE_RES* GetPageRes() const { - return page_res_; - } + TESS_LOCAL const PAGE_RES* GetPageRes() const { return page_res_; } /* @} */ diff --git a/ccmain/thresholder.cpp b/ccmain/thresholder.cpp index fc8111ef7b..77069bc9d9 100644 --- a/ccmain/thresholder.cpp +++ b/ccmain/thresholder.cpp @@ -268,10 +268,10 @@ void ImageThresholder::OtsuThresholdRectToPix(Pix* src_pix, OpenclDevice od; if ((num_channels == 4 || num_channels == 1) && od.selectedDeviceIsOpenCL() && rect_top_ == 0 && rect_left_ == 0 ) { - od.ThresholdRectToPixOCL((unsigned char*)pixGetData(src_pix), - num_channels, pixGetWpl(src_pix) * 4, - thresholds, hi_values, out_pix /*pix_OCL*/, - rect_height_, rect_width_, rect_top_, rect_left_); + od.ThresholdRectToPixOCL((unsigned char*)pixGetData(src_pix), num_channels, + pixGetWpl(src_pix) * 4, thresholds, hi_values, + out_pix /*pix_OCL*/, rect_height_, rect_width_, + rect_top_, rect_left_); } else { #endif ThresholdRectToPix(src_pix, num_channels, thresholds, hi_values, out_pix); diff --git a/ccutil/strngs.h b/ccutil/strngs.h index ea1738895c..2e65463efd 100644 --- a/ccutil/strngs.h +++ b/ccutil/strngs.h @@ -147,9 +147,7 @@ class TESS_API STRING } // returns the string data part of storage - inline char* GetCStr() { - return ((char *)data_) + sizeof(STRING_HEADER); - } + inline char* GetCStr() { return ((char*)data_) + sizeof(STRING_HEADER); } inline const char* GetCStr() const { return ((const char *)data_) + sizeof(STRING_HEADER); diff --git a/lstm/lstmtrainer.cpp b/lstm/lstmtrainer.cpp index cffde1c97d..b3958d3431 100644 --- a/lstm/lstmtrainer.cpp +++ b/lstm/lstmtrainer.cpp @@ -555,7 +555,8 @@ void LSTMTrainer::StartSubtrainer(STRING* log_msg) { // Reduce learning rate so it doesn't diverge this time. sub_trainer_->ReduceLearningRates(this, log_msg); // If it fails again, we will wait twice as long before reverting again. - int stall_offset = learning_iteration() - sub_trainer_->learning_iteration(); + int stall_offset = + learning_iteration() - sub_trainer_->learning_iteration(); stall_iteration_ = learning_iteration() + 2 * stall_offset; sub_trainer_->stall_iteration_ = stall_iteration_; // Re-save the best trainer with the new learning rates and stall iteration. diff --git a/lstm/lstmtrainer.h b/lstm/lstmtrainer.h index c60405076d..2054284dd9 100644 --- a/lstm/lstmtrainer.h +++ b/lstm/lstmtrainer.h @@ -113,7 +113,8 @@ class LSTMTrainer : public LSTMRecognizer { // previously setup UNICHARSET and UnicharCompress. // ctc_mode controls how the truth text is mapped to the network targets. // Note: Call before InitNetwork! - void InitCharSet(const UNICHARSET& unicharset, const UnicharCompress& recoder); + void InitCharSet(const UNICHARSET& unicharset, + const UnicharCompress& recoder); // Initializes the trainer with a network_spec in the network description // net_flags control network behavior according to the NetworkFlags enum. diff --git a/opencl/opencl_device_selection.h b/opencl/opencl_device_selection.h index c5d704ecc6..4c61774fb3 100644 --- a/opencl/opencl_device_selection.h +++ b/opencl/opencl_device_selection.h @@ -68,8 +68,8 @@ typedef struct { typedef ds_status (*ds_score_release)(void* score); static ds_status releaseDSProfile(ds_profile* profile, ds_score_release sr) { ds_status status = DS_SUCCESS; - if (profile!=nullptr) { - if (profile->devices!=nullptr && sr!=nullptr) { + if (profile != nullptr) { + if (profile->devices != nullptr && sr != nullptr) { unsigned int i; for (i = 0; i < profile->numDevices; i++) { free(profile->devices[i].oclDeviceName); @@ -90,18 +90,16 @@ static ds_status initDSProfile(ds_profile** p, const char* version) { int numDevices; cl_uint numPlatforms; cl_platform_id* platforms = nullptr; - cl_device_id* devices = nullptr; + cl_device_id* devices = nullptr; ds_status status = DS_SUCCESS; unsigned int next; unsigned int i; - if (p == nullptr) - return DS_INVALID_PROFILE; + if (p == nullptr) return DS_INVALID_PROFILE; ds_profile* profile = (ds_profile*)malloc(sizeof(ds_profile)); - if (profile == nullptr) - return DS_MEMORY_ERROR; - + if (profile == nullptr) return DS_MEMORY_ERROR; + memset(profile, 0, sizeof(ds_profile)); clGetPlatformIDs(0, nullptr, &numPlatforms); @@ -131,7 +129,8 @@ static ds_status initDSProfile(ds_profile** p, const char* version) { } profile->numDevices = numDevices+1; // +1 to numDevices to include the native CPU - profile->devices = (ds_device*)malloc(profile->numDevices*sizeof(ds_device)); + profile->devices = + (ds_device*)malloc(profile->numDevices * sizeof(ds_device)); if (profile->devices == nullptr) { profile->numDevices = 0; status = DS_MEMORY_ERROR; @@ -151,14 +150,14 @@ static ds_status initDSProfile(ds_profile** p, const char* version) { profile->devices[next].type = DS_DEVICE_OPENCL_DEVICE; profile->devices[next].oclDeviceID = devices[j]; - clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_NAME - , DS_DEVICE_NAME_LENGTH, &buffer, nullptr); + clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_NAME, + DS_DEVICE_NAME_LENGTH, &buffer, nullptr); length = strlen(buffer); profile->devices[next].oclDeviceName = (char*)malloc(length+1); memcpy(profile->devices[next].oclDeviceName, buffer, length+1); - clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DRIVER_VERSION - , DS_DEVICE_NAME_LENGTH, &buffer, nullptr); + clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DRIVER_VERSION, + DS_DEVICE_NAME_LENGTH, &buffer, nullptr); length = strlen(buffer); profile->devices[next].oclDriverVersion = (char*)malloc(length+1); memcpy(profile->devices[next].oclDriverVersion, buffer, length+1); @@ -213,8 +212,7 @@ static ds_status profileDevices(ds_profile* profile, switch (type) { case DS_EVALUATE_NEW_ONLY: - if (profile->devices[i].score != nullptr) - break; + if (profile->devices[i].score != nullptr) break; // else fall through case DS_EVALUATE_ALL: evaluatorStatus = evaluator(profile->devices+i, evaluatorData); @@ -260,11 +258,10 @@ static ds_status writeProfileToFile(ds_profile* profile, const char* file) { ds_status status = DS_SUCCESS; - if (profile == nullptr) - return DS_INVALID_PROFILE; + if (profile == nullptr) return DS_INVALID_PROFILE; FILE* profileFile = fopen(file, "wb"); - if (profileFile==nullptr) { + if (profileFile == nullptr) { status = DS_FILE_ERROR; } else { @@ -327,7 +324,8 @@ static ds_status writeProfileToFile(ds_profile* profile, fwrite(DS_TAG_SCORE, sizeof(char), strlen(DS_TAG_SCORE), profileFile); status = serializer(profile->devices+i, &serializedScore, &serializedScoreSize); - if (status == DS_SUCCESS && serializedScore!=nullptr && serializedScoreSize > 0) { + if (status == DS_SUCCESS && serializedScore != nullptr && + serializedScoreSize > 0) { fwrite(serializedScore, sizeof(char), serializedScoreSize, profileFile); free(serializedScore); } @@ -349,7 +347,7 @@ static ds_status readProFile(const char* fileName, char** content, *content = nullptr; FILE* input = fopen(fileName, "rb"); - if(input == nullptr) { + if (input == nullptr) { return DS_FILE_ERROR; } @@ -357,7 +355,7 @@ static ds_status readProFile(const char* fileName, char** content, size = ftell(input); rewind(input); char* binary = (char*)malloc(size); - if(binary == nullptr) { + if (binary == nullptr) { fclose(input); return DS_FILE_ERROR; } @@ -403,8 +401,7 @@ static ds_status readProfileFromFile(ds_profile* profile, const char* contentEnd = nullptr; size_t contentSize; - if (profile==nullptr) - return DS_INVALID_PROFILE; + if (profile == nullptr) return DS_INVALID_PROFILE; status = readProFile(file, &contentStart, &contentSize); if (status == DS_SUCCESS) { @@ -426,7 +423,7 @@ static ds_status readProfileFromFile(ds_profile* profile, dataStart += strlen(DS_TAG_VERSION); dataEnd = findString(dataStart, contentEnd, DS_TAG_VERSION_END); - if (dataEnd==nullptr) { + if (dataEnd == nullptr) { status = DS_PROFILE_FILE_ERROR; goto cleanup; } @@ -458,27 +455,27 @@ static ds_status readProfileFromFile(ds_profile* profile, const char* deviceDriverEnd; dataStart = findString(currentPosition, contentEnd, DS_TAG_DEVICE); - if (dataStart==nullptr) { + if (dataStart == nullptr) { // nothing useful remain, quit... break; } dataStart+=strlen(DS_TAG_DEVICE); dataEnd = findString(dataStart, contentEnd, DS_TAG_DEVICE_END); - if (dataEnd==nullptr) { + if (dataEnd == nullptr) { status = DS_PROFILE_FILE_ERROR; goto cleanup; } // parse the device type deviceTypeStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_TYPE); - if (deviceTypeStart==nullptr) { + if (deviceTypeStart == nullptr) { status = DS_PROFILE_FILE_ERROR; goto cleanup; } deviceTypeStart+=strlen(DS_TAG_DEVICE_TYPE); deviceTypeEnd = findString(deviceTypeStart, contentEnd, DS_TAG_DEVICE_TYPE_END); - if (deviceTypeEnd==nullptr) { + if (deviceTypeEnd == nullptr) { status = DS_PROFILE_FILE_ERROR; goto cleanup; } @@ -489,14 +486,14 @@ static ds_status readProfileFromFile(ds_profile* profile, if (deviceType == DS_DEVICE_OPENCL_DEVICE) { deviceNameStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_NAME); - if (deviceNameStart==nullptr) { + if (deviceNameStart == nullptr) { status = DS_PROFILE_FILE_ERROR; goto cleanup; } deviceNameStart+=strlen(DS_TAG_DEVICE_NAME); deviceNameEnd = findString(deviceNameStart, contentEnd, DS_TAG_DEVICE_NAME_END); - if (deviceNameEnd==nullptr) { + if (deviceNameEnd == nullptr) { status = DS_PROFILE_FILE_ERROR; goto cleanup; } @@ -504,14 +501,14 @@ static ds_status readProfileFromFile(ds_profile* profile, deviceDriverStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_DRIVER_VERSION); - if (deviceDriverStart==nullptr) { + if (deviceDriverStart == nullptr) { status = DS_PROFILE_FILE_ERROR; goto cleanup; } deviceDriverStart+=strlen(DS_TAG_DEVICE_DRIVER_VERSION); deviceDriverEnd = findString(deviceDriverStart, contentEnd, DS_TAG_DEVICE_DRIVER_VERSION_END); - if (deviceDriverEnd ==nullptr) { + if (deviceDriverEnd == nullptr) { status = DS_PROFILE_FILE_ERROR; goto cleanup; } @@ -532,7 +529,7 @@ static ds_status readProfileFromFile(ds_profile* profile, && strncmp(profile->devices[i].oclDriverVersion, deviceDriverStart, driverVersionLength)==0) { deviceScoreStart = findString(dataStart, contentEnd, DS_TAG_SCORE); - if (deviceNameStart==nullptr) { + if (deviceNameStart == nullptr) { status = DS_PROFILE_FILE_ERROR; goto cleanup; } @@ -554,7 +551,7 @@ static ds_status readProfileFromFile(ds_profile* profile, for (i = 0; i < profile->numDevices; i++) { if (profile->devices[i].type == DS_DEVICE_NATIVE_CPU) { deviceScoreStart = findString(dataStart, contentEnd, DS_TAG_SCORE); - if (deviceScoreStart==nullptr) { + if (deviceScoreStart == nullptr) { status = DS_PROFILE_FILE_ERROR; goto cleanup; } diff --git a/opencl/openclwrapper.cpp b/opencl/openclwrapper.cpp index 7671bf86f0..50bb2f857e 100644 --- a/opencl/openclwrapper.cpp +++ b/opencl/openclwrapper.cpp @@ -112,10 +112,11 @@ void legalizeFileName( char *fileName) { // initial ./ is valid for present directory //if (*pos == '.') pos++; //if (*pos == '/') pos++; - for ( char *pos = strstr(fileName, invalidStr); pos != nullptr; pos = strstr(pos+1, invalidStr)) { - //printf("\tfound: %s, ", pos); - pos[0] = '_'; - //printf("fileName: %s\n", fileName); + for (char *pos = strstr(fileName, invalidStr); pos != nullptr; + pos = strstr(pos + 1, invalidStr)) { + // printf("\tfound: %s, ", pos); + pos[0] = '_'; + // printf("fileName: %s\n", fileName); } } } @@ -128,17 +129,22 @@ void populateGPUEnvFromDevice( GPUEnv *gpuInfo, cl_device_id device ) { gpuInfo->mpDevID = device; gpuInfo->mpArryDevsID = new cl_device_id[1]; gpuInfo->mpArryDevsID[0] = gpuInfo->mpDevID; - clStatus = clGetDeviceInfo(gpuInfo->mpDevID, CL_DEVICE_TYPE , sizeof(cl_device_type), &gpuInfo->mDevType , &size); + clStatus = + clGetDeviceInfo(gpuInfo->mpDevID, CL_DEVICE_TYPE, + sizeof(cl_device_type), &gpuInfo->mDevType, &size); CHECK_OPENCL( clStatus, "populateGPUEnv::getDeviceInfo(TYPE)"); // platform - clStatus = clGetDeviceInfo(gpuInfo->mpDevID, CL_DEVICE_PLATFORM , sizeof(cl_platform_id), &gpuInfo->mpPlatformID , &size); + clStatus = + clGetDeviceInfo(gpuInfo->mpDevID, CL_DEVICE_PLATFORM, + sizeof(cl_platform_id), &gpuInfo->mpPlatformID, &size); CHECK_OPENCL( clStatus, "populateGPUEnv::getDeviceInfo(PLATFORM)"); // context cl_context_properties props[3]; props[0] = CL_CONTEXT_PLATFORM; props[1] = (cl_context_properties) gpuInfo->mpPlatformID; props[2] = 0; - gpuInfo->mpContext = clCreateContext(props, 1, &gpuInfo->mpDevID, nullptr, nullptr, &clStatus); + gpuInfo->mpContext = clCreateContext(props, 1, &gpuInfo->mpDevID, nullptr, + nullptr, &clStatus); CHECK_OPENCL( clStatus, "populateGPUEnv::createContext"); // queue cl_command_queue_properties queueProperties = 0; @@ -149,16 +155,15 @@ void populateGPUEnvFromDevice( GPUEnv *gpuInfo, cl_device_id device ) { int OpenclDevice::LoadOpencl() { #ifdef WIN32 - HINSTANCE HOpenclDll = nullptr; - void * OpenclDll = nullptr; - //fprintf(stderr, " LoadOpenclDllxx... \n"); - OpenclDll = static_cast( HOpenclDll ); - OpenclDll = LoadLibrary( "openCL.dll" ); - if ( !static_cast( OpenclDll ) ) - { - fprintf(stderr, "[OD] Load opencl.dll failed!\n"); - FreeLibrary( static_cast( OpenclDll ) ); - return 0; + HINSTANCE HOpenclDll = nullptr; + void *OpenclDll = nullptr; + // fprintf(stderr, " LoadOpenclDllxx... \n"); + OpenclDll = static_cast(HOpenclDll); + OpenclDll = LoadLibrary("openCL.dll"); + if (!static_cast(OpenclDll)) { + fprintf(stderr, "[OD] Load opencl.dll failed!\n"); + FreeLibrary(static_cast(OpenclDll)); + return 0; } fprintf(stderr, "[OD] Load opencl.dll successful!\n"); #endif @@ -205,7 +210,8 @@ PIX *mapOutputCLBuffer(KernelEnv rEnv, cl_mem clbuffer, PIX *pixd, PIX *pixs, pixSetData(pixd, pValues); } - clEnqueueUnmapMemObject(rEnv.mpkCmdQueue, clbuffer, pValues, 0, nullptr, nullptr); + clEnqueueUnmapMemObject(rEnv.mpkCmdQueue, clbuffer, pValues, 0, nullptr, + nullptr); if (sync) { clFinish(rEnv.mpkCmdQueue); @@ -216,20 +222,21 @@ PIX *mapOutputCLBuffer(KernelEnv rEnv, cl_mem clbuffer, PIX *pixd, PIX *pixs, cl_mem allocateIntBuffer( KernelEnv rEnv, const l_uint32 *_pValues, size_t nElements, cl_int *pStatus , bool sync = false) { - cl_mem xValues = clCreateBuffer( rEnv.mpkContext, (cl_mem_flags) (CL_MEM_READ_WRITE), - nElements * sizeof(l_int32), nullptr, pStatus); + cl_mem xValues = + clCreateBuffer(rEnv.mpkContext, (cl_mem_flags)(CL_MEM_READ_WRITE), + nElements * sizeof(l_int32), nullptr, pStatus); - if (_pValues != nullptr) - { - l_int32 *pValues = (l_int32 *)clEnqueueMapBuffer( rEnv.mpkCmdQueue, xValues, CL_TRUE, CL_MAP_WRITE, 0, - nElements * sizeof(l_int32), 0, nullptr, nullptr, nullptr ); + if (_pValues != nullptr) { + l_int32 *pValues = (l_int32 *)clEnqueueMapBuffer( + rEnv.mpkCmdQueue, xValues, CL_TRUE, CL_MAP_WRITE, 0, + nElements * sizeof(l_int32), 0, nullptr, nullptr, nullptr); - memcpy(pValues, _pValues, nElements * sizeof(l_int32)); + memcpy(pValues, _pValues, nElements * sizeof(l_int32)); - clEnqueueUnmapMemObject(rEnv.mpkCmdQueue,xValues,pValues,0,nullptr,nullptr); + clEnqueueUnmapMemObject(rEnv.mpkCmdQueue, xValues, pValues, 0, nullptr, + nullptr); - if (sync) - clFinish( rEnv.mpkCmdQueue ); + if (sync) clFinish(rEnv.mpkCmdQueue); } return xValues; @@ -238,27 +245,25 @@ PIX *mapOutputCLBuffer(KernelEnv rEnv, cl_mem clbuffer, PIX *pixd, PIX *pixs, void OpenclDevice::releaseMorphCLBuffers() { - if (pixdCLIntermediate != nullptr) - clReleaseMemObject(pixdCLIntermediate); - if (pixsCLBuffer != nullptr) - clReleaseMemObject(pixsCLBuffer); - if (pixdCLBuffer != nullptr) - clReleaseMemObject(pixdCLBuffer); - if (pixThBuffer != nullptr) - clReleaseMemObject(pixThBuffer); - pixdCLIntermediate = pixsCLBuffer = pixdCLBuffer = pixThBuffer = nullptr; + if (pixdCLIntermediate != nullptr) clReleaseMemObject(pixdCLIntermediate); + if (pixsCLBuffer != nullptr) clReleaseMemObject(pixsCLBuffer); + if (pixdCLBuffer != nullptr) clReleaseMemObject(pixdCLBuffer); + if (pixThBuffer != nullptr) clReleaseMemObject(pixThBuffer); + pixdCLIntermediate = pixsCLBuffer = pixdCLBuffer = pixThBuffer = nullptr; } int OpenclDevice::initMorphCLAllocations(l_int32 wpl, l_int32 h, PIX* pixs) { SetKernelEnv( &rEnv ); - if (pixThBuffer != nullptr) - { - pixsCLBuffer = allocateZeroCopyBuffer(rEnv, nullptr, wpl*h, CL_MEM_ALLOC_HOST_PTR, &clStatus); + if (pixThBuffer != nullptr) { + pixsCLBuffer = allocateZeroCopyBuffer(rEnv, nullptr, wpl * h, + CL_MEM_ALLOC_HOST_PTR, &clStatus); - //Get the output from ThresholdToPix operation - clStatus = clEnqueueCopyBuffer(rEnv.mpkCmdQueue, pixThBuffer, pixsCLBuffer, 0, 0, sizeof(l_uint32) * wpl*h, 0, nullptr, nullptr); + // Get the output from ThresholdToPix operation + clStatus = + clEnqueueCopyBuffer(rEnv.mpkCmdQueue, pixThBuffer, pixsCLBuffer, 0, 0, + sizeof(l_uint32) * wpl * h, 0, nullptr, nullptr); } else { @@ -269,9 +274,11 @@ int OpenclDevice::initMorphCLAllocations(l_int32 wpl, l_int32 h, PIX* pixs) pixsCLBuffer = allocateZeroCopyBuffer(rEnv, srcdata, wpl*h, CL_MEM_USE_HOST_PTR, &clStatus); } - pixdCLBuffer = allocateZeroCopyBuffer(rEnv, nullptr, wpl*h, CL_MEM_ALLOC_HOST_PTR, &clStatus); + pixdCLBuffer = allocateZeroCopyBuffer(rEnv, nullptr, wpl * h, + CL_MEM_ALLOC_HOST_PTR, &clStatus); - pixdCLIntermediate = allocateZeroCopyBuffer(rEnv, nullptr, wpl*h, CL_MEM_ALLOC_HOST_PTR, &clStatus); + pixdCLIntermediate = allocateZeroCopyBuffer( + rEnv, nullptr, wpl * h, CL_MEM_ALLOC_HOST_PTR, &clStatus); return (int)clStatus; } @@ -414,9 +421,8 @@ int OpenclDevice::BinaryGenerated( const char * clFileName, FILE ** fhandle ) legalizeFileName(fileName); fd = fopen(fileName, "rb"); status = (fd != nullptr) ? 1 : 0; - if ( fd != nullptr ) - { - *fhandle = fd; + if (fd != nullptr) { + *fhandle = fd; } return status; @@ -428,9 +434,8 @@ int OpenclDevice::CachedOfKernerPrg( const GPUEnv *gpuEnvCached, const char * cl { if ( strcasecmp( gpuEnvCached->mArryKnelSrcFile[i], clFileName ) == 0 ) { - if ( gpuEnvCached->mpArryPrograms[i] != nullptr ) - { - return 1; + if (gpuEnvCached->mpArryPrograms[i] != nullptr) { + return 1; } } } @@ -439,11 +444,10 @@ int OpenclDevice::CachedOfKernerPrg( const GPUEnv *gpuEnvCached, const char * cl } int OpenclDevice::WriteBinaryToFile( const char* fileName, const char* birary, size_t numBytes ) { - FILE *output = nullptr; - output = fopen( fileName, "wb" ); - if ( output == nullptr ) - { - return 0; + FILE *output = nullptr; + output = fopen(fileName, "wb"); + if (output == nullptr) { + return 0; } fwrite( birary, sizeof(char), numBytes, output ); @@ -461,32 +465,32 @@ int OpenclDevice::GeneratBinFromKernelSource( cl_program program, const char * c cl_device_id *mpArryDevsID; char **binaries, *str = nullptr; - clStatus = clGetProgramInfo( program, CL_PROGRAM_NUM_DEVICES, - sizeof(numDevices), &numDevices, nullptr ); + clStatus = clGetProgramInfo(program, CL_PROGRAM_NUM_DEVICES, + sizeof(numDevices), &numDevices, nullptr); CHECK_OPENCL( clStatus, "clGetProgramInfo" ); mpArryDevsID = (cl_device_id*) malloc( sizeof(cl_device_id) * numDevices ); - if ( mpArryDevsID == nullptr ) - { - return 0; + if (mpArryDevsID == nullptr) { + return 0; } /* grab the handles to all of the devices in the program. */ - clStatus = clGetProgramInfo( program, CL_PROGRAM_DEVICES, - sizeof(cl_device_id) * numDevices, mpArryDevsID, nullptr ); + clStatus = clGetProgramInfo(program, CL_PROGRAM_DEVICES, + sizeof(cl_device_id) * numDevices, mpArryDevsID, + nullptr); CHECK_OPENCL( clStatus, "clGetProgramInfo" ); /* figure out the sizes of each of the binaries. */ binarySizes = (size_t*) malloc( sizeof(size_t) * numDevices ); - clStatus = clGetProgramInfo( program, CL_PROGRAM_BINARY_SIZES, - sizeof(size_t) * numDevices, binarySizes, nullptr ); + clStatus = + clGetProgramInfo(program, CL_PROGRAM_BINARY_SIZES, + sizeof(size_t) * numDevices, binarySizes, nullptr); CHECK_OPENCL( clStatus, "clGetProgramInfo" ); /* copy over all of the generated binaries. */ binaries = (char**) malloc( sizeof(char *) * numDevices ); - if ( binaries == nullptr ) - { - return 0; + if (binaries == nullptr) { + return 0; } for ( i = 0; i < numDevices; i++ ) @@ -494,19 +498,18 @@ int OpenclDevice::GeneratBinFromKernelSource( cl_program program, const char * c if ( binarySizes[i] != 0 ) { binaries[i] = (char*) malloc( sizeof(char) * binarySizes[i] ); - if ( binaries[i] == nullptr ) - { - return 0; + if (binaries[i] == nullptr) { + return 0; } } else { - binaries[i] = nullptr; + binaries[i] = nullptr; } } - clStatus = clGetProgramInfo( program, CL_PROGRAM_BINARIES, - sizeof(char *) * numDevices, binaries, nullptr ); + clStatus = clGetProgramInfo(program, CL_PROGRAM_BINARIES, + sizeof(char *) * numDevices, binaries, nullptr); CHECK_OPENCL(clStatus,"clGetProgramInfo"); /* dump out each binary into its own separate file. */ @@ -518,7 +521,7 @@ int OpenclDevice::GeneratBinFromKernelSource( cl_program program, const char * c { char deviceName[1024]; clStatus = clGetDeviceInfo(mpArryDevsID[i], CL_DEVICE_NAME, - sizeof(deviceName), deviceName, nullptr); + sizeof(deviceName), deviceName, nullptr); CHECK_OPENCL( clStatus, "clGetDeviceInfo" ); str = (char*) strstr( clFileName, (char*) ".cl" ); @@ -556,16 +559,16 @@ int OpenclDevice::GeneratBinFromKernelSource( cl_program program, const char * c void copyIntBuffer( KernelEnv rEnv, cl_mem xValues, const l_uint32 *_pValues, size_t nElements, cl_int *pStatus ) { - l_int32 *pValues = (l_int32 *)clEnqueueMapBuffer( rEnv.mpkCmdQueue, xValues, CL_TRUE, CL_MAP_WRITE, 0, - nElements * sizeof(l_int32), 0, nullptr, nullptr, nullptr ); - clFinish( rEnv.mpkCmdQueue ); - if (_pValues != nullptr) - { - for ( int i = 0; i < (int)nElements; i++ ) - pValues[i] = (l_int32)_pValues[i]; + l_int32 *pValues = (l_int32 *)clEnqueueMapBuffer( + rEnv.mpkCmdQueue, xValues, CL_TRUE, CL_MAP_WRITE, 0, + nElements * sizeof(l_int32), 0, nullptr, nullptr, nullptr); + clFinish(rEnv.mpkCmdQueue); + if (_pValues != nullptr) { + for (int i = 0; i < (int)nElements; i++) pValues[i] = (l_int32)_pValues[i]; } - clEnqueueUnmapMemObject(rEnv.mpkCmdQueue,xValues,pValues,0,nullptr,nullptr); + clEnqueueUnmapMemObject(rEnv.mpkCmdQueue, xValues, pValues, 0, nullptr, + nullptr); //clFinish( rEnv.mpkCmdQueue ); return; } @@ -599,14 +602,13 @@ int OpenclDevice::CompileKernelFile( GPUEnv *gpuInfo, const char *buildOption ) //PERF_COUNT_SUB("BinaryGenerated") if ( binaryExisted == 1 ) { - clStatus = clGetContextInfo( gpuInfo->mpContext, CL_CONTEXT_NUM_DEVICES, - sizeof(numDevices), &numDevices, nullptr ); - CHECK_OPENCL( clStatus, "clGetContextInfo" ); + clStatus = clGetContextInfo(gpuInfo->mpContext, CL_CONTEXT_NUM_DEVICES, + sizeof(numDevices), &numDevices, nullptr); + CHECK_OPENCL(clStatus, "clGetContextInfo"); - mpArryDevsID = (cl_device_id*) malloc( sizeof(cl_device_id) * numDevices ); - if ( mpArryDevsID == nullptr ) - { - return 0; + mpArryDevsID = (cl_device_id *)malloc(sizeof(cl_device_id) * numDevices); + if (mpArryDevsID == nullptr) { + return 0; } //PERF_COUNT_SUB("get numDevices") b_error = 0; @@ -633,8 +635,9 @@ int OpenclDevice::CompileKernelFile( GPUEnv *gpuInfo, const char *buildOption ) //PERF_COUNT_SUB("read file") fd = nullptr; // grab the handles to all of the devices in the context. - clStatus = clGetContextInfo( gpuInfo->mpContext, CL_CONTEXT_DEVICES, - sizeof( cl_device_id ) * numDevices, mpArryDevsID, nullptr ); + clStatus = clGetContextInfo(gpuInfo->mpContext, CL_CONTEXT_DEVICES, + sizeof(cl_device_id) * numDevices, + mpArryDevsID, nullptr); CHECK_OPENCL( clStatus, "clGetContextInfo" ); //PERF_COUNT_SUB("get devices") //fprintf(stderr, "[OD] Create kernel from binary\n"); @@ -646,7 +649,7 @@ int OpenclDevice::CompileKernelFile( GPUEnv *gpuInfo, const char *buildOption ) free( binary ); free( mpArryDevsID ); mpArryDevsID = nullptr; -//PERF_COUNT_SUB("binaryExisted") + // PERF_COUNT_SUB("binaryExisted") } else { @@ -658,9 +661,8 @@ int OpenclDevice::CompileKernelFile( GPUEnv *gpuInfo, const char *buildOption ) //PERF_COUNT_SUB("!binaryExisted") } - if ( gpuInfo->mpArryPrograms[idx] == (cl_program) nullptr ) - { - return 0; + if (gpuInfo->mpArryPrograms[idx] == (cl_program) nullptr) { + return 0; } //char options[512]; @@ -669,15 +671,17 @@ int OpenclDevice::CompileKernelFile( GPUEnv *gpuInfo, const char *buildOption ) PERF_COUNT_START("OD::CompileKernel::clBuildProgram") if (!gpuInfo->mnIsUserCreated) { - clStatus = clBuildProgram(gpuInfo->mpArryPrograms[idx], 1, gpuInfo->mpArryDevsID, - buildOption, nullptr, nullptr); -//PERF_COUNT_SUB("clBuildProgram notUserCreated") + clStatus = + clBuildProgram(gpuInfo->mpArryPrograms[idx], 1, gpuInfo->mpArryDevsID, + buildOption, nullptr, nullptr); + // PERF_COUNT_SUB("clBuildProgram notUserCreated") } else { - clStatus = clBuildProgram(gpuInfo->mpArryPrograms[idx], 1, &(gpuInfo->mpDevID), - buildOption, nullptr, nullptr); -//PERF_COUNT_SUB("clBuildProgram isUserCreated") + clStatus = + clBuildProgram(gpuInfo->mpArryPrograms[idx], 1, &(gpuInfo->mpDevID), + buildOption, nullptr, nullptr); + // PERF_COUNT_SUB("clBuildProgram isUserCreated") } PERF_COUNT_END if ( clStatus != CL_SUCCESS ) @@ -685,13 +689,15 @@ PERF_COUNT_END printf ("BuildProgram error!\n"); if ( !gpuInfo->mnIsUserCreated ) { - clStatus = clGetProgramBuildInfo( gpuInfo->mpArryPrograms[idx], gpuInfo->mpArryDevsID[0], - CL_PROGRAM_BUILD_LOG, 0, nullptr, &length ); + clStatus = clGetProgramBuildInfo( + gpuInfo->mpArryPrograms[idx], gpuInfo->mpArryDevsID[0], + CL_PROGRAM_BUILD_LOG, 0, nullptr, &length); } else { - clStatus = clGetProgramBuildInfo( gpuInfo->mpArryPrograms[idx], gpuInfo->mpDevID, - CL_PROGRAM_BUILD_LOG, 0, nullptr, &length); + clStatus = clGetProgramBuildInfo( + gpuInfo->mpArryPrograms[idx], gpuInfo->mpDevID, + CL_PROGRAM_BUILD_LOG, 0, nullptr, &length); } if ( clStatus != CL_SUCCESS ) { @@ -699,9 +705,8 @@ PERF_COUNT_END return 0; } buildLog = (char*) malloc( length ); - if ( buildLog == (char*) nullptr ) - { - return 0; + if (buildLog == (char *)nullptr) { + return 0; } if ( !gpuInfo->mnIsUserCreated ) { @@ -720,10 +725,9 @@ PERF_COUNT_END } fd1 = fopen( "kernel-build.log", "w+" ); - if ( fd1 != nullptr ) - { - fwrite( buildLog, sizeof(char), length, fd1 ); - fclose( fd1 ); + if (fd1 != nullptr) { + fwrite(buildLog, sizeof(char), length, fd1); + fclose(fd1); } free( buildLog ); @@ -766,36 +770,40 @@ PERF_COUNT_START("pixReadFromTiffKernel") l_uint32 *pResult = (l_uint32 *)malloc(w*h * sizeof(l_uint32)); rEnv.mpkKernel = clCreateKernel( rEnv.mpkProgram, "composeRGBPixel", &clStatus ); - CHECK_OPENCL( clStatus, "clCreateKernel composeRGBPixel"); + CHECK_OPENCL(clStatus, "clCreateKernel composeRGBPixel"); //Allocate input and output OCL buffers valuesCl = allocateZeroCopyBuffer(rEnv, tiffdata, w*h, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, &clStatus); outputCl = allocateZeroCopyBuffer(rEnv, pResult, w*h, CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR, &clStatus); //Kernel arguments - clStatus = clSetKernelArg( rEnv.mpkKernel, 0, sizeof(cl_mem), &valuesCl ); + clStatus = clSetKernelArg(rEnv.mpkKernel, 0, sizeof(cl_mem), &valuesCl); CHECK_OPENCL( clStatus, "clSetKernelArg"); - clStatus = clSetKernelArg( rEnv.mpkKernel, 1, sizeof(w), &w ); + clStatus = clSetKernelArg(rEnv.mpkKernel, 1, sizeof(w), &w); CHECK_OPENCL( clStatus, "clSetKernelArg" ); - clStatus = clSetKernelArg( rEnv.mpkKernel, 2, sizeof(h), &h ); + clStatus = clSetKernelArg(rEnv.mpkKernel, 2, sizeof(h), &h); CHECK_OPENCL( clStatus, "clSetKernelArg" ); - clStatus = clSetKernelArg( rEnv.mpkKernel, 3, sizeof(wpl), &wpl ); + clStatus = clSetKernelArg(rEnv.mpkKernel, 3, sizeof(wpl), &wpl); CHECK_OPENCL( clStatus, "clSetKernelArg" ); - clStatus = clSetKernelArg( rEnv.mpkKernel, 4, sizeof(cl_mem), &outputCl ); + clStatus = clSetKernelArg(rEnv.mpkKernel, 4, sizeof(cl_mem), &outputCl); CHECK_OPENCL( clStatus, "clSetKernelArg"); //Kernel enqueue PERF_COUNT_SUB("before") - clStatus = clEnqueueNDRangeKernel( rEnv.mpkCmdQueue, rEnv.mpkKernel, 2, nullptr, globalThreads, localThreads, 0, nullptr, nullptr ); - CHECK_OPENCL( clStatus, "clEnqueueNDRangeKernel" ); - - /* map results back from gpu */ - void *ptr = clEnqueueMapBuffer(rEnv.mpkCmdQueue, outputCl, CL_TRUE, CL_MAP_READ, 0, w*h * sizeof(l_uint32), 0, nullptr, nullptr, &clStatus); - CHECK_OPENCL( clStatus, "clEnqueueMapBuffer outputCl"); - clEnqueueUnmapMemObject(rEnv.mpkCmdQueue, outputCl, ptr, 0, nullptr, nullptr); - - //Sync - clFinish( rEnv.mpkCmdQueue ); +clStatus = + clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, rEnv.mpkKernel, 2, nullptr, + globalThreads, localThreads, 0, nullptr, nullptr); +CHECK_OPENCL(clStatus, "clEnqueueNDRangeKernel"); + +/* map results back from gpu */ +void *ptr = clEnqueueMapBuffer(rEnv.mpkCmdQueue, outputCl, CL_TRUE, CL_MAP_READ, + 0, w * h * sizeof(l_uint32), 0, nullptr, nullptr, + &clStatus); +CHECK_OPENCL(clStatus, "clEnqueueMapBuffer outputCl"); +clEnqueueUnmapMemObject(rEnv.mpkCmdQueue, outputCl, ptr, 0, nullptr, nullptr); + +// Sync +clFinish(rEnv.mpkCmdQueue); PERF_COUNT_SUB("kernel & map") PERF_COUNT_END return pResult; @@ -811,13 +819,13 @@ PIX *pix; PROCNAME("pixReadTiff"); if (!filename) - return (PIX *)ERROR_PTR("filename not defined", procName, nullptr); + return (PIX *)ERROR_PTR("filename not defined", procName, nullptr); if ((fp = fopenReadStream(filename)) == nullptr) - return (PIX *)ERROR_PTR("image file not found", procName, nullptr); + return (PIX *)ERROR_PTR("image file not found", procName, nullptr); if ((pix = pixReadStreamTiffCl(fp, n)) == nullptr) { - fclose(fp); - return (PIX *)ERROR_PTR("pix not read", procName, nullptr); + fclose(fp); + return (PIX *)ERROR_PTR("pix not read", procName, nullptr); } fclose(fp); PERF_COUNT_END @@ -831,13 +839,12 @@ l_int32 fd; PROCNAME("fopenTiff"); - if (!fp) - return (TIFF *)ERROR_PTR("stream not opened", procName, nullptr); + if (!fp) return (TIFF *)ERROR_PTR("stream not opened", procName, nullptr); if (!modestring) - return (TIFF *)ERROR_PTR("modestring not defined", procName, nullptr); + return (TIFF *)ERROR_PTR("modestring not defined", procName, nullptr); if ((fd = fileno(fp)) < 0) - return (TIFF *)ERROR_PTR("invalid file descriptor", procName, nullptr); + return (TIFF *)ERROR_PTR("invalid file descriptor", procName, nullptr); lseek(fd, 0, SEEK_SET); return TIFFFdOpen(fd, "TIFFstream", modestring); @@ -1087,26 +1094,26 @@ size_t *pdatasize) PROCNAME("fopenTiffMemstream"); if (!filename) - return (TIFF *)ERROR_PTR("filename not defined", procName, nullptr); - if (!operation) - return (TIFF *)ERROR_PTR("operation not defined", procName, nullptr); - if (!pdata) - return (TIFF *)ERROR_PTR("&data not defined", procName, nullptr); - if (!pdatasize) - return (TIFF *)ERROR_PTR("&datasize not defined", procName, nullptr); - if (!strcmp(operation, "r") && !strcmp(operation, "w")) - return (TIFF *)ERROR_PTR("operation not 'r' or 'w'}", procName, nullptr); - - if (!strcmp(operation, "r")) - mstream = memstreamCreateForRead(*pdata, *pdatasize); - else - mstream = memstreamCreateForWrite(pdata, pdatasize); + return (TIFF *)ERROR_PTR("filename not defined", procName, nullptr); + if (!operation) + return (TIFF *)ERROR_PTR("operation not defined", procName, nullptr); + if (!pdata) + return (TIFF *)ERROR_PTR("&data not defined", procName, nullptr); + if (!pdatasize) + return (TIFF *)ERROR_PTR("&datasize not defined", procName, nullptr); + if (!strcmp(operation, "r") && !strcmp(operation, "w")) + return (TIFF *)ERROR_PTR("operation not 'r' or 'w'}", procName, + nullptr); + + if (!strcmp(operation, "r")) + mstream = memstreamCreateForRead(*pdata, *pdatasize); + else + mstream = memstreamCreateForWrite(pdata, pdatasize); - return TIFFClientOpen(filename, operation, mstream, - tiffReadCallback, tiffWriteCallback, - tiffSeekCallback, tiffCloseCallback, - tiffSizeCallback, tiffMapCallback, - tiffUnmapCallback); + return TIFFClientOpen(filename, operation, mstream, tiffReadCallback, + tiffWriteCallback, tiffSeekCallback, + tiffCloseCallback, tiffSizeCallback, + tiffMapCallback, tiffUnmapCallback); } @@ -1161,11 +1168,10 @@ TIFF *tif; PROCNAME("pixReadStreamTiff"); - if (!fp) - return (PIX *)ERROR_PTR("stream not defined", procName, nullptr); + if (!fp) return (PIX *)ERROR_PTR("stream not defined", procName, nullptr); if ((tif = fopenTiffCl(fp, "rb")) == nullptr) - return (PIX *)ERROR_PTR("tif not opened", procName, nullptr); + return (PIX *)ERROR_PTR("tif not opened", procName, nullptr); pagefound = FALSE; pix = nullptr; @@ -1173,8 +1179,8 @@ TIFF *tif; if (i == n) { pagefound = TRUE; if ((pix = pixReadFromTiffStreamCl(tif)) == nullptr) { - TIFFCleanup(tif); - return (PIX *)ERROR_PTR("pix not read", procName, nullptr); + TIFFCleanup(tif); + return (PIX *)ERROR_PTR("pix not read", procName, nullptr); } break; } @@ -1254,43 +1260,43 @@ PIXCMAP *cmap; PROCNAME("pixReadFromTiffStream"); - if (!tif) - return (PIX *)ERROR_PTR("tif not defined", procName, nullptr); + if (!tif) return (PIX *)ERROR_PTR("tif not defined", procName, nullptr); TIFFGetFieldDefaulted(tif, TIFFTAG_BITSPERSAMPLE, &bps); TIFFGetFieldDefaulted(tif, TIFFTAG_SAMPLESPERPIXEL, &spp); bpp = bps * spp; if (bpp > 32) - return (PIX *)ERROR_PTR("can't handle bpp > 32", procName, nullptr); + return (PIX *)ERROR_PTR("can't handle bpp > 32", procName, nullptr); if (spp == 1) d = bps; else if (spp == 3 || spp == 4) d = 32; else - return (PIX *)ERROR_PTR("spp not in set {1,3,4}", procName, nullptr); + return (PIX *)ERROR_PTR("spp not in set {1,3,4}", procName, nullptr); TIFFGetField(tif, TIFFTAG_IMAGEWIDTH, &w); TIFFGetField(tif, TIFFTAG_IMAGELENGTH, &h); tiffbpl = TIFFScanlineSize(tif); if ((pix = pixCreate(w, h, d)) == nullptr) - return (PIX *)ERROR_PTR("pix not made", procName, nullptr); + return (PIX *)ERROR_PTR("pix not made", procName, nullptr); data = (l_uint8 *)pixGetData(pix); wpl = pixGetWpl(pix); bpl = 4 * wpl; if (spp == 1) { - if ((linebuf = (l_uint8 *)CALLOC(tiffbpl + 1, sizeof(l_uint8))) == nullptr) - return (PIX *)ERROR_PTR("calloc fail for linebuf", procName, nullptr); - - for (i = 0 ; i < h ; i++) { - if (TIFFReadScanline(tif, linebuf, i, 0) < 0) { - FREE(linebuf); - pixDestroy(&pix); - return (PIX *)ERROR_PTR("line read fail", procName, nullptr); - } - memcpy((char *)data, (char *)linebuf, tiffbpl); - data += bpl; + if ((linebuf = (l_uint8 *)CALLOC(tiffbpl + 1, sizeof(l_uint8))) == + nullptr) + return (PIX *)ERROR_PTR("calloc fail for linebuf", procName, nullptr); + + for (i = 0; i < h; i++) { + if (TIFFReadScanline(tif, linebuf, i, 0) < 0) { + FREE(linebuf); + pixDestroy(&pix); + return (PIX *)ERROR_PTR("line read fail", procName, nullptr); + } + memcpy((char *)data, (char *)linebuf, tiffbpl); + data += bpl; } if (bps <= 8) pixEndianByteSwap(pix); @@ -1331,9 +1337,9 @@ PIXCMAP *cmap; pixSetInputFormat(pix, comptype); if (TIFFGetField(tif, TIFFTAG_COLORMAP, &redmap, &greenmap, &bluemap)) { - if ((cmap = pixcmapCreate(bps)) == nullptr) { - pixDestroy(&pix); - return (PIX *)ERROR_PTR("cmap not made", procName, nullptr); + if ((cmap = pixcmapCreate(bps)) == nullptr) { + pixDestroy(&pix); + return (PIX *)ERROR_PTR("cmap not made", procName, nullptr); } ncolors = 1 << bps; for (i = 0; i < ncolors; i++) @@ -1400,24 +1406,12 @@ pixDilateCL_55(l_int32 wpl, l_int32 h) 1, sizeof(cl_mem), &pixdCLBuffer); - status = clSetKernelArg(rEnv.mpkKernel, - 2, - sizeof(wpl), - &wpl); - status = clSetKernelArg(rEnv.mpkKernel, - 3, - sizeof(h), - &h); - - status = clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, - rEnv.mpkKernel, - 2, - nullptr, - globalThreads, - localThreads, - 0, - nullptr, - nullptr); + status = clSetKernelArg(rEnv.mpkKernel, 2, sizeof(wpl), &wpl); + status = clSetKernelArg(rEnv.mpkKernel, 3, sizeof(h), &h); + + status = clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, rEnv.mpkKernel, 2, + nullptr, globalThreads, localThreads, 0, + nullptr, nullptr); //Swap source and dest buffers pixtemp = pixsCLBuffer; @@ -1443,23 +1437,11 @@ pixDilateCL_55(l_int32 wpl, l_int32 h) 1, sizeof(cl_mem), &pixdCLBuffer); - status = clSetKernelArg(rEnv.mpkKernel, - 2, - sizeof(wpl), - &wpl); - status = clSetKernelArg(rEnv.mpkKernel, - 3, - sizeof(h), - &h); - status = clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, - rEnv.mpkKernel, - 2, - nullptr, - globalThreads, - localThreads, - 0, - nullptr, - nullptr); + status = clSetKernelArg(rEnv.mpkKernel, 2, sizeof(wpl), &wpl); + status = clSetKernelArg(rEnv.mpkKernel, 3, sizeof(h), &h); + status = clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, rEnv.mpkKernel, 2, + nullptr, globalThreads, localThreads, 0, + nullptr, nullptr); return status; } @@ -1496,24 +1478,12 @@ pixErodeCL_55(l_int32 wpl, l_int32 h) 1, sizeof(cl_mem), &pixdCLBuffer); - status = clSetKernelArg(rEnv.mpkKernel, - 2, - sizeof(wpl), - &wpl); - status = clSetKernelArg(rEnv.mpkKernel, - 3, - sizeof(h), - &h); - - status = clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, - rEnv.mpkKernel, - 2, - nullptr, - globalThreads, - localThreads, - 0, - nullptr, - nullptr); + status = clSetKernelArg(rEnv.mpkKernel, 2, sizeof(wpl), &wpl); + status = clSetKernelArg(rEnv.mpkKernel, 3, sizeof(h), &h); + + status = clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, rEnv.mpkKernel, 2, + nullptr, globalThreads, localThreads, 0, + nullptr, nullptr); //Swap source and dest buffers pixtemp = pixsCLBuffer; @@ -1539,31 +1509,13 @@ pixErodeCL_55(l_int32 wpl, l_int32 h) 1, sizeof(cl_mem), &pixdCLBuffer); - status = clSetKernelArg(rEnv.mpkKernel, - 2, - sizeof(wpl), - &wpl); - status = clSetKernelArg(rEnv.mpkKernel, - 3, - sizeof(h), - &h); - status = clSetKernelArg(rEnv.mpkKernel, - 4, - sizeof(fwmask), - &fwmask); - status = clSetKernelArg(rEnv.mpkKernel, - 5, - sizeof(lwmask), - &lwmask); - status = clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, - rEnv.mpkKernel, - 2, - nullptr, - globalThreads, - localThreads, - 0, - nullptr, - nullptr); + status = clSetKernelArg(rEnv.mpkKernel, 2, sizeof(wpl), &wpl); + status = clSetKernelArg(rEnv.mpkKernel, 3, sizeof(h), &h); + status = clSetKernelArg(rEnv.mpkKernel, 4, sizeof(fwmask), &fwmask); + status = clSetKernelArg(rEnv.mpkKernel, 5, sizeof(lwmask), &lwmask); + status = clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, rEnv.mpkKernel, 2, + nullptr, globalThreads, localThreads, 0, + nullptr, nullptr); return status; } @@ -1613,12 +1565,11 @@ pixDilateCL(l_int32 hsize, l_int32 vsize, l_int32 wpl, l_int32 h) status = clSetKernelArg(rEnv.mpkKernel, 1, sizeof(cl_mem), &pixdCLBuffer); status = clSetKernelArg(rEnv.mpkKernel, 2, sizeof(xp), &xp); status = clSetKernelArg(rEnv.mpkKernel, 3, sizeof(xn), &xn); - status = - clSetKernelArg(rEnv.mpkKernel, 4, sizeof(wpl), &wpl); + status = clSetKernelArg(rEnv.mpkKernel, 4, sizeof(wpl), &wpl); status = clSetKernelArg(rEnv.mpkKernel, 5, sizeof(h), &h); - status = - clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, rEnv.mpkKernel, 2, nullptr, - globalThreads, localThreads, 0, nullptr, nullptr); + status = clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, rEnv.mpkKernel, 2, + nullptr, globalThreads, localThreads, 0, + nullptr, nullptr); if (yp > 0 || yn > 0) { pixtemp = pixsCLBuffer; @@ -1637,14 +1588,12 @@ pixDilateCL(l_int32 hsize, l_int32 vsize, l_int32 wpl, l_int32 h) status = clSetKernelArg(rEnv.mpkKernel, 0, sizeof(cl_mem), &pixsCLBuffer); status = clSetKernelArg(rEnv.mpkKernel, 1, sizeof(cl_mem), &pixdCLBuffer); status = clSetKernelArg(rEnv.mpkKernel, 2, sizeof(xp), &xp); - status = - clSetKernelArg(rEnv.mpkKernel, 3, sizeof(wpl), &wpl); + status = clSetKernelArg(rEnv.mpkKernel, 3, sizeof(wpl), &wpl); status = clSetKernelArg(rEnv.mpkKernel, 4, sizeof(h), &h); - status = clSetKernelArg(rEnv.mpkKernel, 5, sizeof(isEven), - &isEven); - status = - clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, rEnv.mpkKernel, 2, nullptr, - globalThreads, localThreads, 0, nullptr, nullptr); + status = clSetKernelArg(rEnv.mpkKernel, 5, sizeof(isEven), &isEven); + status = clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, rEnv.mpkKernel, 2, + nullptr, globalThreads, localThreads, 0, + nullptr, nullptr); if (yp > 0 || yn > 0) { pixtemp = pixsCLBuffer; @@ -1666,31 +1615,13 @@ pixDilateCL(l_int32 hsize, l_int32 vsize, l_int32 wpl, l_int32 h) 1, sizeof(cl_mem), &pixdCLBuffer); - status = clSetKernelArg(rEnv.mpkKernel, - 2, - sizeof(yp), - &yp); - status = clSetKernelArg(rEnv.mpkKernel, - 3, - sizeof(wpl), - &wpl); - status = clSetKernelArg(rEnv.mpkKernel, - 4, - sizeof(h), - &h); - status = clSetKernelArg(rEnv.mpkKernel, - 5, - sizeof(yn), - &yn); - status = clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, - rEnv.mpkKernel, - 2, - nullptr, - globalThreads, - localThreads, - 0, - nullptr, - nullptr); + status = clSetKernelArg(rEnv.mpkKernel, 2, sizeof(yp), &yp); + status = clSetKernelArg(rEnv.mpkKernel, 3, sizeof(wpl), &wpl); + status = clSetKernelArg(rEnv.mpkKernel, 4, sizeof(h), &h); + status = clSetKernelArg(rEnv.mpkKernel, 5, sizeof(yn), &yn); + status = clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, rEnv.mpkKernel, 2, + nullptr, globalThreads, localThreads, 0, + nullptr, nullptr); } return status; @@ -1743,14 +1674,13 @@ cl_int pixErodeCL(l_int32 hsize, l_int32 vsize, l_uint32 wpl, l_uint32 h) { status = clSetKernelArg(rEnv.mpkKernel, 3, sizeof(xn), &xn); status = clSetKernelArg(rEnv.mpkKernel, 4, sizeof(wpl), &wpl); status = clSetKernelArg(rEnv.mpkKernel, 5, sizeof(h), &h); - status = clSetKernelArg(rEnv.mpkKernel, 6, sizeof(isAsymmetric), - &isAsymmetric); - status = clSetKernelArg(rEnv.mpkKernel, 7, sizeof(rwmask), - &rwmask); - status = clSetKernelArg(rEnv.mpkKernel, 8, sizeof(lwmask), - &lwmask); - status = clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, rEnv.mpkKernel, 2, nullptr, - globalThreads, localThreads, 0, nullptr, nullptr); + status = + clSetKernelArg(rEnv.mpkKernel, 6, sizeof(isAsymmetric), &isAsymmetric); + status = clSetKernelArg(rEnv.mpkKernel, 7, sizeof(rwmask), &rwmask); + status = clSetKernelArg(rEnv.mpkKernel, 8, sizeof(lwmask), &lwmask); + status = clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, rEnv.mpkKernel, 2, + nullptr, globalThreads, localThreads, 0, + nullptr, nullptr); if (yp > 0 || yn > 0) { pixtemp = pixsCLBuffer; @@ -1767,16 +1697,14 @@ cl_int pixErodeCL(l_int32 hsize, l_int32 vsize, l_uint32 wpl, l_uint32 h) { status = clSetKernelArg(rEnv.mpkKernel, 2, sizeof(xp), &xp); status = clSetKernelArg(rEnv.mpkKernel, 3, sizeof(wpl), &wpl); status = clSetKernelArg(rEnv.mpkKernel, 4, sizeof(h), &h); - status = clSetKernelArg(rEnv.mpkKernel, 5, sizeof(isAsymmetric), - &isAsymmetric); - status = clSetKernelArg(rEnv.mpkKernel, 6, sizeof(rwmask), - &rwmask); - status = clSetKernelArg(rEnv.mpkKernel, 7, sizeof(lwmask), - &lwmask); - status = clSetKernelArg(rEnv.mpkKernel, 8, sizeof(isEven), - &isEven); - status = clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, rEnv.mpkKernel, 2, nullptr, - globalThreads, localThreads, 0, nullptr, nullptr); + status = + clSetKernelArg(rEnv.mpkKernel, 5, sizeof(isAsymmetric), &isAsymmetric); + status = clSetKernelArg(rEnv.mpkKernel, 6, sizeof(rwmask), &rwmask); + status = clSetKernelArg(rEnv.mpkKernel, 7, sizeof(lwmask), &lwmask); + status = clSetKernelArg(rEnv.mpkKernel, 8, sizeof(isEven), &isEven); + status = clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, rEnv.mpkKernel, 2, + nullptr, globalThreads, localThreads, 0, + nullptr, nullptr); if (yp > 0 || yn > 0) { pixtemp = pixsCLBuffer; @@ -1795,11 +1723,12 @@ cl_int pixErodeCL(l_int32 hsize, l_int32 vsize, l_uint32 wpl, l_uint32 h) { status = clSetKernelArg(rEnv.mpkKernel, 2, sizeof(yp), &yp); status = clSetKernelArg(rEnv.mpkKernel, 3, sizeof(wpl), &wpl); status = clSetKernelArg(rEnv.mpkKernel, 4, sizeof(h), &h); - status = clSetKernelArg(rEnv.mpkKernel, 5, sizeof(isAsymmetric), - &isAsymmetric); + status = + clSetKernelArg(rEnv.mpkKernel, 5, sizeof(isAsymmetric), &isAsymmetric); status = clSetKernelArg(rEnv.mpkKernel, 6, sizeof(yn), &yn); - status = clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, rEnv.mpkKernel, 2, nullptr, - globalThreads, localThreads, 0, nullptr, nullptr); + status = clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, rEnv.mpkKernel, 2, + nullptr, globalThreads, localThreads, 0, + nullptr, nullptr); } return status; @@ -1948,23 +1877,11 @@ pixORCL_work(l_uint32 wpl, l_uint32 h, cl_mem buffer1, cl_mem buffer2, cl_mem ou 2, sizeof(cl_mem), &outbuffer); - status = clSetKernelArg(rEnv.mpkKernel, - 3, - sizeof(wpl), - &wpl); - status = clSetKernelArg(rEnv.mpkKernel, - 4, - sizeof(h), - &h); - status = clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, - rEnv.mpkKernel, - 2, - nullptr, - globalThreads, - localThreads, - 0, - nullptr, - nullptr); + status = clSetKernelArg(rEnv.mpkKernel, 3, sizeof(wpl), &wpl); + status = clSetKernelArg(rEnv.mpkKernel, 4, sizeof(h), &h); + status = clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, rEnv.mpkKernel, 2, + nullptr, globalThreads, localThreads, 0, + nullptr, nullptr); return status; } @@ -1999,87 +1916,50 @@ pixANDCL_work(l_uint32 wpl, l_uint32 h, cl_mem buffer1, cl_mem buffer2, cl_mem o 2, sizeof(cl_mem), &outbuffer); - status = clSetKernelArg(rEnv.mpkKernel, - 3, - sizeof(wpl), - &wpl); - status = clSetKernelArg(rEnv.mpkKernel, - 4, - sizeof(h), - &h); - status = clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, - rEnv.mpkKernel, - 2, - nullptr, - globalThreads, - localThreads, - 0, - nullptr, - nullptr); + status = clSetKernelArg(rEnv.mpkKernel, 3, sizeof(wpl), &wpl); + status = clSetKernelArg(rEnv.mpkKernel, 4, sizeof(h), &h); + status = clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, rEnv.mpkKernel, 2, + nullptr, globalThreads, localThreads, 0, + nullptr, nullptr); return status; } //output = buffer1 & ~(buffer2) -cl_int -pixSubtractCL_work(l_uint32 wpl, l_uint32 h, cl_mem buffer1, cl_mem buffer2, cl_mem outBuffer = nullptr) -{ - cl_int status; - size_t globalThreads[2]; - int gsize; - size_t localThreads[] = {GROUPSIZE_X, GROUPSIZE_Y}; +cl_int pixSubtractCL_work(l_uint32 wpl, l_uint32 h, cl_mem buffer1, + cl_mem buffer2, cl_mem outBuffer = nullptr) { + cl_int status; + size_t globalThreads[2]; + int gsize; + size_t localThreads[] = {GROUPSIZE_X, GROUPSIZE_Y}; - gsize = (wpl + GROUPSIZE_X - 1)/ GROUPSIZE_X * GROUPSIZE_X; - globalThreads[0] = gsize; - gsize = (h + GROUPSIZE_Y - 1)/ GROUPSIZE_Y * GROUPSIZE_Y; - globalThreads[1] = gsize; + gsize = (wpl + GROUPSIZE_X - 1) / GROUPSIZE_X * GROUPSIZE_X; + globalThreads[0] = gsize; + gsize = (h + GROUPSIZE_Y - 1) / GROUPSIZE_Y * GROUPSIZE_Y; + globalThreads[1] = gsize; - if (outBuffer != nullptr) - { - rEnv.mpkKernel = clCreateKernel( rEnv.mpkProgram, "pixSubtract", &status ); - CHECK_OPENCL(status, "clCreateKernel pixSubtract"); - } - else - { - rEnv.mpkKernel = clCreateKernel( rEnv.mpkProgram, "pixSubtract_inplace", &status ); - CHECK_OPENCL(status, "clCreateKernel pixSubtract_inplace"); - } + if (outBuffer != nullptr) { + rEnv.mpkKernel = clCreateKernel(rEnv.mpkProgram, "pixSubtract", &status); + CHECK_OPENCL(status, "clCreateKernel pixSubtract"); + } else { + rEnv.mpkKernel = + clCreateKernel(rEnv.mpkProgram, "pixSubtract_inplace", &status); + CHECK_OPENCL(status, "clCreateKernel pixSubtract_inplace"); + } - // Enqueue a kernel run call. - status = clSetKernelArg(rEnv.mpkKernel, - 0, - sizeof(cl_mem), - &buffer1); - status = clSetKernelArg(rEnv.mpkKernel, - 1, - sizeof(cl_mem), - &buffer2); - status = clSetKernelArg(rEnv.mpkKernel, - 2, - sizeof(wpl), - &wpl); - status = clSetKernelArg(rEnv.mpkKernel, - 3, - sizeof(h), - &h); - if (outBuffer != nullptr) - { - status = clSetKernelArg(rEnv.mpkKernel, - 4, - sizeof(cl_mem), - &outBuffer); - } - status = clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, - rEnv.mpkKernel, - 2, - nullptr, - globalThreads, - localThreads, - 0, - nullptr, - nullptr); + // Enqueue a kernel run call. + status = clSetKernelArg(rEnv.mpkKernel, 0, sizeof(cl_mem), &buffer1); + status = clSetKernelArg(rEnv.mpkKernel, 1, sizeof(cl_mem), &buffer2); + status = clSetKernelArg(rEnv.mpkKernel, 2, sizeof(wpl), &wpl); + status = clSetKernelArg(rEnv.mpkKernel, 3, sizeof(h), &h); + if (outBuffer != nullptr) { + status = clSetKernelArg(rEnv.mpkKernel, 4, sizeof(cl_mem), &outBuffer); + } + status = + clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, rEnv.mpkKernel, 2, nullptr, + globalThreads, localThreads, 0, nullptr, nullptr); - return status; + return status; } // OpenCL implementation of Subtract pix @@ -2260,8 +2140,7 @@ int OpenclDevice::HistogramRectOCL(unsigned char *imageData, // using a garlic bus memory type cl_mem imageBuffer = clCreateBuffer( histKern.mpkContext, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, - width * height * bytes_per_pixel * sizeof(char), imageData, - &clStatus); + width * height * bytes_per_pixel * sizeof(char), imageData, &clStatus); CHECK_OPENCL(clStatus, "clCreateBuffer imageBuffer"); /* setup work group size parameters */ @@ -2284,8 +2163,8 @@ int OpenclDevice::HistogramRectOCL(unsigned char *imageData, cl_mem histogramBuffer = clCreateBuffer( histKern.mpkContext, CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR, - kHistogramSize * bytes_per_pixel * sizeof(int), - histogramAllChannels, &clStatus); + kHistogramSize * bytes_per_pixel * sizeof(int), histogramAllChannels, + &clStatus); CHECK_OPENCL(clStatus, "clCreateBuffer histogramBuffer"); /* intermediate histogram buffer */ @@ -2327,48 +2206,53 @@ int OpenclDevice::HistogramRectOCL(unsigned char *imageData, void *ptr; //Initialize tmpHistogramBuffer buffer - ptr = clEnqueueMapBuffer(histKern.mpkCmdQueue, tmpHistogramBuffer, CL_TRUE, CL_MAP_WRITE, 0, tmpHistogramBins*sizeof(cl_uint), 0, nullptr, nullptr, &clStatus); + ptr = clEnqueueMapBuffer( + histKern.mpkCmdQueue, tmpHistogramBuffer, CL_TRUE, CL_MAP_WRITE, 0, + tmpHistogramBins * sizeof(cl_uint), 0, nullptr, nullptr, &clStatus); CHECK_OPENCL( clStatus, "clEnqueueMapBuffer tmpHistogramBuffer"); memset(ptr, 0, tmpHistogramBins*sizeof(cl_uint)); - clEnqueueUnmapMemObject(histKern.mpkCmdQueue, tmpHistogramBuffer, ptr, 0, nullptr, nullptr); + clEnqueueUnmapMemObject(histKern.mpkCmdQueue, tmpHistogramBuffer, ptr, 0, + nullptr, nullptr); /* set kernel 1 arguments */ - clStatus = clSetKernelArg( histKern.mpkKernel, 0, sizeof(cl_mem), &imageBuffer ); + clStatus = + clSetKernelArg(histKern.mpkKernel, 0, sizeof(cl_mem), &imageBuffer); CHECK_OPENCL( clStatus, "clSetKernelArg imageBuffer"); cl_uint numPixels = width*height; - clStatus = clSetKernelArg( histKern.mpkKernel, 1, sizeof(cl_uint), &numPixels ); + clStatus = + clSetKernelArg(histKern.mpkKernel, 1, sizeof(cl_uint), &numPixels); CHECK_OPENCL( clStatus, "clSetKernelArg numPixels" ); - clStatus = clSetKernelArg( histKern.mpkKernel, 2, sizeof(cl_mem), &tmpHistogramBuffer ); + clStatus = clSetKernelArg(histKern.mpkKernel, 2, sizeof(cl_mem), + &tmpHistogramBuffer); CHECK_OPENCL( clStatus, "clSetKernelArg tmpHistogramBuffer"); /* set kernel 2 arguments */ int n = numThreads/bytes_per_pixel; - clStatus = clSetKernelArg( histRedKern.mpkKernel, 0, sizeof(cl_int), &n ); + clStatus = clSetKernelArg(histRedKern.mpkKernel, 0, sizeof(cl_int), &n); CHECK_OPENCL( clStatus, "clSetKernelArg imageBuffer"); - clStatus = clSetKernelArg( histRedKern.mpkKernel, 1, sizeof(cl_mem), &tmpHistogramBuffer ); + clStatus = clSetKernelArg(histRedKern.mpkKernel, 1, sizeof(cl_mem), + &tmpHistogramBuffer); CHECK_OPENCL( clStatus, "clSetKernelArg tmpHistogramBuffer"); - clStatus = clSetKernelArg( histRedKern.mpkKernel, 2, sizeof(cl_mem), &histogramBuffer ); + clStatus = clSetKernelArg(histRedKern.mpkKernel, 2, sizeof(cl_mem), + &histogramBuffer); CHECK_OPENCL( clStatus, "clSetKernelArg histogramBuffer"); /* launch histogram */ PERF_COUNT_SUB("before") - clStatus = clEnqueueNDRangeKernel( - histKern.mpkCmdQueue, - histKern.mpkKernel, - 1, nullptr, global_work_size, local_work_size, - 0, nullptr, nullptr ); - CHECK_OPENCL( clStatus, "clEnqueueNDRangeKernel kernel_HistogramRectAllChannels" ); - clFinish( histKern.mpkCmdQueue ); - if (clStatus != 0) { - retVal = -1; +clStatus = clEnqueueNDRangeKernel(histKern.mpkCmdQueue, histKern.mpkKernel, 1, + nullptr, global_work_size, local_work_size, 0, + nullptr, nullptr); +CHECK_OPENCL(clStatus, + "clEnqueueNDRangeKernel kernel_HistogramRectAllChannels"); +clFinish(histKern.mpkCmdQueue); +if (clStatus != 0) { + retVal = -1; } /* launch histogram */ clStatus = clEnqueueNDRangeKernel( - histRedKern.mpkCmdQueue, - histRedKern.mpkKernel, - 1, nullptr, red_global_work_size, local_work_size, - 0, nullptr, nullptr ); + histRedKern.mpkCmdQueue, histRedKern.mpkKernel, 1, nullptr, + red_global_work_size, local_work_size, 0, nullptr, nullptr); CHECK_OPENCL( clStatus, "clEnqueueNDRangeKernel kernel_HistogramRectAllChannelsReduction" ); clFinish( histRedKern.mpkCmdQueue ); if (clStatus != 0) { @@ -2377,12 +2261,16 @@ PERF_COUNT_SUB("before") PERF_COUNT_SUB("redKernel") /* map results back from gpu */ - ptr = clEnqueueMapBuffer(histRedKern.mpkCmdQueue, histogramBuffer, CL_TRUE, CL_MAP_READ, 0, kHistogramSize*bytes_per_pixel*sizeof(int), 0, nullptr, nullptr, &clStatus); + ptr = clEnqueueMapBuffer(histRedKern.mpkCmdQueue, histogramBuffer, CL_TRUE, + CL_MAP_READ, 0, + kHistogramSize * bytes_per_pixel * sizeof(int), 0, + nullptr, nullptr, &clStatus); CHECK_OPENCL( clStatus, "clEnqueueMapBuffer histogramBuffer"); if (clStatus != 0) { retVal = -1; } - clEnqueueUnmapMemObject(histRedKern.mpkCmdQueue, histogramBuffer, ptr, 0, nullptr, nullptr); + clEnqueueUnmapMemObject(histRedKern.mpkCmdQueue, histogramBuffer, ptr, 0, + nullptr, nullptr); clReleaseMemObject(histogramBuffer); clReleaseMemObject(imageBuffer); @@ -2398,10 +2286,9 @@ return retVal; ************************************************************************/ int OpenclDevice::ThresholdRectToPixOCL(unsigned char *imageData, int bytes_per_pixel, int bytes_per_line, - int *thresholds, - int *hi_values, Pix **pix, - int height, int width, int top, - int left) { + int *thresholds, int *hi_values, + Pix **pix, int height, int width, + int top, int left) { PERF_COUNT_START("ThresholdRectToPixOCL") int retVal = 0; /* create pix result buffer */ @@ -2432,10 +2319,9 @@ int OpenclDevice::ThresholdRectToPixOCL(unsigned char *imageData, // coherent which we don't need. // faster option would be to allocate initial image buffer // using a garlic bus memory type - cl_mem imageBuffer = - clCreateBuffer(rEnv.mpkContext, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, - width * height * bytes_per_pixel * sizeof(char), - imageData, &clStatus); + cl_mem imageBuffer = clCreateBuffer( + rEnv.mpkContext, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, + width * height * bytes_per_pixel * sizeof(char), imageData, &clStatus); CHECK_OPENCL(clStatus, "clCreateBuffer imageBuffer"); /* map pix as write only */ @@ -2445,13 +2331,13 @@ int OpenclDevice::ThresholdRectToPixOCL(unsigned char *imageData, CHECK_OPENCL(clStatus, "clCreateBuffer pix"); /* map thresholds and hi_values */ - cl_mem thresholdsBuffer = clCreateBuffer( - rEnv.mpkContext, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, - bytes_per_pixel * sizeof(int), thresholds, &clStatus); + cl_mem thresholdsBuffer = + clCreateBuffer(rEnv.mpkContext, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, + bytes_per_pixel * sizeof(int), thresholds, &clStatus); CHECK_OPENCL(clStatus, "clCreateBuffer thresholdBuffer"); - cl_mem hiValuesBuffer = clCreateBuffer( - rEnv.mpkContext, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, - bytes_per_pixel * sizeof(int), hi_values, &clStatus); + cl_mem hiValuesBuffer = + clCreateBuffer(rEnv.mpkContext, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, + bytes_per_pixel * sizeof(int), hi_values, &clStatus); CHECK_OPENCL(clStatus, "clCreateBuffer hiValuesBuffer"); /* compile kernel */ @@ -2466,8 +2352,7 @@ int OpenclDevice::ThresholdRectToPixOCL(unsigned char *imageData, } /* set kernel arguments */ - clStatus = - clSetKernelArg(rEnv.mpkKernel, 0, sizeof(cl_mem), &imageBuffer); + clStatus = clSetKernelArg(rEnv.mpkKernel, 0, sizeof(cl_mem), &imageBuffer); CHECK_OPENCL(clStatus, "clSetKernelArg imageBuffer"); cl_uint numPixels = width * height; clStatus = clSetKernelArg(rEnv.mpkKernel, 1, sizeof(int), &height); @@ -2476,21 +2361,19 @@ int OpenclDevice::ThresholdRectToPixOCL(unsigned char *imageData, CHECK_OPENCL(clStatus, "clSetKernelArg width"); clStatus = clSetKernelArg(rEnv.mpkKernel, 3, sizeof(int), &wpl); CHECK_OPENCL(clStatus, "clSetKernelArg wpl"); - clStatus = clSetKernelArg(rEnv.mpkKernel, 4, sizeof(cl_mem), - &thresholdsBuffer); + clStatus = + clSetKernelArg(rEnv.mpkKernel, 4, sizeof(cl_mem), &thresholdsBuffer); CHECK_OPENCL(clStatus, "clSetKernelArg thresholdsBuffer"); - clStatus = clSetKernelArg(rEnv.mpkKernel, 5, sizeof(cl_mem), - &hiValuesBuffer); + clStatus = clSetKernelArg(rEnv.mpkKernel, 5, sizeof(cl_mem), &hiValuesBuffer); CHECK_OPENCL(clStatus, "clSetKernelArg hiValuesBuffer"); - clStatus = - clSetKernelArg(rEnv.mpkKernel, 6, sizeof(cl_mem), &pixThBuffer); + clStatus = clSetKernelArg(rEnv.mpkKernel, 6, sizeof(cl_mem), &pixThBuffer); CHECK_OPENCL(clStatus, "clSetKernelArg pixThBuffer"); /* launch kernel & wait */ PERF_COUNT_SUB("before") - clStatus = - clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, rEnv.mpkKernel, 1, nullptr, - global_work_size, local_work_size, 0, nullptr, nullptr); + clStatus = clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, rEnv.mpkKernel, 1, + nullptr, global_work_size, local_work_size, + 0, nullptr, nullptr); CHECK_OPENCL(clStatus, "clEnqueueNDRangeKernel kernel_ThresholdRectToPix"); clFinish(rEnv.mpkCmdQueue); PERF_COUNT_SUB("kernel") @@ -2503,7 +2386,8 @@ int OpenclDevice::ThresholdRectToPixOCL(unsigned char *imageData, clEnqueueMapBuffer(rEnv.mpkCmdQueue, pixThBuffer, CL_TRUE, CL_MAP_READ, 0, pixSize, 0, nullptr, nullptr, &clStatus); CHECK_OPENCL(clStatus, "clEnqueueMapBuffer histogramBuffer"); - clEnqueueUnmapMemObject(rEnv.mpkCmdQueue, pixThBuffer, ptr, 0, nullptr, nullptr); + clEnqueueUnmapMemObject(rEnv.mpkCmdQueue, pixThBuffer, ptr, 0, nullptr, + nullptr); clReleaseMemObject(imageBuffer); clReleaseMemObject(thresholdsBuffer); @@ -2644,7 +2528,8 @@ double composeRGBPixelMicroBench( GPUEnv *env, TessScoreEvaluationInputData inpu OpenclDevice::gpuEnv = *env; int wpl = pixGetWpl(input.pix); - OpenclDevice::pixReadFromTiffKernel(tiffdata, input.width, input.height, wpl, nullptr); + OpenclDevice::pixReadFromTiffKernel(tiffdata, input.width, input.height, + wpl, nullptr); #if ON_WINDOWS QueryPerformanceCounter(&time_funct_end); time = (time_funct_end.QuadPart-time_funct_start.QuadPart)/(double)(freq.QuadPart); @@ -2796,13 +2681,13 @@ void ThresholdRectToPix_Native(const unsigned char* imagedata, int height = pixGetHeight(*pix); *pix = pixCreate(width, height, 1); - uint32_t* pixdata = pixGetData(*pix); + uint32_t *pixdata = pixGetData(*pix); int wpl = pixGetWpl(*pix); const unsigned char* srcdata = imagedata + top * bytes_per_line + left * bytes_per_pixel; for (int y = 0; y < height; ++y) { - const uint8_t* linedata = srcdata; - uint32_t* pixline = pixdata + y * wpl; + const uint8_t *linedata = srcdata; + uint32_t *pixline = pixdata + y * wpl; for (int x = 0; x < width; ++x, linedata += bytes_per_pixel) { bool white_result = true; for (int ch = 0; ch < bytes_per_pixel; ++ch) { @@ -2953,7 +2838,10 @@ double getLineMasksMorphMicroBench( GPUEnv *env, TessScoreEvaluationInputData in OpenclDevice::gpuEnv = *env; OpenclDevice::initMorphCLAllocations(wpl, input.height, input.pix); Pix *pix_vline = nullptr, *pix_hline = nullptr, *pix_closed = nullptr; - OpenclDevice::pixGetLinesCL(nullptr, input.pix, &pix_vline, &pix_hline, &pix_closed, true, closing_brick, closing_brick, max_line_width, max_line_width, min_line_length, min_line_length); + OpenclDevice::pixGetLinesCL( + nullptr, input.pix, &pix_vline, &pix_hline, &pix_closed, true, + closing_brick, closing_brick, max_line_width, max_line_width, + min_line_length, min_line_length); OpenclDevice::releaseMorphCLBuffers(); @@ -2978,8 +2866,10 @@ double getLineMasksMorphMicroBench( GPUEnv *env, TessScoreEvaluationInputData in // native serial code Pix *src_pix = input.pix; - Pix *pix_closed = pixCloseBrick(nullptr, src_pix, closing_brick, closing_brick); - Pix *pix_solid = pixOpenBrick(nullptr, pix_closed, max_line_width, max_line_width); + Pix *pix_closed = + pixCloseBrick(nullptr, src_pix, closing_brick, closing_brick); + Pix *pix_solid = + pixOpenBrick(nullptr, pix_closed, max_line_width, max_line_width); Pix *pix_hollow = pixSubtract(nullptr, pix_closed, pix_solid); pixDestroy(&pix_solid); Pix *pix_vline = pixOpenBrick(nullptr, pix_hollow, 1, min_line_length); @@ -3112,9 +3002,8 @@ ds_device OpenclDevice::getDeviceSelection( ) { // PERF_COUNT_SUB("populateTessScoreEvaluationInputData") // perform evaluations unsigned int numUpdates; - status = - profileDevices(profile, DS_EVALUATE_ALL, evaluateScoreForDevice, - &input, &numUpdates); + status = profileDevices(profile, DS_EVALUATE_ALL, + evaluateScoreForDevice, &input, &numUpdates); PERF_COUNT_SUB("profileDevices") // write scores to file if (status == DS_SUCCESS) { @@ -3301,11 +3190,9 @@ Pix *OpenclDevice::pixConvertRGBToGrayOCL(Pix *srcPix, // 32-bit source CHECK_OPENCL(clStatus, "clCreateKernel kernel_RGBToGray"); /* set kernel arguments */ - clStatus = - clSetKernelArg(kEnv.mpkKernel, 0, sizeof(cl_mem), &srcBuffer); + clStatus = clSetKernelArg(kEnv.mpkKernel, 0, sizeof(cl_mem), &srcBuffer); CHECK_OPENCL(clStatus, "clSetKernelArg srcBuffer"); - clStatus = - clSetKernelArg(kEnv.mpkKernel, 1, sizeof(cl_mem), &dstBuffer); + clStatus = clSetKernelArg(kEnv.mpkKernel, 1, sizeof(cl_mem), &dstBuffer); CHECK_OPENCL(clStatus, "clSetKernelArg dstBuffer"); clStatus = clSetKernelArg(kEnv.mpkKernel, 2, sizeof(int), &srcWPL); CHECK_OPENCL(clStatus, "clSetKernelArg srcWPL"); @@ -3324,9 +3211,9 @@ Pix *OpenclDevice::pixConvertRGBToGrayOCL(Pix *srcPix, // 32-bit source /* launch kernel & wait */ PERF_COUNT_SUB("before") - clStatus = - clEnqueueNDRangeKernel(kEnv.mpkCmdQueue, kEnv.mpkKernel, 1, nullptr, - global_work_size, local_work_size, 0, nullptr, nullptr); + clStatus = clEnqueueNDRangeKernel(kEnv.mpkCmdQueue, kEnv.mpkKernel, 1, + nullptr, global_work_size, local_work_size, + 0, nullptr, nullptr); CHECK_OPENCL(clStatus, "clEnqueueNDRangeKernel kernel_RGBToGray"); clFinish(kEnv.mpkCmdQueue); PERF_COUNT_SUB("kernel") @@ -3336,7 +3223,8 @@ Pix *OpenclDevice::pixConvertRGBToGrayOCL(Pix *srcPix, // 32-bit source clEnqueueMapBuffer(kEnv.mpkCmdQueue, dstBuffer, CL_TRUE, CL_MAP_READ, 0, dstSize, 0, nullptr, nullptr, &clStatus); CHECK_OPENCL(clStatus, "clEnqueueMapBuffer dstBuffer"); - clEnqueueUnmapMemObject(rEnv.mpkCmdQueue, dstBuffer, ptr, 0, nullptr, nullptr); + clEnqueueUnmapMemObject(rEnv.mpkCmdQueue, dstBuffer, ptr, 0, nullptr, + nullptr); #if 0 // validate: compute on cpu diff --git a/opencl/openclwrapper.h b/opencl/openclwrapper.h index 0fbaf89d7d..e09e371dbb 100644 --- a/opencl/openclwrapper.h +++ b/opencl/openclwrapper.h @@ -298,15 +298,14 @@ class OpenclDevice inline static int AddKernelConfig( int kCount, const char *kName ); /* for binarization */ - static int HistogramRectOCL(unsigned char *imagedata, - int bytes_per_pixel, int bytes_per_line, - int left, int top, int width, int height, - int kHistogramSize, int *histogramAllChannels); + static int HistogramRectOCL(unsigned char *imagedata, int bytes_per_pixel, + int bytes_per_line, int left, int top, + int width, int height, int kHistogramSize, + int *histogramAllChannels); static int ThresholdRectToPixOCL(unsigned char *imagedata, int bytes_per_pixel, int bytes_per_line, - int *thresholds, - int *hi_values, Pix **pix, + int *thresholds, int *hi_values, Pix **pix, int rect_height, int rect_width, int rect_top, int rect_left); diff --git a/textord/blkocc.h b/textord/blkocc.h index d80afe25f3..f27bb9a5a5 100644 --- a/textord/blkocc.h +++ b/textord/blkocc.h @@ -52,9 +52,8 @@ class REGION_OCC:public ELIST_LINK float max_x; //Highest x in region inT16 region_type; //Type of crossing - REGION_OCC() { - } //constructor used - //only in COPIER etc + REGION_OCC() {} // constructor used + // only in COPIER etc REGION_OCC( //constructor float min, float max,