mirror/tesseract
1
0
Fork 0
mirror of https://github.com/tesseract-ocr/tesseract.git synced 2025-06-11 12:43:17 +08:00
Code Issues Actions 10 Packages Projects Releases Wiki Activity

OpenCL: Remove PERF_COUNT framework

Browse Source 
It was rarely used, but added a lot of code and an unconditional
dependency on openclwrapper.h.

Signed-off-by: Stefan Weil <sw@weilnetz.de>
This commit is contained in:
Stefan Weil 2019-02-09 10:58:15 +01:00
parent ec8f02c0de
commit d42413dd17
7 changed files with 33 additions and 213 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

15
src/api/baseapi.cpp
View File

@ -70,7 +70,9 @@
#include "mutableiterator.h" // for MutableIterator #include "mutableiterator.h" // for MutableIterator
#include "normalis.h" // for kBlnBaselineOffset, kBlnXHeight #include "normalis.h" // for kBlnBaselineOffset, kBlnXHeight
#include "ocrclass.h" // for ETEXT_DESC #include "ocrclass.h" // for ETEXT_DESC
#include "openclwrapper.h" // for PERF_COUNT_END, PERF_COUNT_START, PERF... #if defined(USE_OPENCL)
#include "openclwrapper.h" // for OpenclDevice
#endif
#include "osdetect.h" // for OSResults, OSBestResult, OrientationId... #include "osdetect.h" // for OSResults, OSBestResult, OrientationId...
#include "pageres.h" // for PAGE_RES_IT, WERD_RES, PAGE_RES, CR_DE... #include "pageres.h" // for PAGE_RES_IT, WERD_RES, PAGE_RES, CR_DE...
#include "paragraphs.h" // for DetectParagraphs #include "paragraphs.h" // for DetectParagraphs
@ -360,7 +362,6 @@ int TessBaseAPI::Init(const char* data, int data_size, const char* language,
const GenericVector<STRING>* vars_vec, const GenericVector<STRING>* vars_vec,
const GenericVector<STRING>* vars_values, const GenericVector<STRING>* vars_values,
bool set_only_non_debug_params, FileReader reader) { bool set_only_non_debug_params, FileReader reader) {
PERF_COUNT_START("TessBaseAPI::Init")
// Default language is "eng". // Default language is "eng".
if (language == nullptr) language = "eng"; if (language == nullptr) language = "eng";
STRING datapath = data_size == 0 ? data : language; STRING datapath = data_size == 0 ? data : language;
@ -376,12 +377,10 @@ int TessBaseAPI::Init(const char* data, int data_size, const char* language,
delete tesseract_; delete tesseract_;
tesseract_ = nullptr; tesseract_ = nullptr;
} }
// PERF_COUNT_SUB("delete tesseract_")
#ifdef USE_OPENCL #ifdef USE_OPENCL
OpenclDevice od; OpenclDevice od;
od.InitEnv(); od.InitEnv();
#endif #endif
PERF_COUNT_SUB("OD::InitEnv()")
bool reset_classifier = true; bool reset_classifier = true;
if (tesseract_ == nullptr) { if (tesseract_ == nullptr) {
reset_classifier = false; reset_classifier = false;
@ -400,7 +399,6 @@ int TessBaseAPI::Init(const char* data, int data_size, const char* language,
} }
} }
PERF_COUNT_SUB("update tesseract_")
// Update datapath and language requested for the last valid initialization. // Update datapath and language requested for the last valid initialization.
if (datapath_ == nullptr) if (datapath_ == nullptr)
datapath_ = new STRING(datapath); datapath_ = new STRING(datapath);
@ -417,14 +415,11 @@ int TessBaseAPI::Init(const char* data, int data_size, const char* language,
last_oem_requested_ = oem; last_oem_requested_ = oem;
#ifndef DISABLED_LEGACY_ENGINE #ifndef DISABLED_LEGACY_ENGINE
// PERF_COUNT_SUB("update last_oem_requested_")
// For same language and datapath, just reset the adaptive classifier. // For same language and datapath, just reset the adaptive classifier.
if (reset_classifier) { if (reset_classifier) {
tesseract_->ResetAdaptiveClassifier(); tesseract_->ResetAdaptiveClassifier();
PERF_COUNT_SUB("tesseract_->ResetAdaptiveClassifier()")
} }
#endif // ndef DISABLED_LEGACY_ENGINE #endif // ndef DISABLED_LEGACY_ENGINE
PERF_COUNT_END
return 0; return 0;
} }
@ -1106,7 +1101,6 @@ bool TessBaseAPI::ProcessPagesInternal(const char* filename,
const char* retry_config, const char* retry_config,
int timeout_millisec, int timeout_millisec,
TessResultRenderer* renderer) { TessResultRenderer* renderer) {
PERF_COUNT_START("ProcessPages")
bool stdInput = !strcmp(filename, "stdin") || !strcmp(filename, "-"); bool stdInput = !strcmp(filename, "stdin") || !strcmp(filename, "-");
if (stdInput) { if (stdInput) {
#ifdef WIN32 #ifdef WIN32
@ -1199,14 +1193,12 @@ bool TessBaseAPI::ProcessPagesInternal(const char* filename,
if (!r || (renderer && !renderer->EndDocument())) { if (!r || (renderer && !renderer->EndDocument())) {
return false; return false;
} }
PERF_COUNT_END
return true; return true;
} }
bool TessBaseAPI::ProcessPage(Pix* pix, int page_index, const char* filename, bool TessBaseAPI::ProcessPage(Pix* pix, int page_index, const char* filename,
const char* retry_config, int timeout_millisec, const char* retry_config, int timeout_millisec,
TessResultRenderer* renderer) { TessResultRenderer* renderer) {
PERF_COUNT_START("ProcessPage")
SetInputName(filename); SetInputName(filename);
SetImage(pix); SetImage(pix);
bool failed = false; bool failed = false;
@ -1264,7 +1256,6 @@ bool TessBaseAPI::ProcessPage(Pix* pix, int page_index, const char* filename,
failed = !renderer->AddImage(this); failed = !renderer->AddImage(this);
} }
PERF_COUNT_END
return !failed; return !failed;
} }

8
src/api/tesseractmain.cpp
View File

@ -28,7 +28,9 @@
#include "baseapi.h" #include "baseapi.h"
#include "basedir.h" #include "basedir.h"
#include "dict.h" #include "dict.h"
#include "openclwrapper.h" // for PERF_COUNT_START, ... #if defined(USE_OPENCL)
#include "openclwrapper.h" // for OpenclDevice
#endif
#include "osdetect.h" #include "osdetect.h"
#include "renderer.h" #include "renderer.h"
#include "simddetect.h" #include "simddetect.h"
@ -582,8 +584,6 @@ int main(int argc, char** argv) {
if (image == nullptr && !list_langs && !print_parameters) if (image == nullptr && !list_langs && !print_parameters)
return EXIT_SUCCESS; return EXIT_SUCCESS;
PERF_COUNT_START("Tesseract:main")
// Call GlobalDawgCache here to create the global DawgCache object before // Call GlobalDawgCache here to create the global DawgCache object before
// the TessBaseAPI object. This fixes the order of destructor calls: // the TessBaseAPI object. This fixes the order of destructor calls:
// first TessBaseAPI must be destructed, DawgCache must be the last object. // first TessBaseAPI must be destructed, DawgCache must be the last object.
@ -721,7 +721,5 @@ int main(int argc, char** argv) {
} }
} }
PERF_COUNT_END
return EXIT_SUCCESS; return EXIT_SUCCESS;
} }

11
src/ccmain/thresholder.cpp
View File

@ -25,8 +25,11 @@
#include <cstring> #include <cstring>
#include "otsuthr.h" #include "otsuthr.h"
#include "tprintf.h" // for tprintf
#include "openclwrapper.h" // for PERF_COUNT_START, ... #if defined(USE_OPENCL)
#include "openclwrapper.h" // for OpenclDevice
#endif
namespace tesseract { namespace tesseract {
@ -263,7 +266,6 @@ Pix* ImageThresholder::GetPixRectGrey() {
// Otsu thresholds the rectangle, taking the rectangle from *this. // Otsu thresholds the rectangle, taking the rectangle from *this.
void ImageThresholder::OtsuThresholdRectToPix(Pix* src_pix, void ImageThresholder::OtsuThresholdRectToPix(Pix* src_pix,
Pix** out_pix) const { Pix** out_pix) const {
PERF_COUNT_START("OtsuThresholdRectToPix")
int* thresholds; int* thresholds;
int* hi_values; int* hi_values;
@ -286,8 +288,6 @@ void ImageThresholder::OtsuThresholdRectToPix(Pix* src_pix,
#endif #endif
delete [] thresholds; delete [] thresholds;
delete [] hi_values; delete [] hi_values;
PERF_COUNT_END
} }
/// Threshold the rectangle, taking everything except the src_pix /// Threshold the rectangle, taking everything except the src_pix
@ -299,7 +299,6 @@ void ImageThresholder::ThresholdRectToPix(Pix* src_pix,
const int* thresholds, const int* thresholds,
const int* hi_values, const int* hi_values,
Pix** pix) const { Pix** pix) const {
PERF_COUNT_START("ThresholdRectToPix")
*pix = pixCreate(rect_width_, rect_height_, 1); *pix = pixCreate(rect_width_, rect_height_, 1);
uint32_t* pixdata = pixGetData(*pix); uint32_t* pixdata = pixGetData(*pix);
int wpl = pixGetWpl(*pix); int wpl = pixGetWpl(*pix);
@ -325,8 +324,6 @@ void ImageThresholder::ThresholdRectToPix(Pix* src_pix,
SET_DATA_BIT(pixline, x); SET_DATA_BIT(pixline, x);
} }
} }
PERF_COUNT_END
} }
} // namespace tesseract. } // namespace tesseract.

8
src/ccstruct/otsuthr.cpp
View File

@ -22,7 +22,9 @@
#include <cstring> #include <cstring>
#include "allheaders.h" #include "allheaders.h"
#include "helpers.h" #include "helpers.h"
#include "openclwrapper.h" // for PERF_COUNT_START, ... #if defined(USE_OPENCL)
#include "openclwrapper.h" // for OpenclDevice
#endif
namespace tesseract { namespace tesseract {
@ -40,7 +42,6 @@ int OtsuThreshold(Pix* src_pix, int left, int top, int width, int height,
int num_channels = pixGetDepth(src_pix) / 8; int num_channels = pixGetDepth(src_pix) / 8;
// Of all channels with no good hi_value, keep the best so we can always // Of all channels with no good hi_value, keep the best so we can always
// produce at least one answer. // produce at least one answer.
PERF_COUNT_START("OtsuThreshold")
int best_hi_value = 1; int best_hi_value = 1;
int best_hi_index = 0; int best_hi_index = 0;
bool any_good_hivalue = false; bool any_good_hivalue = false;
@ -139,7 +140,6 @@ int OtsuThreshold(Pix* src_pix, int left, int top, int width, int height,
// Use the best of the ones that were not good enough. // Use the best of the ones that were not good enough.
(*hi_values)[best_hi_index] = best_hi_value; (*hi_values)[best_hi_index] = best_hi_value;
} }
PERF_COUNT_END
return num_channels; return num_channels;
} }
@ -150,7 +150,6 @@ int OtsuThreshold(Pix* src_pix, int left, int top, int width, int height,
void HistogramRect(Pix* src_pix, int channel, void HistogramRect(Pix* src_pix, int channel,
int left, int top, int width, int height, int left, int top, int width, int height,
int* histogram) { int* histogram) {
PERF_COUNT_START("HistogramRect")
int num_channels = pixGetDepth(src_pix) / 8; int num_channels = pixGetDepth(src_pix) / 8;
channel = ClipToRange(channel, 0, num_channels - 1); channel = ClipToRange(channel, 0, num_channels - 1);
int bottom = top + height; int bottom = top + height;
@ -164,7 +163,6 @@ void HistogramRect(Pix* src_pix, int channel,
++histogram[pixel]; ++histogram[pixel];
} }
} }
PERF_COUNT_END
} }
// Computes the Otsu threshold(s) for the given histogram. // Computes the Otsu threshold(s) for the given histogram.

66
src/opencl/openclwrapper.cpp
View File

@ -14,6 +14,7 @@
#include <unistd.h> #include <unistd.h>
#endif #endif
#include <cfloat> #include <cfloat>
#include <ctime> // for clock_gettime
#include "oclkernels.h" #include "oclkernels.h"
#include "openclwrapper.h" #include "openclwrapper.h"
@ -22,6 +23,22 @@
#include "otsuthr.h" #include "otsuthr.h"
#include "thresholder.h" #include "thresholder.h"
// platform preprocessor commands
#if defined(WIN32) || defined(__WIN32__) || defined(_WIN32) || \
defined(__CYGWIN__) || defined(__MINGW32__)
#define ON_WINDOWS 1
#define ON_APPLE 0
#elif defined(__linux__)
#define ON_WINDOWS 0
#define ON_APPLE 0
#elif defined(__APPLE__)
#define ON_WINDOWS 0
#define ON_APPLE 1
#else
#define ON_WINDOWS 0
#define ON_APPLE 0
#endif
#if ON_APPLE #if ON_APPLE
#include <mach/mach_time.h> #include <mach/mach_time.h>
#endif #endif
@ -712,19 +729,15 @@ int OpenclDevice::initMorphCLAllocations(l_int32 wpl, l_int32 h, Pix* pixs) {
} }
int OpenclDevice::InitEnv() { int OpenclDevice::InitEnv() {
// PERF_COUNT_START("OD::InitEnv")
// tprintf("[OD] OpenclDevice::InitEnv()\n"); // tprintf("[OD] OpenclDevice::InitEnv()\n");
#ifdef SAL_WIN32 #ifdef SAL_WIN32
while (1) { while (1) {
if (1 == LoadOpencl()) break; if (1 == LoadOpencl()) break;
} }
PERF_COUNT_SUB("LoadOpencl")
#endif #endif
// sets up environment, compiles programs // sets up environment, compiles programs
InitOpenclRunEnv_DeviceSelection(0); InitOpenclRunEnv_DeviceSelection(0);
// PERF_COUNT_SUB("called InitOpenclRunEnv_DS")
// PERF_COUNT_END
return 1; return 1;
} }
@ -755,11 +768,9 @@ int OpenclDevice::RegistOpenclKernel() {
} }
int OpenclDevice::InitOpenclRunEnv_DeviceSelection(int argc) { int OpenclDevice::InitOpenclRunEnv_DeviceSelection(int argc) {
// PERF_COUNT_START("InitOpenclRunEnv_DS")
if (!isInited) { if (!isInited) {
// after programs compiled, selects best device // after programs compiled, selects best device
ds_device bestDevice_DS = getDeviceSelection(); ds_device bestDevice_DS = getDeviceSelection();
// PERF_COUNT_SUB("called getDeviceSelection()")
cl_device_id bestDevice = bestDevice_DS.oclDeviceID; cl_device_id bestDevice = bestDevice_DS.oclDeviceID;
// overwrite global static GPUEnv with new device // overwrite global static GPUEnv with new device
if (selectedDeviceIsOpenCL()) { if (selectedDeviceIsOpenCL()) {
@ -768,16 +779,13 @@ int OpenclDevice::InitOpenclRunEnv_DeviceSelection(int argc) {
populateGPUEnvFromDevice(&gpuEnv, bestDevice); populateGPUEnvFromDevice(&gpuEnv, bestDevice);
gpuEnv.mnFileCount = 0; // argc; gpuEnv.mnFileCount = 0; // argc;
gpuEnv.mnKernelCount = 0UL; gpuEnv.mnKernelCount = 0UL;
// PERF_COUNT_SUB("populate gpuEnv")
CompileKernelFile(&gpuEnv, ""); CompileKernelFile(&gpuEnv, "");
// PERF_COUNT_SUB("CompileKernelFile")
} else { } else {
// tprintf("[DS] InitOpenclRunEnv_DS::Skipping populateGPUEnvFromDevice() // tprintf("[DS] InitOpenclRunEnv_DS::Skipping populateGPUEnvFromDevice()
// b/c native cpu selected\n"); // b/c native cpu selected\n");
} }
isInited = 1; isInited = 1;
} }
// PERF_COUNT_END
return 0; return 0;
} }
@ -941,7 +949,6 @@ int OpenclDevice::GeneratBinFromKernelSource(cl_program program,
} }
int OpenclDevice::CompileKernelFile(GPUEnv* gpuInfo, const char* buildOption) { int OpenclDevice::CompileKernelFile(GPUEnv* gpuInfo, const char* buildOption) {
// PERF_COUNT_START("CompileKernelFile")
cl_int clStatus = 0; cl_int clStatus = 0;
const char* source; const char* source;
size_t source_size[1]; size_t source_size[1];
@ -962,14 +969,12 @@ int OpenclDevice::CompileKernelFile(GPUEnv* gpuInfo, const char* buildOption) {
binaryExisted = 0; binaryExisted = 0;
binaryExisted = BinaryGenerated( binaryExisted = BinaryGenerated(
filename, &fd); // don't check for binary during microbenchmark filename, &fd); // don't check for binary during microbenchmark
// PERF_COUNT_SUB("BinaryGenerated")
if (binaryExisted == 1) { if (binaryExisted == 1) {
clStatus = clGetContextInfo(gpuInfo->mpContext, CL_CONTEXT_NUM_DEVICES, clStatus = clGetContextInfo(gpuInfo->mpContext, CL_CONTEXT_NUM_DEVICES,
sizeof(numDevices), &numDevices, nullptr); sizeof(numDevices), &numDevices, nullptr);
CHECK_OPENCL(clStatus, "clGetContextInfo"); CHECK_OPENCL(clStatus, "clGetContextInfo");
std::vector<cl_device_id> mpArryDevsID(numDevices); std::vector<cl_device_id> mpArryDevsID(numDevices);
// PERF_COUNT_SUB("get numDevices")
bool b_error = fseek(fd, 0, SEEK_END) < 0; bool b_error = fseek(fd, 0, SEEK_END) < 0;
long pos = ftell(fd); long pos = ftell(fd);
b_error |= (pos <= 0); b_error |= (pos <= 0);
@ -986,29 +991,24 @@ int OpenclDevice::CompileKernelFile(GPUEnv* gpuInfo, const char* buildOption) {
b_error |= fread(&binary[0], 1, length, fd) != length; b_error |= fread(&binary[0], 1, length, fd) != length;
fclose(fd); fclose(fd);
// PERF_COUNT_SUB("read file")
fd = nullptr; fd = nullptr;
// grab the handles to all of the devices in the context. // grab the handles to all of the devices in the context.
clStatus = clGetContextInfo(gpuInfo->mpContext, CL_CONTEXT_DEVICES, clStatus = clGetContextInfo(gpuInfo->mpContext, CL_CONTEXT_DEVICES,
sizeof(cl_device_id) * numDevices, sizeof(cl_device_id) * numDevices,
&mpArryDevsID[0], nullptr); &mpArryDevsID[0], nullptr);
CHECK_OPENCL(clStatus, "clGetContextInfo"); CHECK_OPENCL(clStatus, "clGetContextInfo");
// PERF_COUNT_SUB("get devices")
// fprintf(stderr, "[OD] Create kernel from binary\n"); // fprintf(stderr, "[OD] Create kernel from binary\n");
const uint8_t* c_binary = &binary[0]; const uint8_t* c_binary = &binary[0];
gpuInfo->mpArryPrograms[idx] = clCreateProgramWithBinary( gpuInfo->mpArryPrograms[idx] = clCreateProgramWithBinary(
gpuInfo->mpContext, numDevices, &mpArryDevsID[0], &length, &c_binary, gpuInfo->mpContext, numDevices, &mpArryDevsID[0], &length, &c_binary,
&binary_status, &clStatus); &binary_status, &clStatus);
CHECK_OPENCL(clStatus, "clCreateProgramWithBinary"); CHECK_OPENCL(clStatus, "clCreateProgramWithBinary");
// PERF_COUNT_SUB("clCreateProgramWithBinary")
// PERF_COUNT_SUB("binaryExisted")
} else { } else {
// create a CL program using the kernel source // create a CL program using the kernel source
// fprintf(stderr, "[OD] Create kernel from source\n"); // fprintf(stderr, "[OD] Create kernel from source\n");
gpuInfo->mpArryPrograms[idx] = clCreateProgramWithSource( gpuInfo->mpArryPrograms[idx] = clCreateProgramWithSource(
gpuInfo->mpContext, 1, &source, source_size, &clStatus); gpuInfo->mpContext, 1, &source, source_size, &clStatus);
CHECK_OPENCL(clStatus, "clCreateProgramWithSource"); CHECK_OPENCL(clStatus, "clCreateProgramWithSource");
// PERF_COUNT_SUB("!binaryExisted")
} }
if (gpuInfo->mpArryPrograms[idx] == (cl_program) nullptr) { if (gpuInfo->mpArryPrograms[idx] == (cl_program) nullptr) {
@ -1018,19 +1018,15 @@ int OpenclDevice::CompileKernelFile(GPUEnv* gpuInfo, const char* buildOption) {
// char options[512]; // char options[512];
// create a cl program executable for all the devices specified // create a cl program executable for all the devices specified
// tprintf("[OD] BuildProgram.\n"); // tprintf("[OD] BuildProgram.\n");
PERF_COUNT_START("OD::CompileKernel::clBuildProgram")
if (!gpuInfo->mnIsUserCreated) { if (!gpuInfo->mnIsUserCreated) {
clStatus = clStatus =
clBuildProgram(gpuInfo->mpArryPrograms[idx], 1, gpuInfo->mpArryDevsID, clBuildProgram(gpuInfo->mpArryPrograms[idx], 1, gpuInfo->mpArryDevsID,
buildOption, nullptr, nullptr); buildOption, nullptr, nullptr);
// PERF_COUNT_SUB("clBuildProgram notUserCreated")
} else { } else {
clStatus = clStatus =
clBuildProgram(gpuInfo->mpArryPrograms[idx], 1, &(gpuInfo->mpDevID), clBuildProgram(gpuInfo->mpArryPrograms[idx], 1, &(gpuInfo->mpDevID),
buildOption, nullptr, nullptr); buildOption, nullptr, nullptr);
// PERF_COUNT_SUB("clBuildProgram isUserCreated")
} }
PERF_COUNT_END
if (clStatus != CL_SUCCESS) { if (clStatus != CL_SUCCESS) {
tprintf("BuildProgram error!\n"); tprintf("BuildProgram error!\n");
size_t length; size_t length;
@ -1068,26 +1064,21 @@ int OpenclDevice::CompileKernelFile(GPUEnv* gpuInfo, const char* buildOption) {
fclose(fd1); fclose(fd1);
} }
// PERF_COUNT_SUB("build error log")
return 0; return 0;
} }
strcpy(gpuInfo->mArryKnelSrcFile[idx], filename); strcpy(gpuInfo->mArryKnelSrcFile[idx], filename);
// PERF_COUNT_SUB("strcpy")
if (binaryExisted == 0) { if (binaryExisted == 0) {
GeneratBinFromKernelSource(gpuInfo->mpArryPrograms[idx], filename); GeneratBinFromKernelSource(gpuInfo->mpArryPrograms[idx], filename);
PERF_COUNT_SUB("GenerateBinFromKernelSource")
} }
gpuInfo->mnFileCount += 1; gpuInfo->mnFileCount += 1;
// PERF_COUNT_END
return 1; return 1;
} }
l_uint32* OpenclDevice::pixReadFromTiffKernel(l_uint32* tiffdata, l_int32 w, l_uint32* OpenclDevice::pixReadFromTiffKernel(l_uint32* tiffdata, l_int32 w,
l_int32 h, l_int32 wpl, l_int32 h, l_int32 wpl,
l_uint32* line) { l_uint32* line) {
PERF_COUNT_START("pixReadFromTiffKernel")
cl_int clStatus; cl_int clStatus;
KernelEnv rEnv; KernelEnv rEnv;
size_t globalThreads[2]; size_t globalThreads[2];
@ -1130,7 +1121,6 @@ l_uint32* OpenclDevice::pixReadFromTiffKernel(l_uint32* tiffdata, l_int32 w,
CHECK_OPENCL(clStatus, "clSetKernelArg"); CHECK_OPENCL(clStatus, "clSetKernelArg");
// Kernel enqueue // Kernel enqueue
PERF_COUNT_SUB("before")
clStatus = clStatus =
clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, rEnv.mpkKernel, 2, nullptr, clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, rEnv.mpkKernel, 2, nullptr,
globalThreads, localThreads, 0, nullptr, nullptr); globalThreads, localThreads, 0, nullptr, nullptr);
@ -1145,8 +1135,6 @@ l_uint32* OpenclDevice::pixReadFromTiffKernel(l_uint32* tiffdata, l_int32 w,
// Sync // Sync
clFinish(rEnv.mpkCmdQueue); clFinish(rEnv.mpkCmdQueue);
PERF_COUNT_SUB("kernel & map")
PERF_COUNT_END
return pResult; return pResult;
} }
@ -1627,7 +1615,6 @@ int OpenclDevice::HistogramRectOCL(void* imageData,
int top, // always 0 int top, // always 0
int width, int height, int kHistogramSize, int width, int height, int kHistogramSize,
int* histogramAllChannels) { int* histogramAllChannels) {
PERF_COUNT_START("HistogramRectOCL")
cl_int clStatus; cl_int clStatus;
int retVal = 0; int retVal = 0;
KernelEnv histKern; KernelEnv histKern;
@ -1742,7 +1729,6 @@ int OpenclDevice::HistogramRectOCL(void* imageData,
CHECK_OPENCL(clStatus, "clSetKernelArg histogramBuffer"); CHECK_OPENCL(clStatus, "clSetKernelArg histogramBuffer");
/* launch histogram */ /* launch histogram */
PERF_COUNT_SUB("before")
clStatus = clEnqueueNDRangeKernel(histKern.mpkCmdQueue, histKern.mpkKernel, 1, clStatus = clEnqueueNDRangeKernel(histKern.mpkCmdQueue, histKern.mpkKernel, 1,
nullptr, global_work_size, local_work_size, nullptr, global_work_size, local_work_size,
0, nullptr, nullptr); 0, nullptr, nullptr);
@ -1763,7 +1749,6 @@ int OpenclDevice::HistogramRectOCL(void* imageData,
if (clStatus != 0) { if (clStatus != 0) {
retVal = -1; retVal = -1;
} }
PERF_COUNT_SUB("redKernel")
/* map results back from gpu */ /* map results back from gpu */
ptr = clEnqueueMapBuffer(histRedKern.mpkCmdQueue, histogramBuffer, CL_TRUE, ptr = clEnqueueMapBuffer(histRedKern.mpkCmdQueue, histogramBuffer, CL_TRUE,
@ -1779,8 +1764,6 @@ int OpenclDevice::HistogramRectOCL(void* imageData,
clReleaseMemObject(histogramBuffer); clReleaseMemObject(histogramBuffer);
clReleaseMemObject(imageBuffer); clReleaseMemObject(imageBuffer);
PERF_COUNT_SUB("after")
PERF_COUNT_END
return retVal; return retVal;
} }
@ -1794,7 +1777,6 @@ int OpenclDevice::ThresholdRectToPixOCL(unsigned char* imageData,
int* thresholds, int* hi_values, int* thresholds, int* hi_values,
Pix** pix, int height, int width, Pix** pix, int height, int width,
int top, int left) { int top, int left) {
PERF_COUNT_START("ThresholdRectToPixOCL")
int retVal = 0; int retVal = 0;
/* create pix result buffer */ /* create pix result buffer */
*pix = pixCreate(width, height, 1); *pix = pixCreate(width, height, 1);
@ -1874,13 +1856,11 @@ int OpenclDevice::ThresholdRectToPixOCL(unsigned char* imageData,
CHECK_OPENCL(clStatus, "clSetKernelArg pixThBuffer"); CHECK_OPENCL(clStatus, "clSetKernelArg pixThBuffer");
/* launch kernel & wait */ /* launch kernel & wait */
PERF_COUNT_SUB("before")
clStatus = clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, rEnv.mpkKernel, 1, clStatus = clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, rEnv.mpkKernel, 1,
nullptr, global_work_size, local_work_size, nullptr, global_work_size, local_work_size,
0, nullptr, nullptr); 0, nullptr, nullptr);
CHECK_OPENCL(clStatus, "clEnqueueNDRangeKernel kernel_ThresholdRectToPix"); CHECK_OPENCL(clStatus, "clEnqueueNDRangeKernel kernel_ThresholdRectToPix");
clFinish(rEnv.mpkCmdQueue); clFinish(rEnv.mpkCmdQueue);
PERF_COUNT_SUB("kernel")
if (clStatus != 0) { if (clStatus != 0) {
tprintf("Setting return value to -1\n"); tprintf("Setting return value to -1\n");
retVal = -1; retVal = -1;
@ -1897,8 +1877,6 @@ int OpenclDevice::ThresholdRectToPixOCL(unsigned char* imageData,
clReleaseMemObject(thresholdsBuffer); clReleaseMemObject(thresholdsBuffer);
clReleaseMemObject(hiValuesBuffer); clReleaseMemObject(hiValuesBuffer);
PERF_COUNT_SUB("after")
PERF_COUNT_END
return retVal; return retVal;
} }
@ -2499,16 +2477,13 @@ static ds_status evaluateScoreForDevice(ds_device* device, void* inputData) {
// initial call to select device // initial call to select device
ds_device OpenclDevice::getDeviceSelection() { ds_device OpenclDevice::getDeviceSelection() {
if (!deviceIsSelected) { if (!deviceIsSelected) {
PERF_COUNT_START("getDeviceSelection")
// check if opencl is available at runtime // check if opencl is available at runtime
if (1 == LoadOpencl()) { if (1 == LoadOpencl()) {
// opencl is available // opencl is available
// PERF_COUNT_SUB("LoadOpencl")
// setup devices // setup devices
ds_status status; ds_status status;
ds_profile* profile; ds_profile* profile;
status = initDSProfile(&profile, "v0.1"); status = initDSProfile(&profile, "v0.1");
PERF_COUNT_SUB("initDSProfile")
// try reading scores from file // try reading scores from file
const char* fileName = "tesseract_opencl_profile_devices.dat"; const char* fileName = "tesseract_opencl_profile_devices.dat";
status = readProfileFromFile(profile, deserializeScore, fileName); status = readProfileFromFile(profile, deserializeScore, fileName);
@ -2520,16 +2495,13 @@ ds_device OpenclDevice::getDeviceSelection() {
// create input data // create input data
TessScoreEvaluationInputData input; TessScoreEvaluationInputData input;
populateTessScoreEvaluationInputData(&input); populateTessScoreEvaluationInputData(&input);
// PERF_COUNT_SUB("populateTessScoreEvaluationInputData")
// perform evaluations // perform evaluations
unsigned int numUpdates; unsigned int numUpdates;
status = profileDevices(profile, DS_EVALUATE_ALL, status = profileDevices(profile, DS_EVALUATE_ALL,
evaluateScoreForDevice, &input, &numUpdates); evaluateScoreForDevice, &input, &numUpdates);
PERF_COUNT_SUB("profileDevices")
// write scores to file // write scores to file
if (status == DS_SUCCESS) { if (status == DS_SUCCESS) {
status = writeProfileToFile(profile, serializeScore, fileName); status = writeProfileToFile(profile, serializeScore, fileName);
PERF_COUNT_SUB("writeProfileToFile")
if (status == DS_SUCCESS) { if (status == DS_SUCCESS) {
tprintf("[DS] Scores written to file (%s).\n", fileName); tprintf("[DS] Scores written to file (%s).\n", fileName);
} else { } else {
@ -2544,7 +2516,6 @@ ds_device OpenclDevice::getDeviceSelection() {
"file.\n"); "file.\n");
} }
} else { } else {
PERF_COUNT_SUB("readProfileFromFile")
tprintf("[DS] Profile read from file (%s).\n", fileName); tprintf("[DS] Profile read from file (%s).\n", fileName);
} }
@ -2614,10 +2585,7 @@ ds_device OpenclDevice::getDeviceSelection() {
selectedDevice.oclDriverVersion = nullptr; selectedDevice.oclDriverVersion = nullptr;
} }
deviceIsSelected = true; deviceIsSelected = true;
PERF_COUNT_SUB("select from Profile")
PERF_COUNT_END
} }
// PERF_COUNT_END
return selectedDevice; return selectedDevice;
} }

130
src/opencl/openclwrapper.h
View File

@ -18,136 +18,6 @@
// including CL/cl.h doesn't occur until USE_OPENCL defined below // including CL/cl.h doesn't occur until USE_OPENCL defined below
// platform preprocessor commands
#if defined(WIN32) || defined(__WIN32__) || defined(_WIN32) || \
defined(__CYGWIN__) || defined(__MINGW32__)
#define ON_WINDOWS 1
#define ON_LINUX 0
#define ON_APPLE 0
#define ON_OTHER 0
#define IF_WINDOWS(X) X
#define IF_LINUX(X)
#define IF_APPLE(X)
#define IF_OTHER(X)
#define NOT_WINDOWS(X)
#elif defined(__linux__)
#define ON_WINDOWS 0
#define ON_LINUX 1
#define ON_APPLE 0
#define ON_OTHER 0
#define IF_WINDOWS(X)
#define IF_LINUX(X) X
#define IF_APPLE(X)
#define IF_OTHER(X)
#define NOT_WINDOWS(X) X
#elif defined(__APPLE__)
#define ON_WINDOWS 0
#define ON_LINUX 0
#define ON_APPLE 1
#define ON_OTHER 0
#define IF_WINDOWS(X)
#define IF_LINUX(X)
#define IF_APPLE(X) X
#define IF_OTHER(X)
#define NOT_WINDOWS(X) X
#else
#define ON_WINDOWS 0
#define ON_LINUX 0
#define ON_APPLE 0
#define ON_OTHER 1
#define IF_WINDOWS(X)
#define IF_LINUX(X)
#define IF_APPLE(X)
#define IF_OTHER(X) X
#define NOT_WINDOWS(X) X
#endif
#if ON_LINUX
#include <ctime>
#endif
/************************************************************************************
* enable/disable reporting of performance
* PERF_REPORT_LEVEL
* 0 - no reporting
* 1 - no reporting
* 2 - report total function call time for functions we're tracking
* 3 - optionally report breakdown of function calls (kernel launch, kernel
*time, data copies)
************************************************************************************/
#define PERF_COUNT_VERBOSE 1
#define PERF_COUNT_REPORT_STR "[%36s], %24s, %11.6f\n"
#if ON_WINDOWS
#if PERF_COUNT_VERBOSE >= 2
#define PERF_COUNT_START(FUNCT_NAME) \
char* funct_name = FUNCT_NAME; \
double elapsed_time_sec; \
LARGE_INTEGER freq, time_funct_start, time_funct_end, time_sub_start, \
time_sub_end; \
QueryPerformanceFrequency(&freq); \
QueryPerformanceCounter(&time_funct_start); \
time_sub_start = time_funct_start; \
time_sub_end = time_funct_start;
#define PERF_COUNT_END \
QueryPerformanceCounter(&time_funct_end); \
elapsed_time_sec = (time_funct_end.QuadPart - time_funct_start.QuadPart) / \
(double)(freq.QuadPart); \
tprintf(PERF_COUNT_REPORT_STR, funct_name, "total", elapsed_time_sec);
#else
#define PERF_COUNT_START(FUNCT_NAME)
#define PERF_COUNT_END
#endif
#if PERF_COUNT_VERBOSE >= 3
#define PERF_COUNT_SUB(SUB) \
QueryPerformanceCounter(&time_sub_end); \
elapsed_time_sec = (time_sub_end.QuadPart - time_sub_start.QuadPart) / \
(double)(freq.QuadPart); \
tprintf(PERF_COUNT_REPORT_STR, funct_name, SUB, elapsed_time_sec); \
time_sub_start = time_sub_end;
#else
#define PERF_COUNT_SUB(SUB)
#endif
// not on windows
#else
#if PERF_COUNT_VERBOSE >= 2
#define PERF_COUNT_START(FUNCT_NAME) \
char* funct_name = FUNCT_NAME; \
double elapsed_time_sec; \
timespec time_funct_start, time_funct_end, time_sub_start, time_sub_end; \
clock_gettime(CLOCK_MONOTONIC, &time_funct_start); \
time_sub_start = time_funct_start; \
time_sub_end = time_funct_start;
#define PERF_COUNT_END \
clock_gettime(CLOCK_MONOTONIC, &time_funct_end); \
elapsed_time_sec = \
(time_funct_end.tv_sec - time_funct_start.tv_sec) * 1.0 + \
(time_funct_end.tv_nsec - time_funct_start.tv_nsec) / 1000000000.0; \
tprintf(PERF_COUNT_REPORT_STR, funct_name, "total", elapsed_time_sec);
#else
#define PERF_COUNT_START(FUNCT_NAME)
#define PERF_COUNT_END
#endif
#if PERF_COUNT_VERBOSE >= 3
#define PERF_COUNT_SUB(SUB) \
clock_gettime(CLOCK_MONOTONIC, &time_sub_end); \
elapsed_time_sec = \
(time_sub_end.tv_sec - time_sub_start.tv_sec) * 1.0 + \
(time_sub_end.tv_nsec - time_sub_start.tv_nsec) / 1000000000.0; \
tprintf(PERF_COUNT_REPORT_STR, funct_name, SUB, elapsed_time_sec); \
time_sub_start = time_sub_end;
#else
#define PERF_COUNT_SUB(SUB)
#endif
#endif
/************************************************************************** /**************************************************************************
* enable/disable use of OpenCL * enable/disable use of OpenCL
**************************************************************************/ **************************************************************************/

8
src/textord/linefind.cpp
View File

@ -27,7 +27,9 @@
#include "tabvector.h" #include "tabvector.h"
#include "blobbox.h" #include "blobbox.h"
#include "edgblob.h" #include "edgblob.h"
#include "openclwrapper.h" // for PERF_COUNT_START, ... #if defined(USE_OPENCL)
#include "openclwrapper.h" // for OpenclDevice
#endif
#include "allheaders.h" #include "allheaders.h"
@ -243,7 +245,6 @@ void LineFinder::FindAndRemoveLines(int resolution, bool debug, Pix* pix,
Pix** pix_music_mask, Pix** pix_music_mask,
TabVector_LIST* v_lines, TabVector_LIST* v_lines,
TabVector_LIST* h_lines) { TabVector_LIST* h_lines) {
PERF_COUNT_START("FindAndRemoveLines")
if (pix == nullptr || vertical_x == nullptr || vertical_y == nullptr) { if (pix == nullptr || vertical_x == nullptr || vertical_y == nullptr) {
tprintf("Error in parameters for LineFinder::FindAndRemoveLines\n"); tprintf("Error in parameters for LineFinder::FindAndRemoveLines\n");
return; return;
@ -308,7 +309,6 @@ void LineFinder::FindAndRemoveLines(int resolution, bool debug, Pix* pix,
"vhlinefinding.pdf"); "vhlinefinding.pdf");
pixaDestroy(&pixa_display); pixaDestroy(&pixa_display);
} }
PERF_COUNT_END
} }
// Converts the Boxa array to a list of C_BLOB, getting rid of severely // Converts the Boxa array to a list of C_BLOB, getting rid of severely
@ -583,7 +583,6 @@ void LineFinder::GetLineMasks(int resolution, Pix* src_pix,
} }
int closing_brick = max_line_width / 3; int closing_brick = max_line_width / 3;
PERF_COUNT_START("GetLineMasksMorph")
// only use opencl if compiled w/ OpenCL and selected device is opencl // only use opencl if compiled w/ OpenCL and selected device is opencl
#ifdef USE_OPENCL #ifdef USE_OPENCL
if (OpenclDevice::selectedDeviceIsOpenCL()) { if (OpenclDevice::selectedDeviceIsOpenCL()) {
@ -625,7 +624,6 @@ void LineFinder::GetLineMasks(int resolution, Pix* src_pix,
#ifdef USE_OPENCL #ifdef USE_OPENCL
} }
#endif #endif
PERF_COUNT_END
// Lines are sufficiently rare, that it is worth checking for a zero image. // Lines are sufficiently rare, that it is worth checking for a zero image.
l_int32 v_empty = 0; l_int32 v_empty = 0;