diff --git a/.gitignore b/.gitignore index 3d403b38..81b9995a 100644 --- a/.gitignore +++ b/.gitignore @@ -83,10 +83,6 @@ __pycache__ *.traineddata tessdata_* -# OpenCL -tesseract_opencl_profile_devices.dat -kernel*.bin - # build dirs /build* /*.dll diff --git a/CMakeLists.txt b/CMakeLists.txt index 68da6c53..367f1098 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -88,7 +88,6 @@ option(GRAPHICS_DISABLED "Disable disable graphics (ScrollView)" OFF) option(DISABLED_LEGACY_ENGINE "Disable the legacy OCR engine" OFF) option(ENABLE_LTO "Enable link-time optimization" OFF) option(FAST_FLOAT "Enable float for LSTM" ON) -option(ENABLE_OPENCL "Enable unsupported experimental OpenCL support" OFF) option(ENABLE_NATIVE "Enable optimization for host CPU (could break HW compatibility)" OFF) # see @@ -460,18 +459,6 @@ else() endif(DISABLE_CURL) endif() -if(ENABLE_OPENCL) - find_package(OpenCL) - if(OpenCL_FOUND) - include_directories(${OpenCL_INCLUDE_DIRS}) - message(STATUS "OpenCL_INCLUDE_DIRS: ${OpenCL_INCLUDE_DIRS}") - message(STATUS "OpenCL_LIBRARY: ${OpenCL_LIBRARY}") - set(USE_OPENCL ON) - else() - set(USE_OPENCL OFF) - endif(OpenCL_FOUND) -endif(ENABLE_OPENCL) - # ############################################################################## # # configure @@ -565,11 +552,6 @@ message(STATUS "Disable the legacy OCR engine [DISABLED_LEGACY_ENGINE]: " message(STATUS "Build training tools [BUILD_TRAINING_TOOLS]: " "${BUILD_TRAINING_TOOLS}") message(STATUS "Build tests [BUILD_TESTS]: ${BUILD_TESTS}") -if(ENABLE_OPENCL) - message( - STATUS - "Enable unsupported experimental OpenCL [ENABLE_OPENCL]: ${USE_OPENCL}") -endif(ENABLE_OPENCL) message(STATUS "Use system ICU Library [USE_SYSTEM_ICU]: ${USE_SYSTEM_ICU}") message( STATUS "Install tesseract configs [INSTALL_CONFIGS]: ${INSTALL_CONFIGS}") @@ -608,7 +590,6 @@ file( src/cutil/*.cpp src/dict/*.cpp src/lstm/*.cpp - src/opencl/*.cpp src/textord/*.cpp src/viewer/*.cpp src/wordrec/*.cpp) @@ -749,7 +730,6 @@ file( src/cutil/*.h src/dict/*.h src/lstm/*.h - src/opencl/*.h src/textord/*.h src/viewer/*.h src/wordrec/*.h) @@ -816,7 +796,6 @@ target_include_directories( $ $ $ - $ $ $ $ @@ -832,9 +811,6 @@ target_link_libraries(libtesseract PRIVATE ${LIB_Ws2_32} ${LIB_pthread}) if(OpenMP_CXX_FOUND) target_link_libraries(libtesseract PUBLIC OpenMP::OpenMP_CXX) endif() -if(OpenCL_FOUND) - target_link_libraries(libtesseract PUBLIC OpenCL::OpenCL) -endif() if(LibArchive_FOUND) target_link_libraries(libtesseract PUBLIC ${LibArchive_LIBRARIES}) endif(LibArchive_FOUND) diff --git a/Makefile.am b/Makefile.am index b0e0baee..b51857f9 100644 --- a/Makefile.am +++ b/Makefile.am @@ -78,7 +78,6 @@ if VISIBILITY AM_CPPFLAGS += -DTESS_EXPORTS AM_CPPFLAGS += -fvisibility=hidden -fvisibility-inlines-hidden -fPIC endif -AM_CPPFLAGS += $(OPENCL_CPPFLAGS) AM_CXXFLAGS = $(OPENMP_CXXFLAGS) @@ -94,9 +93,6 @@ libtesseract_la_CPPFLAGS += -I$(top_srcdir)/src/classify libtesseract_la_CPPFLAGS += -I$(top_srcdir)/src/cutil libtesseract_la_CPPFLAGS += -I$(top_srcdir)/src/dict libtesseract_la_CPPFLAGS += -I$(top_srcdir)/src/lstm -if OPENCL -libtesseract_la_CPPFLAGS += -I$(top_srcdir)/src/opencl -endif libtesseract_la_CPPFLAGS += -I$(top_srcdir)/src/textord libtesseract_la_CPPFLAGS += -I$(top_srcdir)/src/training/common libtesseract_la_CPPFLAGS += -I$(top_srcdir)/src/viewer @@ -104,7 +100,7 @@ libtesseract_la_CPPFLAGS += -I$(top_srcdir)/src/wordrec libtesseract_la_CPPFLAGS += $(libcurl_CFLAGS) lib_LTLIBRARIES = libtesseract.la -libtesseract_la_LDFLAGS = $(LEPTONICA_LIBS) $(OPENCL_LDFLAGS) +libtesseract_la_LDFLAGS = $(LEPTONICA_LIBS) libtesseract_la_LDFLAGS += $(libarchive_LIBS) libtesseract_la_LDFLAGS += $(libcurl_LIBS) libtesseract_la_LDFLAGS += $(TENSORFLOW_LIBS) @@ -127,9 +123,6 @@ libtesseract_la_SOURCES += src/api/wordstrboxrenderer.cpp libtesseract_la_LIBADD = libtesseract_ccutil.la libtesseract_la_LIBADD += libtesseract_lstm.la libtesseract_la_LIBADD += libtesseract_native.la -if OPENCL -libtesseract_la_LIBADD += libtesseract_opencl.la -endif # Rules for src/arch. @@ -561,23 +554,6 @@ if TENSORFLOW libtesseract_lstm_la_SOURCES += src/lstm/tfnetwork.pb.cc endif -# Rules for src/opencl. - -if OPENCL -libtesseract_opencl_la_CPPFLAGS = $(AM_CPPFLAGS) -libtesseract_opencl_la_CPPFLAGS += $(OPENCL_CFLAGS) -libtesseract_opencl_la_CPPFLAGS += -I$(top_srcdir)/src/ccutil -libtesseract_opencl_la_CPPFLAGS += -I$(top_srcdir)/src/ccstruct -libtesseract_opencl_la_CPPFLAGS += -I$(top_srcdir)/src/ccmain - -noinst_HEADERS += src/opencl/openclwrapper.h -noinst_HEADERS += src/opencl/oclkernels.h - -noinst_LTLIBRARIES += libtesseract_opencl.la - -libtesseract_opencl_la_SOURCES = src/opencl/openclwrapper.cpp -endif - # Rules for src/textord. noinst_HEADERS += src/textord/alignedblob.h @@ -743,16 +719,12 @@ tesseract_CPPFLAGS += -I$(top_srcdir)/src/dict tesseract_CPPFLAGS += -I$(top_srcdir)/src/textord tesseract_CPPFLAGS += -I$(top_srcdir)/src/viewer tesseract_CPPFLAGS += -I$(top_srcdir)/src/wordrec -if OPENCL -tesseract_CPPFLAGS += -I$(top_srcdir)/src/opencl -endif tesseract_CPPFLAGS += $(AM_CPPFLAGS) if VISIBILITY tesseract_CPPFLAGS += -DTESS_IMPORTS endif -tesseract_LDFLAGS = $(OPENCL_LDFLAGS) -tesseract_LDFLAGS += $(OPENMP_CXXFLAGS) +tesseract_LDFLAGS = $(OPENMP_CXXFLAGS) tesseract_LDADD = libtesseract.la tesseract_LDADD += $(LEPTONICA_LIBS) @@ -927,7 +899,6 @@ EXTRA_PROGRAMS += $(trainingtools) extralib = libtesseract.la extralib += $(libarchive_LIBS) extralib += $(LEPTONICA_LIBS) -extralib += $(OPENCL_LDFLAGS) extralib += $(TENSORFLOW_LIBS) if T_WIN extralib += -lws2_32 @@ -1308,7 +1279,7 @@ TESTS = $(check_PROGRAMS) apiexample_test_SOURCES = unittest/apiexample_test.cc apiexample_test_CPPFLAGS = $(unittest_CPPFLAGS) -apiexample_test_LDFLAGS = $(OPENCL_LDFLAGS) $(LEPTONICA_LIBS) +apiexample_test_LDFLAGS = $(LEPTONICA_LIBS) apiexample_test_LDADD = $(TESS_LIBS) $(LEPTONICA_LIBS) if !DISABLED_LEGACY_ENGINE @@ -1494,7 +1465,7 @@ endif # !DISABLED_LEGACY_ENGINE progress_test_SOURCES = unittest/progress_test.cc progress_test_CPPFLAGS = $(unittest_CPPFLAGS) -progress_test_LDFLAGS = $(OPENCL_LDFLAGS) $(LEPTONICA_LIBS) +progress_test_LDFLAGS = $(LEPTONICA_LIBS) progress_test_LDADD = $(GTEST_LIBS) $(GMOCK_LIBS) $(TESS_LIBS) $(LEPTONICA_LIBS) qrsequence_test_SOURCES = unittest/qrsequence_test.cc diff --git a/cmake/Configure.cmake b/cmake/Configure.cmake index 1e0ee2a3..d211d117 100644 --- a/cmake/Configure.cmake +++ b/cmake/Configure.cmake @@ -91,8 +91,6 @@ set(include_files_list unistd.h cairo/cairo-version.h - CL/cl.h - OpenCL/cl.h pango-1.0/pango/pango-features.h unicode/uchar.h ) @@ -121,7 +119,6 @@ file(APPEND ${AUTOCONFIG_SRC} " #cmakedefine HAVE_NEON ${HAVE_NEON} #cmakedefine HAVE_LIBARCHIVE ${HAVE_LIBARCHIVE} #cmakedefine HAVE_LIBCURL ${HAVE_LIBCURL} -#cmakedefine USE_OPENCL ${USE_OPENCL} ") if(TESSDATA_PREFIX) diff --git a/cmake/SourceGroups.cmake b/cmake/SourceGroups.cmake index b21763c3..34e9e081 100644 --- a/cmake/SourceGroups.cmake +++ b/cmake/SourceGroups.cmake @@ -31,7 +31,6 @@ source_group("classify" "${SSRC}/classify/${H_CPP}") source_group("cutil" "${SSRC}/cutil/${H_CPP}") source_group("dict" "${SSRC}/dict/${H_CPP}") source_group("lstm" "${SSRC}/lstm/${H_CPP}") -source_group("opencl" "${SSRC}/opencl/${H_CPP}") source_group("textord" "${SSRC}/textord/${H_CPP}") source_group("viewer" "${SSRC}/viewer/${H_CPP}") source_group("wordrec" "${SSRC}/wordrec/${H_CPP}") diff --git a/configure.ac b/configure.ac index 4e257a01..9c44180b 100644 --- a/configure.ac +++ b/configure.ac @@ -79,8 +79,6 @@ AC_SUBST([AM_CPPFLAGS]) # Can be overridden with `configure --disable-silent-rules` or with `make V=1`. AM_SILENT_RULES([yes]) -OPENCL_INC="/opt/AMDAPP/include" -OPENCL_LIBS="-lOpenCL" ############################# # # Platform specific setup @@ -102,15 +100,12 @@ case "${host_os}" in AM_CONDITIONAL([ADD_RT], true) ;; *darwin*) - OPENCL_LIBS="" - OPENCL_INC="" AM_CONDITIONAL([ADD_RT], false) ;; *android*|openbsd*) AM_CONDITIONAL([ADD_RT], false) ;; powerpc-*-darwin*) - OPENCL_LIBS="" ;; *) # default @@ -263,19 +258,6 @@ have_tiff=false # Note that the first usage of AC_CHECK_HEADERS must be unconditional. AC_CHECK_HEADERS([tiffio.h], [have_tiff=true], [have_tiff=false]) -# check whether to build opencl version -AC_MSG_CHECKING([--enable-opencl argument]) -AC_ARG_ENABLE([opencl], - AS_HELP_STRING([--enable-opencl], [enable opencl build [default=no]])) -AC_MSG_RESULT([$enable_opencl]) -# check for opencl header -have_opencl=false -if test "$enable_opencl" = "yes"; then - AC_CHECK_HEADERS([CL/cl.h], [have_opencl=true], [ - AC_CHECK_HEADERS(OpenCL/cl.h, [have_opencl=true], [have_opencl=false]) - ]) -fi - # Configure arguments which allow disabling some optional libraries. AC_ARG_WITH([archive], AS_HELP_STRING([--with-archive], @@ -323,9 +305,6 @@ m4_define([MY_CHECK_FRAMEWORK], fi] ) -have_opencl_lib=false -OPENCL_CPPFLAGS='' -OPENCL_LDFLAGS='' case "${host_os}" in *darwin* | *-macos10*) MY_CHECK_FRAMEWORK([Accelerate]) @@ -333,36 +312,11 @@ case "${host_os}" in AM_CPPFLAGS="-DHAVE_FRAMEWORK_ACCELERATE $AM_CPPFLAGS" AM_LDFLAGS="$AM_LDFLAGS -framework Accelerate" fi - MY_CHECK_FRAMEWORK([OpenCL]) - if test "$enable_opencl" = "yes"; then - if test $my_cv_framework_OpenCL = no; then - AC_MSG_ERROR([Required OpenCL library not found!]) - fi - AM_CPPFLAGS="-DUSE_OPENCL $AM_CPPFLAGS" - OPENCL_CPPFLAGS="" - OPENCL_LDFLAGS="-framework OpenCL" - fi ;; *) # default - if test "$enable_opencl" = "yes"; then - AC_CHECK_LIB([OpenCL], [clGetPlatformIDs], - [have_opencl_lib=true], [have_opencl_lib=false]) - if !($have_opencl); then - AC_MSG_ERROR([Required OpenCL headers not found!]) - fi - if !($have_opencl_lib); then - AC_MSG_ERROR([Required OpenCL library not found!]) - fi - AM_CPPFLAGS="-DUSE_OPENCL $AM_CPPFLAGS" - OPENCL_CPPFLAGS="-I${OPENCL_INC}" - OPENCL_LDFLAGS="${OPENCL_LIBS}" - fi ;; esac -AM_CONDITIONAL([OPENCL], [test "$enable_opencl" = "yes"]) -AC_SUBST([OPENCL_CPPFLAGS]) -AC_SUBST([OPENCL_LDFLAGS]) # check whether to build tesseract with -fvisibility=hidden -fvisibility-inlines-hidden # http://gcc.gnu.org/wiki/Visibility diff --git a/include/tesseract/baseapi.h b/include/tesseract/baseapi.h index 535f20a1..7aae3a8e 100644 --- a/include/tesseract/baseapi.h +++ b/include/tesseract/baseapi.h @@ -86,15 +86,6 @@ public: */ static const char *Version(); - /** - * If compiled with OpenCL AND an available OpenCL - * device is deemed faster than serial code, then - * "device" is populated with the cl_device_id - * and returns sizeof(cl_device_id) - * otherwise *device=nullptr and returns 0. - */ - static size_t getOpenCLDevice(void **device); - /** * Set the name of the input file. Needed for training and * reading a UNLV zone file, and for searchable PDF output. diff --git a/include/tesseract/capi.h b/include/tesseract/capi.h index 40f4856a..cc95e30e 100644 --- a/include/tesseract/capi.h +++ b/include/tesseract/capi.h @@ -186,8 +186,6 @@ TESS_API int TessResultRendererImageNum(TessResultRenderer *renderer); TESS_API TessBaseAPI *TessBaseAPICreate(); TESS_API void TessBaseAPIDelete(TessBaseAPI *handle); -TESS_API size_t TessBaseAPIGetOpenCLDevice(TessBaseAPI *handle, void **device); - TESS_API void TessBaseAPISetInputName(TessBaseAPI *handle, const char *name); TESS_API const char *TessBaseAPIGetInputName(TessBaseAPI *handle); diff --git a/src/api/baseapi.cpp b/src/api/baseapi.cpp index c87d6eb6..11892772 100644 --- a/src/api/baseapi.cpp +++ b/src/api/baseapi.cpp @@ -41,9 +41,6 @@ #endif #include "mutableiterator.h" // for MutableIterator #include "normalis.h" // for kBlnBaselineOffset, kBlnXHeight -#if defined(USE_OPENCL) -# include "openclwrapper.h" // for OpenclDevice -#endif #include "pageres.h" // for PAGE_RES_IT, WERD_RES, PAGE_RES, CR_DE... #include "paragraphs.h" // for DetectParagraphs #include "params.h" // for BoolParam, IntParam, DoubleParam, Stri... @@ -243,27 +240,6 @@ const char *TessBaseAPI::Version() { return TESSERACT_VERSION_STR; } -/** - * If compiled with OpenCL AND an available OpenCL - * device is deemed faster than serial code, then - * "device" is populated with the cl_device_id - * and returns sizeof(cl_device_id) - * otherwise *device=nullptr and returns 0. - */ -size_t TessBaseAPI::getOpenCLDevice(void **data) { -#ifdef USE_OPENCL - ds_device device = OpenclDevice::getDeviceSelection(); - if (device.type == DS_DEVICE_OPENCL_DEVICE) { - *data = new cl_device_id; - memcpy(*data, &device.oclDeviceID, sizeof(cl_device_id)); - return sizeof(cl_device_id); - } -#endif - - *data = nullptr; - return 0; -} - /** * Set the name of the input file. Needed only for training and * loading a UNLV zone file. @@ -398,10 +374,6 @@ int TessBaseAPI::Init(const char *data, int data_size, const char *language, Ocr delete tesseract_; tesseract_ = nullptr; } -#ifdef USE_OPENCL - OpenclDevice od; - od.InitEnv(); -#endif bool reset_classifier = true; if (tesseract_ == nullptr) { reset_classifier = false; diff --git a/src/api/capi.cpp b/src/api/capi.cpp index 65c0e6e5..e16fa93b 100644 --- a/src/api/capi.cpp +++ b/src/api/capi.cpp @@ -137,10 +137,6 @@ void TessBaseAPIDelete(TessBaseAPI *handle) { delete handle; } -size_t TessBaseAPIGetOpenCLDevice(TessBaseAPI * /*handle*/, void **device) { - return TessBaseAPI::getOpenCLDevice(device); -} - void TessBaseAPISetInputName(TessBaseAPI *handle, const char *name) { handle->SetInputName(name); } diff --git a/src/ccmain/thresholder.cpp b/src/ccmain/thresholder.cpp index fa8e5688..44cf29b4 100644 --- a/src/ccmain/thresholder.cpp +++ b/src/ccmain/thresholder.cpp @@ -25,10 +25,6 @@ #include "thresholder.h" #include "tprintf.h" // for tprintf -#if defined(USE_OPENCL) -# include "openclwrapper.h" // for OpenclDevice -#endif - #include #include // for api->GetIntVariable() @@ -388,19 +384,7 @@ void ImageThresholder::OtsuThresholdRectToPix(Image src_pix, Image *out_pix) con int num_channels = OtsuThreshold(src_pix, rect_left_, rect_top_, rect_width_, rect_height_, thresholds, hi_values); - // only use opencl if compiled w/ OpenCL and selected device is opencl -#ifdef USE_OPENCL - OpenclDevice od; - if (num_channels == 4 && od.selectedDeviceIsOpenCL() && rect_top_ == 0 && rect_left_ == 0) { - od.ThresholdRectToPixOCL((unsigned char *)pixGetData(src_pix), num_channels, - pixGetWpl(src_pix) * 4, &thresholds[0], &hi_values[0], out_pix /*pix_OCL*/, - rect_height_, rect_width_, rect_top_, rect_left_); - } else { -#endif - ThresholdRectToPix(src_pix, num_channels, thresholds, hi_values, out_pix); -#ifdef USE_OPENCL - } -#endif + ThresholdRectToPix(src_pix, num_channels, thresholds, hi_values, out_pix); } /// Threshold the rectangle, taking everything except the src_pix diff --git a/src/ccstruct/otsuthr.cpp b/src/ccstruct/otsuthr.cpp index ddceecfb..e1a6ffe6 100644 --- a/src/ccstruct/otsuthr.cpp +++ b/src/ccstruct/otsuthr.cpp @@ -21,9 +21,6 @@ #include #include #include "helpers.h" -#if defined(USE_OPENCL) -# include "openclwrapper.h" // for OpenclDevice -#endif namespace tesseract { @@ -47,90 +44,40 @@ int OtsuThreshold(Image src_pix, int left, int top, int width, int height, std:: thresholds.resize(num_channels); hi_values.resize(num_channels); - // only use opencl if compiled w/ OpenCL and selected device is opencl -#ifdef USE_OPENCL - // all of channel 0 then all of channel 1... - std::vector histogramAllChannels(kHistogramSize * num_channels); - - // Calculate Histogram on GPU - OpenclDevice od; - if (od.selectedDeviceIsOpenCL() && (num_channels == 1 || num_channels == 4) && top == 0 && - left == 0) { - od.HistogramRectOCL(pixGetData(src_pix), num_channels, pixGetWpl(src_pix) * 4, left, top, width, - height, kHistogramSize, &histogramAllChannels[0]); - - // Calculate Threshold from Histogram on cpu - for (int ch = 0; ch < num_channels; ++ch) { - thresholds[ch] = -1; - hi_values[ch] = -1; - int *histogram = &histogramAllChannels[kHistogramSize * ch]; - int H; - int best_omega_0; - int best_t = OtsuStats(histogram, &H, &best_omega_0); - if (best_omega_0 == 0 || best_omega_0 == H) { - // This channel is empty. - continue; - } - // To be a convincing foreground we must have a small fraction of H - // or to be a convincing background we must have a large fraction of H. - // In between we assume this channel contains no thresholding information. - int hi_value = best_omega_0 < H * 0.5; - thresholds[ch] = best_t; - if (best_omega_0 > H * 0.75) { - any_good_hivalue = true; - hi_values[ch] = 0; - } else if (best_omega_0 < H * 0.25) { - any_good_hivalue = true; - hi_values[ch] = 1; - } else { - // In case all channels are like this, keep the best of the bad lot. - double hi_dist = hi_value ? (H - best_omega_0) : best_omega_0; - if (hi_dist > best_hi_dist) { - best_hi_dist = hi_dist; - best_hi_value = hi_value; - best_hi_index = ch; - } + for (int ch = 0; ch < num_channels; ++ch) { + thresholds[ch] = -1; + hi_values[ch] = -1; + // Compute the histogram of the image rectangle. + int histogram[kHistogramSize]; + HistogramRect(src_pix, ch, left, top, width, height, histogram); + int H; + int best_omega_0; + int best_t = OtsuStats(histogram, &H, &best_omega_0); + if (best_omega_0 == 0 || best_omega_0 == H) { + // This channel is empty. + continue; + } + // To be a convincing foreground we must have a small fraction of H + // or to be a convincing background we must have a large fraction of H. + // In between we assume this channel contains no thresholding information. + int hi_value = best_omega_0 < H * 0.5; + thresholds[ch] = best_t; + if (best_omega_0 > H * 0.75) { + any_good_hivalue = true; + hi_values[ch] = 0; + } else if (best_omega_0 < H * 0.25) { + any_good_hivalue = true; + hi_values[ch] = 1; + } else { + // In case all channels are like this, keep the best of the bad lot. + double hi_dist = hi_value ? (H - best_omega_0) : best_omega_0; + if (hi_dist > best_hi_dist) { + best_hi_dist = hi_dist; + best_hi_value = hi_value; + best_hi_index = ch; } } - } else { -#endif - for (int ch = 0; ch < num_channels; ++ch) { - thresholds[ch] = -1; - hi_values[ch] = -1; - // Compute the histogram of the image rectangle. - int histogram[kHistogramSize]; - HistogramRect(src_pix, ch, left, top, width, height, histogram); - int H; - int best_omega_0; - int best_t = OtsuStats(histogram, &H, &best_omega_0); - if (best_omega_0 == 0 || best_omega_0 == H) { - // This channel is empty. - continue; - } - // To be a convincing foreground we must have a small fraction of H - // or to be a convincing background we must have a large fraction of H. - // In between we assume this channel contains no thresholding information. - int hi_value = best_omega_0 < H * 0.5; - thresholds[ch] = best_t; - if (best_omega_0 > H * 0.75) { - any_good_hivalue = true; - hi_values[ch] = 0; - } else if (best_omega_0 < H * 0.25) { - any_good_hivalue = true; - hi_values[ch] = 1; - } else { - // In case all channels are like this, keep the best of the bad lot. - double hi_dist = hi_value ? (H - best_omega_0) : best_omega_0; - if (hi_dist > best_hi_dist) { - best_hi_dist = hi_dist; - best_hi_value = hi_value; - best_hi_index = ch; - } - } - } -#ifdef USE_OPENCL } -#endif // USE_OPENCL if (!any_good_hivalue) { // Use the best of the ones that were not good enough. diff --git a/src/opencl/oclkernels.h b/src/opencl/oclkernels.h deleted file mode 100644 index 87e5af50..00000000 --- a/src/opencl/oclkernels.h +++ /dev/null @@ -1,926 +0,0 @@ -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// http://www.apache.org/licenses/LICENSE-2.0 -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. - -#ifndef TESSERACT_OPENCL_OCLKERNELS_H_ -# define TESSERACT_OPENCL_OCLKERNELS_H_ - -# ifndef USE_EXTERNAL_KERNEL -# define KERNEL(...) # __VA_ARGS__ "\n" -// Double precision is a default of spreadsheets -// cl_khr_fp64: Khronos extension -// cl_amd_fp64: AMD extension -// use build option outside to define fp_t -///////////////////////////////////////////// -static const char *kernel_src = KERNEL( -\n #ifdef KHR_DP_EXTENSION\n -\n #pragma OPENCL EXTENSION cl_khr_fp64 - : enable\n -\n #elif AMD_DP_EXTENSION\n -\n #pragma OPENCL EXTENSION cl_amd_fp64 - : enable\n -\n #else \n -\n #endif \n - __kernel void composeRGBPixel(__global uint *tiffdata, int w, int h, int wpl, __global uint *output) { - int i = get_global_id(1); - int j = get_global_id(0); - int tiffword, rval, gval, bval; - - //Ignore the excess - if ((i >= h) || (j >= w)) - return; - - tiffword = tiffdata[i * w + j]; - rval = ((tiffword)&0xff); - gval = (((tiffword) >> 8) & 0xff); - bval = (((tiffword) >> 16) & 0xff); - output[i * wpl + j] = (rval << (8 * (sizeof(uint) - 1 - 0))) | (gval << (8 * (sizeof(uint) - 1 - 1))) | (bval << (8 * (sizeof(uint) - 1 - 2))); - }) - - KERNEL( -\n__kernel void pixSubtract_inplace(__global int *dword, __global int *sword, const int wpl, const int h) { - const unsigned int row = get_global_id(1); - const unsigned int col = get_global_id(0); - const unsigned int pos = row * wpl + col; - - //Ignore the execss - if (row >= h || col >= wpl) - return; - - *(dword + pos) &= ~(*(sword + pos)); - }\n) - - KERNEL( -\n__kernel void morphoDilateHor_5x5(__global int *sword, __global int *dword, const int wpl, const int h) { - const unsigned int pos = get_global_id(0); - unsigned int prevword, nextword, currword, tempword; - unsigned int destword; - const int col = pos % wpl; - - //Ignore the execss - if (pos >= (wpl * h)) - return; - - currword = *(sword + pos); - destword = currword; - - //Handle boundary conditions - if (col == 0) - prevword = 0; - else - prevword = *(sword + pos - 1); - - if (col == (wpl - 1)) - nextword = 0; - else - nextword = *(sword + pos + 1); - - //Loop unrolled - - //1 bit to left and 1 bit to right - //Get the max value on LHS of every pixel - tempword = (prevword << (31)) | ((currword >> 1)); - destword |= tempword; - //Get max value on RHS of every pixel - tempword = (currword << 1) | (nextword >> (31)); - destword |= tempword; - - //2 bit to left and 2 bit to right - //Get the max value on LHS of every pixel - tempword = (prevword << (30)) | ((currword >> 2)); - destword |= tempword; - //Get max value on RHS of every pixel - tempword = (currword << 2) | (nextword >> (30)); - destword |= tempword; - - *(dword + pos) = destword; - }\n) - - KERNEL( -\n__kernel void morphoDilateVer_5x5(__global int *sword, __global int *dword, const int wpl, const int h) { - const int col = get_global_id(0); - const int row = get_global_id(1); - const unsigned int pos = row * wpl + col; - unsigned int tempword; - unsigned int destword; - int i; - - //Ignore the execss - if (row >= h || col >= wpl) - return; - - destword = *(sword + pos); - - //2 words above - i = (row - 2) < 0 ? row : (row - 2); - tempword = *(sword + i * wpl + col); - destword |= tempword; - - //1 word above - i = (row - 1) < 0 ? row : (row - 1); - tempword = *(sword + i * wpl + col); - destword |= tempword; - - //1 word below - i = (row >= (h - 1)) ? row : (row + 1); - tempword = *(sword + i * wpl + col); - destword |= tempword; - - //2 words below - i = (row >= (h - 2)) ? row : (row + 2); - tempword = *(sword + i * wpl + col); - destword |= tempword; - - *(dword + pos) = destword; - }\n) - - KERNEL( -\n__kernel void morphoDilateHor(__global int *sword, __global int *dword, const int xp, const int xn, const int wpl, const int h) { - const int col = get_global_id(0); - const int row = get_global_id(1); - const unsigned int pos = row * wpl + col; - unsigned int parbitsxp, parbitsxn, nwords; - unsigned int destword, tempword, lastword, currword; - unsigned int lnextword, lprevword, rnextword, rprevword, firstword, secondword; - int i, j, siter, eiter; - - //Ignore the execss - if (pos >= (wpl * h) || (xn < 1 && xp < 1)) - return; - - currword = *(sword + pos); - destword = currword; - - parbitsxp = xp & 31; - parbitsxn = xn & 31; - nwords = xp >> 5; - - if (parbitsxp > 0) - nwords += 1; - else - parbitsxp = 31; - - siter = (col - nwords); - eiter = (col + nwords); - - //Get prev word - if (col == 0) - firstword = 0x0; - else - firstword = *(sword + pos - 1); - - //Get next word - if (col == (wpl - 1)) - secondword = 0x0; - else - secondword = *(sword + pos + 1); - - //Last partial bits on either side - for (i = 1; i <= parbitsxp; i++) { - //Get the max value on LHS of every pixel - tempword = ((i == parbitsxp) && (parbitsxp != parbitsxn)) ? 0x0 : (firstword << (32 - i)) | ((currword >> i)); - - destword |= tempword; - - //Get max value on RHS of every pixel - tempword = (currword << i) | (secondword >> (32 - i)); - destword |= tempword; - } - - //Return if halfwidth <= 1 word - if (nwords == 1) { - if (xn == 32) { - destword |= firstword; - } - if (xp == 32) { - destword |= secondword; - } - - *(dword + pos) = destword; - return; - } - - if (siter < 0) - firstword = 0x0; - else - firstword = *(sword + row * wpl + siter); - - if (eiter >= wpl) - lastword = 0x0; - else - lastword = *(sword + row * wpl + eiter); - - for (i = 1; i < nwords; i++) { - //Gets LHS words - if ((siter + i) < 0) - secondword = 0x0; - else - secondword = *(sword + row * wpl + siter + i); - - lprevword = firstword << (32 - parbitsxn) | secondword >> parbitsxn; - - firstword = secondword; - - if ((siter + i + 1) < 0) - secondword = 0x0; - else - secondword = *(sword + row * wpl + siter + i + 1); - - lnextword = firstword << (32 - parbitsxn) | secondword >> parbitsxn; - - //Gets RHS words - if ((eiter - i) >= wpl) - firstword = 0x0; - else - firstword = *(sword + row * wpl + eiter - i); - - rnextword = firstword << parbitsxp | lastword >> (32 - parbitsxp); - - lastword = firstword; - if ((eiter - i - 1) >= wpl) - firstword = 0x0; - else - firstword = *(sword + row * wpl + eiter - i - 1); - - rprevword = firstword << parbitsxp | lastword >> (32 - parbitsxp); - - for (j = 1; j < 32; j++) { - //OR LHS full words - tempword = (lprevword << j) | (lnextword >> (32 - j)); - destword |= tempword; - - //OR RHS full words - tempword = (rprevword << j) | (rnextword >> (32 - j)); - destword |= tempword; - } - - destword |= lprevword; - destword |= lnextword; - destword |= rprevword; - destword |= rnextword; - - lastword = firstword; - firstword = secondword; - } - - *(dword + pos) = destword; - }\n) - - KERNEL( -\n__kernel void morphoDilateHor_32word(__global int *sword, __global int *dword, const int halfwidth, const int wpl, const int h, const char isEven) { - const int col = get_global_id(0); - const int row = get_global_id(1); - const unsigned int pos = row * wpl + col; - unsigned int prevword, nextword, currword, tempword; - unsigned int destword; - int i; - - //Ignore the execss - if (pos >= (wpl * h)) - return; - - currword = *(sword + pos); - destword = currword; - - //Handle boundary conditions - if (col == 0) - prevword = 0; - else - prevword = *(sword + pos - 1); - - if (col == (wpl - 1)) - nextword = 0; - else - nextword = *(sword + pos + 1); - - for (i = 1; i <= halfwidth; i++) { - //Get the max value on LHS of every pixel - if (i == halfwidth && isEven) { - tempword = 0x0; - } else { - tempword = (prevword << (32 - i)) | ((currword >> i)); - } - - destword |= tempword; - - //Get max value on RHS of every pixel - tempword = (currword << i) | (nextword >> (32 - i)); - - destword |= tempword; - } - - *(dword + pos) = destword; - }\n) - - KERNEL( -\n__kernel void morphoDilateVer(__global int *sword, __global int *dword, const int yp, const int wpl, const int h, const int yn) { - const int col = get_global_id(0); - const int row = get_global_id(1); - const unsigned int pos = row * wpl + col; - unsigned int tempword; - unsigned int destword; - int i, siter, eiter; - - //Ignore the execss - if (row >= h || col >= wpl) - return; - - destword = *(sword + pos); - - //Set start position and end position considering the boundary conditions - siter = (row - yn) < 0 ? 0 : (row - yn); - eiter = (row >= (h - yp)) ? (h - 1) : (row + yp); - - for (i = siter; i <= eiter; i++) { - tempword = *(sword + i * wpl + col); - - destword |= tempword; - } - - *(dword + pos) = destword; - }\n) - - KERNEL( -\n__kernel void morphoErodeHor_5x5(__global int *sword, __global int *dword, const int wpl, const int h) { - const unsigned int pos = get_global_id(0); - unsigned int prevword, nextword, currword, tempword; - unsigned int destword; - const int col = pos % wpl; - - //Ignore the execss - if (pos >= (wpl * h)) - return; - - currword = *(sword + pos); - destword = currword; - - //Handle boundary conditions - if (col == 0) - prevword = 0xffffffff; - else - prevword = *(sword + pos - 1); - - if (col == (wpl - 1)) - nextword = 0xffffffff; - else - nextword = *(sword + pos + 1); - - //Loop unrolled - - //1 bit to left and 1 bit to right - //Get the min value on LHS of every pixel - tempword = (prevword << (31)) | ((currword >> 1)); - destword &= tempword; - //Get min value on RHS of every pixel - tempword = (currword << 1) | (nextword >> (31)); - destword &= tempword; - - //2 bit to left and 2 bit to right - //Get the min value on LHS of every pixel - tempword = (prevword << (30)) | ((currword >> 2)); - destword &= tempword; - //Get min value on RHS of every pixel - tempword = (currword << 2) | (nextword >> (30)); - destword &= tempword; - - *(dword + pos) = destword; - }\n) - - KERNEL( -\n__kernel void morphoErodeVer_5x5(__global int *sword, __global int *dword, const int wpl, const int h, const int fwmask, const int lwmask) { - const int col = get_global_id(0); - const int row = get_global_id(1); - const unsigned int pos = row * wpl + col; - unsigned int tempword; - unsigned int destword; - int i; - - //Ignore the execss - if (row >= h || col >= wpl) - return; - - destword = *(sword + pos); - - if (row < 2 || row >= (h - 2)) { - destword = 0x0; - } else { - //2 words above - //i = (row - 2) < 0 ? row : (row - 2); - i = (row - 2); - tempword = *(sword + i * wpl + col); - destword &= tempword; - - //1 word above - //i = (row - 1) < 0 ? row : (row - 1); - i = (row - 1); - tempword = *(sword + i * wpl + col); - destword &= tempword; - - //1 word below - //i = (row >= (h - 1)) ? row : (row + 1); - i = (row + 1); - tempword = *(sword + i * wpl + col); - destword &= tempword; - - //2 words below - //i = (row >= (h - 2)) ? row : (row + 2); - i = (row + 2); - tempword = *(sword + i * wpl + col); - destword &= tempword; - - if (col == 0) { - destword &= fwmask; - } - if (col == (wpl - 1)) { - destword &= lwmask; - } - } - - *(dword + pos) = destword; - }\n) - - KERNEL( -\n__kernel void morphoErodeHor(__global int *sword, __global int *dword, const int xp, const int xn, const int wpl, const int h, const char isAsymmetric, const int rwmask, const int lwmask) { - const int col = get_global_id(0); - const int row = get_global_id(1); - const unsigned int pos = row * wpl + col; - unsigned int parbitsxp, parbitsxn, nwords; - unsigned int destword, tempword, lastword, currword; - unsigned int lnextword, lprevword, rnextword, rprevword, firstword, secondword; - int i, j, siter, eiter; - - //Ignore the execss - if (pos >= (wpl * h) || (xn < 1 && xp < 1)) - return; - - currword = *(sword + pos); - destword = currword; - - parbitsxp = xp & 31; - parbitsxn = xn & 31; - nwords = xp >> 5; - - if (parbitsxp > 0) - nwords += 1; - else - parbitsxp = 31; - - siter = (col - nwords); - eiter = (col + nwords); - - //Get prev word - if (col == 0) - firstword = 0xffffffff; - else - firstword = *(sword + pos - 1); - - //Get next word - if (col == (wpl - 1)) - secondword = 0xffffffff; - else - secondword = *(sword + pos + 1); - - //Last partial bits on either side - for (i = 1; i <= parbitsxp; i++) { - //Get the max value on LHS of every pixel - tempword = (firstword << (32 - i)) | ((currword >> i)); - destword &= tempword; - - //Get max value on RHS of every pixel - tempword = ((i == parbitsxp) && (parbitsxp != parbitsxn)) ? 0xffffffff : (currword << i) | (secondword >> (32 - i)); - - //tempword = (currword << i) | (secondword >> (32 - i)); - destword &= tempword; - } - - //Return if halfwidth <= 1 word - if (nwords == 1) { - if (xp == 32) { - destword &= firstword; - } - if (xn == 32) { - destword &= secondword; - } - - //Clear boundary pixels - if (isAsymmetric) { - if (col == 0) - destword &= rwmask; - if (col == (wpl - 1)) - destword &= lwmask; - } - - *(dword + pos) = destword; - return; - } - - if (siter < 0) - firstword = 0xffffffff; - else - firstword = *(sword + row * wpl + siter); - - if (eiter >= wpl) - lastword = 0xffffffff; - else - lastword = *(sword + row * wpl + eiter); - - for (i = 1; i < nwords; i++) { - //Gets LHS words - if ((siter + i) < 0) - secondword = 0xffffffff; - else - secondword = *(sword + row * wpl + siter + i); - - lprevword = firstword << (32 - parbitsxp) | secondword >> (parbitsxp); - - firstword = secondword; - - if ((siter + i + 1) < 0) - secondword = 0xffffffff; - else - secondword = *(sword + row * wpl + siter + i + 1); - - lnextword = firstword << (32 - parbitsxp) | secondword >> (parbitsxp); - - //Gets RHS words - if ((eiter - i) >= wpl) - firstword = 0xffffffff; - else - firstword = *(sword + row * wpl + eiter - i); - - rnextword = firstword << parbitsxn | lastword >> (32 - parbitsxn); - - lastword = firstword; - if ((eiter - i - 1) >= wpl) - firstword = 0xffffffff; - else - firstword = *(sword + row * wpl + eiter - i - 1); - - rprevword = firstword << parbitsxn | lastword >> (32 - parbitsxn); - - for (j = 0; j < 32; j++) { - //OR LHS full words - tempword = (lprevword << j) | (lnextword >> (32 - j)); - destword &= tempword; - - //OR RHS full words - tempword = (rprevword << j) | (rnextword >> (32 - j)); - destword &= tempword; - } - - destword &= lprevword; - destword &= lnextword; - destword &= rprevword; - destword &= rnextword; - - lastword = firstword; - firstword = secondword; - } - - if (isAsymmetric) { - //Clear boundary pixels - if (col < (nwords - 1)) - destword = 0x0; - else if (col == (nwords - 1)) - destword &= rwmask; - else if (col > (wpl - nwords)) - destword = 0x0; - else if (col == (wpl - nwords)) - destword &= lwmask; - } - - *(dword + pos) = destword; - }\n) - - KERNEL( -\n__kernel void morphoErodeHor_32word(__global int *sword, __global int *dword, const int halfwidth, const int wpl, const int h, const char clearBoundPixH, const int rwmask, const int lwmask, const char isEven) { - const int col = get_global_id(0); - const int row = get_global_id(1); - const unsigned int pos = row * wpl + col; - unsigned int prevword, nextword, currword, tempword, destword; - int i; - - //Ignore the execss - if (pos >= (wpl * h)) - return; - - currword = *(sword + pos); - destword = currword; - - //Handle boundary conditions - if (col == 0) - prevword = 0xffffffff; - else - prevword = *(sword + pos - 1); - - if (col == (wpl - 1)) - nextword = 0xffffffff; - else - nextword = *(sword + pos + 1); - - for (i = 1; i <= halfwidth; i++) { - //Get the min value on LHS of every pixel - tempword = (prevword << (32 - i)) | ((currword >> i)); - - destword &= tempword; - - //Get min value on RHS of every pixel - if (i == halfwidth && isEven) { - tempword = 0xffffffff; - } else { - tempword = (currword << i) | (nextword >> (32 - i)); - } - - destword &= tempword; - } - - if (clearBoundPixH) { - if (col == 0) { - destword &= rwmask; - } else if (col == (wpl - 1)) { - destword &= lwmask; - } - } - - *(dword + pos) = destword; - }\n) - - KERNEL( -\n__kernel void morphoErodeVer(__global int *sword, __global int *dword, const int yp, const int wpl, const int h, const char clearBoundPixV, const int yn) { - const int col = get_global_id(0); - const int row = get_global_id(1); - const unsigned int pos = row * wpl + col; - unsigned int tempword, destword; - int i, siter, eiter; - - //Ignore the execss - if (row >= h || col >= wpl) - return; - - destword = *(sword + pos); - - //Set start position and end position considering the boundary conditions - siter = (row - yp) < 0 ? 0 : (row - yp); - eiter = (row >= (h - yn)) ? (h - 1) : (row + yn); - - for (i = siter; i <= eiter; i++) { - tempword = *(sword + i * wpl + col); - - destword &= tempword; - } - - //Clear boundary pixels - if (clearBoundPixV && ((row < yp) || ((h - row) <= yn))) { - destword = 0x0; - } - - *(dword + pos) = destword; - }\n) - - // HistogramRect Kernel: Accumulate - // assumes 4 channels, i.e., bytes_per_pixel = 4 - // assumes number of pixels is multiple of 8 - // data is laid out as - // ch0 ch1 ... - // bin0 bin1 bin2... bin0... - // rpt0,1,2...256 rpt0,1,2... - KERNEL( -\n #define HIST_REDUNDANCY 256\n -\n #define GROUP_SIZE 256\n -\n #define HIST_SIZE 256\n -\n #define NUM_CHANNELS 4\n -\n #define HR_UNROLL_SIZE 8 \n -\n #define HR_UNROLL_TYPE uchar8 \n - - __attribute__((reqd_work_group_size(256, 1, 1))) __kernel void kernel_HistogramRectAllChannels(__global const uchar8 *data, uint numPixels, __global uint *histBuffer) { - // declare variables - uchar8 pixels; - int threadOffset = get_global_id(0) % HIST_REDUNDANCY; - - // for each pixel/channel, accumulate in global memory - for (uint pc = get_global_id(0); pc < numPixels * NUM_CHANNELS / HR_UNROLL_SIZE; pc += get_global_size(0)) { - pixels = data[pc]; - // channel bin thread - atomic_inc(&histBuffer[0 * HIST_SIZE * HIST_REDUNDANCY + pixels.s0 * HIST_REDUNDANCY + threadOffset]); // ch0 - atomic_inc(&histBuffer[0 * HIST_SIZE * HIST_REDUNDANCY + pixels.s4 * HIST_REDUNDANCY + threadOffset]); // ch0 - atomic_inc(&histBuffer[1 * HIST_SIZE * HIST_REDUNDANCY + pixels.s1 * HIST_REDUNDANCY + threadOffset]); // ch1 - atomic_inc(&histBuffer[1 * HIST_SIZE * HIST_REDUNDANCY + pixels.s5 * HIST_REDUNDANCY + threadOffset]); // ch1 - atomic_inc(&histBuffer[2 * HIST_SIZE * HIST_REDUNDANCY + pixels.s2 * HIST_REDUNDANCY + threadOffset]); // ch2 - atomic_inc(&histBuffer[2 * HIST_SIZE * HIST_REDUNDANCY + pixels.s6 * HIST_REDUNDANCY + threadOffset]); // ch2 - atomic_inc(&histBuffer[3 * HIST_SIZE * HIST_REDUNDANCY + pixels.s3 * HIST_REDUNDANCY + threadOffset]); // ch3 - atomic_inc(&histBuffer[3 * HIST_SIZE * HIST_REDUNDANCY + pixels.s7 * HIST_REDUNDANCY + threadOffset]); // ch3 - } - }) - - KERNEL( - // NUM_CHANNELS = 1 - __attribute__((reqd_work_group_size(256, 1, 1))) __kernel void kernel_HistogramRectOneChannel(__global const uchar8 *data, uint numPixels, __global uint *histBuffer) { - // declare variables - uchar8 pixels; - int threadOffset = get_global_id(0) % HIST_REDUNDANCY; - - // for each pixel/channel, accumulate in global memory - for (uint pc = get_global_id(0); pc < numPixels / HR_UNROLL_SIZE; pc += get_global_size(0)) { - pixels = data[pc]; - // bin thread - atomic_inc(&histBuffer[pixels.s0 * HIST_REDUNDANCY + threadOffset]); - atomic_inc(&histBuffer[pixels.s1 * HIST_REDUNDANCY + threadOffset]); - atomic_inc(&histBuffer[pixels.s2 * HIST_REDUNDANCY + threadOffset]); - atomic_inc(&histBuffer[pixels.s3 * HIST_REDUNDANCY + threadOffset]); - atomic_inc(&histBuffer[pixels.s4 * HIST_REDUNDANCY + threadOffset]); - atomic_inc(&histBuffer[pixels.s5 * HIST_REDUNDANCY + threadOffset]); - atomic_inc(&histBuffer[pixels.s6 * HIST_REDUNDANCY + threadOffset]); - atomic_inc(&histBuffer[pixels.s7 * HIST_REDUNDANCY + threadOffset]); - } - }) - - // HistogramRect Kernel: Reduction - // only supports 4 channels - // each work group handles a single channel of a single histogram bin - KERNEL(__attribute__((reqd_work_group_size(256, 1, 1))) __kernel void kernel_HistogramRectAllChannelsReduction(int n, // unused pixel redundancy - __global uint *histBuffer, __global int *histResult) { - // declare variables - int channel = get_group_id(0) / HIST_SIZE; - int bin = get_group_id(0) % HIST_SIZE; - int value = 0; - - // accumulate in register - for (uint i = get_local_id(0); i < HIST_REDUNDANCY; i += GROUP_SIZE) { - value += histBuffer[channel * HIST_SIZE * HIST_REDUNDANCY + bin * HIST_REDUNDANCY + i]; - } - - // reduction in local memory - __local int localHist[GROUP_SIZE]; - localHist[get_local_id(0)] = value; - barrier(CLK_LOCAL_MEM_FENCE); - for (int stride = GROUP_SIZE / 2; stride >= 1; stride /= 2) { - if (get_local_id(0) < stride) { - value = localHist[get_local_id(0) + stride]; - } - barrier(CLK_LOCAL_MEM_FENCE); - if (get_local_id(0) < stride) { - localHist[get_local_id(0)] += value; - } - barrier(CLK_LOCAL_MEM_FENCE); - } - - // write reduction to final result - if (get_local_id(0) == 0) { - histResult[get_group_id(0)] = localHist[0]; - } - } // kernel_HistogramRectAllChannels - ) - - KERNEL( - // NUM_CHANNELS = 1 - __attribute__((reqd_work_group_size(256, 1, 1))) __kernel void kernel_HistogramRectOneChannelReduction(int n, // unused pixel redundancy - __global uint *histBuffer, __global int *histResult) { - // declare variables - // int channel = get_group_id(0)/HIST_SIZE; - int bin = get_group_id(0) % HIST_SIZE; - int value = 0; - - // accumulate in register - for (int i = get_local_id(0); i < HIST_REDUNDANCY; i += GROUP_SIZE) { - value += histBuffer[bin * HIST_REDUNDANCY + i]; - } - - // reduction in local memory - __local int localHist[GROUP_SIZE]; - localHist[get_local_id(0)] = value; - barrier(CLK_LOCAL_MEM_FENCE); - for (int stride = GROUP_SIZE / 2; stride >= 1; stride /= 2) { - if (get_local_id(0) < stride) { - value = localHist[get_local_id(0) + stride]; - } - barrier(CLK_LOCAL_MEM_FENCE); - if (get_local_id(0) < stride) { - localHist[get_local_id(0)] += value; - } - barrier(CLK_LOCAL_MEM_FENCE); - } - - // write reduction to final result - if (get_local_id(0) == 0) { - histResult[get_group_id(0)] = localHist[0]; - } - } // kernel_HistogramRectOneChannelReduction - ) - - // ThresholdRectToPix Kernel - // only supports 4 channels - // imageData is input image (24-bits/pixel) - // pix is output image (1-bit/pixel) - KERNEL( -\n #define CHAR_VEC_WIDTH 4 \n -\n #define PIXELS_PER_WORD 32 \n -\n #define PIXELS_PER_BURST 8 \n -\n #define BURSTS_PER_WORD (PIXELS_PER_WORD)/(PIXELS_PER_BURST) \n -\n typedef union { - uchar s[PIXELS_PER_BURST * NUM_CHANNELS]; - uchar4 v[(PIXELS_PER_BURST * NUM_CHANNELS) / CHAR_VEC_WIDTH]; - } charVec; - - __attribute__((reqd_work_group_size(256, 1, 1))) __kernel void kernel_ThresholdRectToPix(__global const uchar4 *imageData, int height, int width, - int wpl, // words per line - __global int *thresholds, __global int *hi_values, __global int *pix) { - // declare variables - int pThresholds[NUM_CHANNELS]; - int pHi_Values[NUM_CHANNELS]; - for (int i = 0; i < NUM_CHANNELS; i++) { - pThresholds[i] = thresholds[i]; - pHi_Values[i] = hi_values[i]; - } - - // for each word (32 pixels) in output image - for (uint w = get_global_id(0); w < wpl * height; w += get_global_size(0)) { - unsigned int word = 0; // all bits start at zero - // for each burst in word - for (int b = 0; b < BURSTS_PER_WORD; b++) { - // load burst - charVec pixels; - int offset = (w / wpl) * width; - offset += (w % wpl) * PIXELS_PER_WORD; - offset += b * PIXELS_PER_BURST; - - for (int i = 0; i < PIXELS_PER_BURST; ++i) - pixels.v[i] = imageData[offset + i]; - - // for each pixel in burst - for (int p = 0; p < PIXELS_PER_BURST; p++) { - for (int c = 0; c < NUM_CHANNELS; c++) { - unsigned char pixChan = pixels.s[p * NUM_CHANNELS + c]; - if (pHi_Values[c] >= 0 && (pixChan > pThresholds[c]) == (pHi_Values[c] == 0)) { - const uint kTopBit = 0x80000000; - word |= (kTopBit >> ((b * PIXELS_PER_BURST + p) & 31)); - } - } - } - } - pix[w] = word; - } - } - -\n #define CHAR_VEC_WIDTH 8 \n -\n #define PIXELS_PER_WORD 32 \n -\n #define PIXELS_PER_BURST 8 \n -\n #define BURSTS_PER_WORD (PIXELS_PER_WORD) / (PIXELS_PER_BURST) \n -\n typedef union { - uchar s[PIXELS_PER_BURST * 1]; - uchar8 v[(PIXELS_PER_BURST * 1) / CHAR_VEC_WIDTH]; - } charVec1; - - __attribute__((reqd_work_group_size(256, 1, 1))) __kernel void kernel_ThresholdRectToPix_OneChan(__global const uchar8 *imageData, int height, int width, - int wpl, // words per line of output image - __global int *thresholds, __global int *hi_values, __global int *pix) { - // declare variables - int pThresholds[1]; - int pHi_Values[1]; - for (int i = 0; i < 1; i++) { - pThresholds[i] = thresholds[i]; - pHi_Values[i] = hi_values[i]; - } - - // for each word (32 pixels) in output image - for (uint w = get_global_id(0); w < wpl * height; w += get_global_size(0)) { - unsigned int word = 0; // all bits start at zero - - // for each burst in word - for (int b = 0; b < BURSTS_PER_WORD; b++) { - // load burst - charVec1 pixels; - // for each char8 in burst - pixels.v[0] = imageData[w * BURSTS_PER_WORD + b + 0]; - - // for each pixel in burst - for (int p = 0; p < PIXELS_PER_BURST; p++) { - //int littleEndianIdx = p ^ 3; - //int bigEndianIdx = p; - int idx = -\n #ifdef __ENDIAN_LITTLE__\n - p ^ - 3; - \n #else \n - p; - \n #endif \n unsigned char pixChan = pixels.s[idx]; - if (pHi_Values[0] >= 0 && (pixChan > pThresholds[0]) == (pHi_Values[0] == 0)) { - const uint kTopBit = 0x80000000; - word |= (kTopBit >> ((b * PIXELS_PER_BURST + p) & 31)); - } - } - } - pix[w] = word; - } - }) - - ; // close char* - -# endif // USE_EXTERNAL_KERNEL -#endif // TESSERACT_OPENCL_OCLKERNELS_H_ -/* vim:set shiftwidth=4 softtabstop=4 expandtab: */ diff --git a/src/opencl/openclwrapper.cpp b/src/opencl/openclwrapper.cpp deleted file mode 100644 index 8377d402..00000000 --- a/src/opencl/openclwrapper.cpp +++ /dev/null @@ -1,2466 +0,0 @@ -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// http://www.apache.org/licenses/LICENSE-2.0 -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. - -// Include automatically generated configuration file -#ifdef HAVE_CONFIG_H -# include "config_auto.h" -#endif - -#ifdef USE_OPENCL - -# ifdef _WIN32 -# include -# include -# else -# include -# include -# endif -# include -# include // for clock_gettime - -# include "oclkernels.h" -# include "openclwrapper.h" - -// for micro-benchmark -# include "otsuthr.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 -# include -# endif - -# include -# include -# include // for memset, strcpy, ... -# include - -# include "errcode.h" // for ASSERT_HOST -# include "image.h" // for Image - -namespace tesseract { - -GPUEnv OpenclDevice::gpuEnv; - -bool OpenclDevice::deviceIsSelected = false; -ds_device OpenclDevice::selectedDevice; - -int OpenclDevice::isInited = 0; - -static l_int32 MORPH_BC = ASYMMETRIC_MORPH_BC; - -static const l_uint32 lmask32[] = { - 0x80000000, 0xc0000000, 0xe0000000, 0xf0000000, 0xf8000000, 0xfc000000, 0xfe000000, 0xff000000, - 0xff800000, 0xffc00000, 0xffe00000, 0xfff00000, 0xfff80000, 0xfffc0000, 0xfffe0000, 0xffff0000, - 0xffff8000, 0xffffc000, 0xffffe000, 0xfffff000, 0xfffff800, 0xfffffc00, 0xfffffe00, 0xffffff00, - 0xffffff80, 0xffffffc0, 0xffffffe0, 0xfffffff0, 0xfffffff8, 0xfffffffc, 0xfffffffe, 0xffffffff}; - -static const l_uint32 rmask32[] = { - 0x00000001, 0x00000003, 0x00000007, 0x0000000f, 0x0000001f, 0x0000003f, 0x0000007f, 0x000000ff, - 0x000001ff, 0x000003ff, 0x000007ff, 0x00000fff, 0x00001fff, 0x00003fff, 0x00007fff, 0x0000ffff, - 0x0001ffff, 0x0003ffff, 0x0007ffff, 0x000fffff, 0x001fffff, 0x003fffff, 0x007fffff, 0x00ffffff, - 0x01ffffff, 0x03ffffff, 0x07ffffff, 0x0fffffff, 0x1fffffff, 0x3fffffff, 0x7fffffff, 0xffffffff}; - -static cl_mem pixsCLBuffer, pixdCLBuffer, - pixdCLIntermediate; // Morph operations buffers -static cl_mem pixThBuffer; // output from thresholdtopix calculation -static cl_int clStatus; -static KernelEnv rEnv; - -# define DS_TAG_VERSION "" -# define DS_TAG_VERSION_END "" -# define DS_TAG_DEVICE "" -# define DS_TAG_DEVICE_END "" -# define DS_TAG_SCORE "" -# define DS_TAG_SCORE_END "" -# define DS_TAG_DEVICE_TYPE "" -# define DS_TAG_DEVICE_TYPE_END "" -# define DS_TAG_DEVICE_NAME "" -# define DS_TAG_DEVICE_NAME_END "" -# define DS_TAG_DEVICE_DRIVER_VERSION "" -# define DS_TAG_DEVICE_DRIVER_VERSION_END "" - -# define DS_DEVICE_NATIVE_CPU_STRING "native_cpu" - -# define DS_DEVICE_NAME_LENGTH 256 - -enum ds_evaluation_type { DS_EVALUATE_ALL, DS_EVALUATE_NEW_ONLY }; - -struct ds_profile { - std::vector devices; - unsigned int numDevices; - const char *version; -}; - -enum ds_status { - DS_SUCCESS = 0, - DS_INVALID_PROFILE = 1000, - DS_MEMORY_ERROR, - DS_INVALID_PERF_EVALUATOR_TYPE, - DS_INVALID_PERF_EVALUATOR, - DS_PERF_EVALUATOR_ERROR, - DS_FILE_ERROR, - DS_UNKNOWN_DEVICE_TYPE, - DS_PROFILE_FILE_ERROR, - DS_SCORE_SERIALIZER_ERROR, - DS_SCORE_DESERIALIZER_ERROR -}; - -// Pointer to a function that calculates the score of a device (ex: -// device->score) update the data size of score. The encoding and the format -// of the score data is implementation defined. The function should return -// DS_SUCCESS if there's no error to be reported. -typedef ds_status (*ds_perf_evaluator)(ds_device *device, void *data); - -// deallocate memory used by score -typedef ds_status (*ds_score_release)(TessDeviceScore *score); - -static ds_status releaseDSProfile(ds_profile *profile, ds_score_release sr) { - ds_status status = DS_SUCCESS; - if (profile != nullptr) { - if (sr != nullptr) { - unsigned int i; - for (i = 0; i < profile->numDevices; i++) { - free(profile->devices[i].oclDeviceName); - free(profile->devices[i].oclDriverVersion); - status = sr(profile->devices[i].score); - if (status != DS_SUCCESS) - break; - } - } - delete profile; - } - return status; -} - -static ds_status initDSProfile(ds_profile **p, const char *version) { - int numDevices; - cl_uint numPlatforms; - std::vector platforms; - std::vector devices; - ds_status status = DS_SUCCESS; - unsigned int next; - unsigned int i; - - if (p == nullptr) - return DS_INVALID_PROFILE; - - ds_profile *profile = new ds_profile; - - memset(profile, 0, sizeof(ds_profile)); - - clGetPlatformIDs(0, nullptr, &numPlatforms); - - if (numPlatforms > 0) { - platforms.resize(numPlatforms); - clGetPlatformIDs(numPlatforms, platforms.data(), nullptr); - } - - numDevices = 0; - for (i = 0; i < numPlatforms; i++) { - cl_uint num; - clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_ALL, 0, nullptr, &num); - numDevices += num; - } - - if (numDevices > 0) { - devices.resize(numDevices); - } - - profile->numDevices = numDevices + 1; // +1 to numDevices to include the native CPU - profile->devices.resize(profile->numDevices); - - next = 0; - for (i = 0; i < numPlatforms; i++) { - cl_uint num; - unsigned j; - clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_ALL, numDevices, &devices[0], &num); - for (j = 0; j < num; j++, next++) { - char buffer[DS_DEVICE_NAME_LENGTH]; - size_t length; - - 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); - 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); - length = strlen(buffer); - profile->devices[next].oclDriverVersion = (char *)malloc(length + 1); - memcpy(profile->devices[next].oclDriverVersion, buffer, length + 1); - } - } - profile->devices[next].type = DS_DEVICE_NATIVE_CPU; - profile->version = version; - - *p = profile; - return status; -} - -static ds_status profileDevices(ds_profile *profile, const ds_evaluation_type type, - ds_perf_evaluator evaluator, void *evaluatorData, - unsigned int *numUpdates) { - ds_status status = DS_SUCCESS; - unsigned int i; - unsigned int updates = 0; - - if (profile == nullptr) { - return DS_INVALID_PROFILE; - } - if (evaluator == nullptr) { - return DS_INVALID_PERF_EVALUATOR; - } - - for (i = 0; i < profile->numDevices; i++) { - ds_status evaluatorStatus; - - switch (type) { - case DS_EVALUATE_NEW_ONLY: - if (profile->devices[i].score != nullptr) - break; - // else fall through - case DS_EVALUATE_ALL: - evaluatorStatus = evaluator(&profile->devices[i], evaluatorData); - if (evaluatorStatus != DS_SUCCESS) { - status = evaluatorStatus; - return status; - } - updates++; - break; - default: - return DS_INVALID_PERF_EVALUATOR_TYPE; - break; - }; - } - if (numUpdates) - *numUpdates = updates; - return status; -} - -static const char *findString(const char *contentStart, const char *contentEnd, - const char *string) { - size_t stringLength; - const char *currentPosition; - const char *found = nullptr; - stringLength = strlen(string); - currentPosition = contentStart; - for (currentPosition = contentStart; currentPosition < contentEnd; currentPosition++) { - if (*currentPosition == string[0]) { - if (currentPosition + stringLength < contentEnd) { - if (strncmp(currentPosition, string, stringLength) == 0) { - found = currentPosition; - break; - } - } - } - } - return found; -} - -static ds_status readProFile(const char *fileName, char **content, size_t *contentSize) { - *contentSize = 0; - *content = nullptr; - ds_status status = DS_SUCCESS; - FILE *input = fopen(fileName, "rb"); - if (input == nullptr) { - status = DS_FILE_ERROR; - } else { - fseek(input, 0L, SEEK_END); - auto pos = std::ftell(input); - rewind(input); - if (pos > 0) { - size_t size = pos; - char *binary = new char[size]; - if (fread(binary, sizeof(char), size, input) != size) { - status = DS_FILE_ERROR; - delete[] binary; - } else { - *contentSize = size; - *content = binary; - } - } - fclose(input); - } - return status; -} - -typedef ds_status (*ds_score_deserializer)(ds_device *device, const uint8_t *serializedScore, - unsigned int serializedScoreSize); - -static ds_status readProfileFromFile(ds_profile *profile, ds_score_deserializer deserializer, - const char *file) { - ds_status status = DS_SUCCESS; - char *contentStart; - size_t contentSize; - - if (profile == nullptr) - return DS_INVALID_PROFILE; - - status = readProFile(file, &contentStart, &contentSize); - if (status == DS_SUCCESS) { - const char *currentPosition; - const char *dataStart; - const char *dataEnd; - - const char *contentEnd = contentStart + contentSize; - currentPosition = contentStart; - - // parse the version string - dataStart = findString(currentPosition, contentEnd, DS_TAG_VERSION); - if (dataStart == nullptr) { - status = DS_PROFILE_FILE_ERROR; - goto cleanup; - } - dataStart += strlen(DS_TAG_VERSION); - - dataEnd = findString(dataStart, contentEnd, DS_TAG_VERSION_END); - if (dataEnd == nullptr) { - status = DS_PROFILE_FILE_ERROR; - goto cleanup; - } - - size_t versionStringLength = strlen(profile->version); - if (versionStringLength + dataStart != dataEnd || - strncmp(profile->version, dataStart, versionStringLength) != 0) { - // version mismatch - status = DS_PROFILE_FILE_ERROR; - goto cleanup; - } - currentPosition = dataEnd + strlen(DS_TAG_VERSION_END); - - // parse the device information - while (1) { - unsigned int i; - - const char *deviceTypeStart; - const char *deviceTypeEnd; - ds_device_type deviceType; - - const char *deviceNameStart; - const char *deviceNameEnd; - - const char *deviceScoreStart; - const char *deviceScoreEnd; - - const char *deviceDriverStart; - const char *deviceDriverEnd; - - dataStart = findString(currentPosition, contentEnd, DS_TAG_DEVICE); - if (dataStart == nullptr) { - // nothing useful remain, quit... - break; - } - dataStart += strlen(DS_TAG_DEVICE); - dataEnd = findString(dataStart, contentEnd, DS_TAG_DEVICE_END); - 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) { - 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) { - status = DS_PROFILE_FILE_ERROR; - goto cleanup; - } - memcpy(&deviceType, deviceTypeStart, sizeof(ds_device_type)); - - // parse the device name - if (deviceType == DS_DEVICE_OPENCL_DEVICE) { - deviceNameStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_NAME); - 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) { - status = DS_PROFILE_FILE_ERROR; - goto cleanup; - } - - deviceDriverStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_DRIVER_VERSION); - 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) { - status = DS_PROFILE_FILE_ERROR; - goto cleanup; - } - - // check if this device is on the system - for (i = 0; i < profile->numDevices; i++) { - if (profile->devices[i].type == DS_DEVICE_OPENCL_DEVICE) { - size_t actualDeviceNameLength; - size_t driverVersionLength; - - actualDeviceNameLength = strlen(profile->devices[i].oclDeviceName); - driverVersionLength = strlen(profile->devices[i].oclDriverVersion); - if (deviceNameStart + actualDeviceNameLength == deviceNameEnd && - deviceDriverStart + driverVersionLength == deviceDriverEnd && - strncmp(profile->devices[i].oclDeviceName, deviceNameStart, - actualDeviceNameLength) == 0 && - strncmp(profile->devices[i].oclDriverVersion, deviceDriverStart, - driverVersionLength) == 0) { - deviceScoreStart = findString(dataStart, contentEnd, DS_TAG_SCORE); - deviceScoreStart += strlen(DS_TAG_SCORE); - deviceScoreEnd = findString(deviceScoreStart, contentEnd, DS_TAG_SCORE_END); - status = deserializer(&profile->devices[i], (const unsigned char *)deviceScoreStart, - deviceScoreEnd - deviceScoreStart); - if (status != DS_SUCCESS) { - goto cleanup; - } - } - } - } - } else if (deviceType == DS_DEVICE_NATIVE_CPU) { - 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) { - status = DS_PROFILE_FILE_ERROR; - goto cleanup; - } - deviceScoreStart += strlen(DS_TAG_SCORE); - deviceScoreEnd = findString(deviceScoreStart, contentEnd, DS_TAG_SCORE_END); - status = deserializer(&profile->devices[i], (const unsigned char *)deviceScoreStart, - deviceScoreEnd - deviceScoreStart); - if (status != DS_SUCCESS) { - goto cleanup; - } - } - } - } - - // skip over the current one to find the next device - currentPosition = dataEnd + strlen(DS_TAG_DEVICE_END); - } - } -cleanup: - delete[] contentStart; - return status; -} - -typedef ds_status (*ds_score_serializer)(ds_device *device, uint8_t **serializedScore, - unsigned int *serializedScoreSize); -static ds_status writeProfileToFile(ds_profile *profile, ds_score_serializer serializer, - const char *file) { - ds_status status = DS_SUCCESS; - - if (profile == nullptr) - return DS_INVALID_PROFILE; - - FILE *profileFile = fopen(file, "wb"); - if (profileFile == nullptr) { - status = DS_FILE_ERROR; - } else { - unsigned int i; - - // write version string - fwrite(DS_TAG_VERSION, sizeof(char), strlen(DS_TAG_VERSION), profileFile); - fwrite(profile->version, sizeof(char), strlen(profile->version), profileFile); - fwrite(DS_TAG_VERSION_END, sizeof(char), strlen(DS_TAG_VERSION_END), profileFile); - fwrite("\n", sizeof(char), 1, profileFile); - - for (i = 0; i < profile->numDevices && status == DS_SUCCESS; i++) { - uint8_t *serializedScore; - unsigned int serializedScoreSize; - - fwrite(DS_TAG_DEVICE, sizeof(char), strlen(DS_TAG_DEVICE), profileFile); - - fwrite(DS_TAG_DEVICE_TYPE, sizeof(char), strlen(DS_TAG_DEVICE_TYPE), profileFile); - fwrite(&profile->devices[i].type, sizeof(ds_device_type), 1, profileFile); - fwrite(DS_TAG_DEVICE_TYPE_END, sizeof(char), strlen(DS_TAG_DEVICE_TYPE_END), profileFile); - - switch (profile->devices[i].type) { - case DS_DEVICE_NATIVE_CPU: { - // There's no need to emit a device name for the native CPU device. - /* -fwrite(DS_TAG_DEVICE_NAME, sizeof(char), strlen(DS_TAG_DEVICE_NAME), - profileFile); -fwrite(DS_DEVICE_NATIVE_CPU_STRING,sizeof(char), - strlen(DS_DEVICE_NATIVE_CPU_STRING), profileFile); -fwrite(DS_TAG_DEVICE_NAME_END, sizeof(char), - strlen(DS_TAG_DEVICE_NAME_END), profileFile); -*/ - } break; - case DS_DEVICE_OPENCL_DEVICE: { - fwrite(DS_TAG_DEVICE_NAME, sizeof(char), strlen(DS_TAG_DEVICE_NAME), profileFile); - fwrite(profile->devices[i].oclDeviceName, sizeof(char), - strlen(profile->devices[i].oclDeviceName), profileFile); - fwrite(DS_TAG_DEVICE_NAME_END, sizeof(char), strlen(DS_TAG_DEVICE_NAME_END), profileFile); - - fwrite(DS_TAG_DEVICE_DRIVER_VERSION, sizeof(char), strlen(DS_TAG_DEVICE_DRIVER_VERSION), - profileFile); - fwrite(profile->devices[i].oclDriverVersion, sizeof(char), - strlen(profile->devices[i].oclDriverVersion), profileFile); - fwrite(DS_TAG_DEVICE_DRIVER_VERSION_END, sizeof(char), - strlen(DS_TAG_DEVICE_DRIVER_VERSION_END), profileFile); - } break; - default: - status = DS_UNKNOWN_DEVICE_TYPE; - continue; - }; - - 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) { - fwrite(serializedScore, sizeof(char), serializedScoreSize, profileFile); - delete[] serializedScore; - } - fwrite(DS_TAG_SCORE_END, sizeof(char), strlen(DS_TAG_SCORE_END), profileFile); - fwrite(DS_TAG_DEVICE_END, sizeof(char), strlen(DS_TAG_DEVICE_END), profileFile); - fwrite("\n", sizeof(char), 1, profileFile); - } - fclose(profileFile); - } - return status; -} - -// substitute invalid characters in device name with _ -static void legalizeFileName(char *fileName) { - // tprintf("fileName: %s\n", fileName); - const char *invalidChars = "/\?:*\"><| "; // space is valid but can cause headaches - // for each invalid char - for (unsigned i = 0; i < strlen(invalidChars); i++) { - char invalidStr[4]; - invalidStr[0] = invalidChars[i]; - invalidStr[1] = '\0'; - // tprintf("eliminating %s\n", invalidStr); - // char *pos = strstr(fileName, invalidStr); - // 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)) { - // tprintf("\tfound: %s, ", pos); - pos[0] = '_'; - // tprintf("fileName: %s\n", fileName); - } - } -} - -static void populateGPUEnvFromDevice(GPUEnv *gpuInfo, cl_device_id device) { - // tprintf("[DS] populateGPUEnvFromDevice\n"); - size_t size; - gpuInfo->mnIsUserCreated = 1; - // 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); - CHECK_OPENCL(clStatus, "populateGPUEnv::getDeviceInfo(TYPE)"); - // platform - 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); - CHECK_OPENCL(clStatus, "populateGPUEnv::createContext"); - // queue - cl_command_queue_properties queueProperties = 0; - gpuInfo->mpCmdQueue = - clCreateCommandQueue(gpuInfo->mpContext, gpuInfo->mpDevID, queueProperties, &clStatus); - CHECK_OPENCL(clStatus, "populateGPUEnv::createCommandQueue"); -} - -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; - } - fprintf(stderr, "[OD] Load opencl.dll successful!\n"); -# endif - return 1; -} -int OpenclDevice::SetKernelEnv(KernelEnv *envInfo) { - envInfo->mpkContext = gpuEnv.mpContext; - envInfo->mpkCmdQueue = gpuEnv.mpCmdQueue; - envInfo->mpkProgram = gpuEnv.mpArryPrograms[0]; - - return 1; -} - -static cl_mem allocateZeroCopyBuffer(const KernelEnv &rEnv, l_uint32 *hostbuffer, size_t nElements, - cl_mem_flags flags, cl_int *pStatus) { - cl_mem membuffer = clCreateBuffer(rEnv.mpkContext, (cl_mem_flags)(flags), - nElements * sizeof(l_uint32), hostbuffer, pStatus); - - return membuffer; -} - -static Image mapOutputCLBuffer(const KernelEnv &rEnv, cl_mem clbuffer, Image pixd, Image pixs, - int elements, cl_mem_flags flags, bool memcopy = false, - bool sync = true) { - if (!pixd) { - if (memcopy) { - if ((pixd = pixCreateTemplate(pixs)) == nullptr) - tprintf("pixd not made\n"); - } else { - if ((pixd = pixCreateHeader(pixGetWidth(pixs), pixGetHeight(pixs), pixGetDepth(pixs))) == - nullptr) - tprintf("pixd not made\n"); - } - } - l_uint32 *pValues = - (l_uint32 *)clEnqueueMapBuffer(rEnv.mpkCmdQueue, clbuffer, CL_TRUE, flags, 0, - elements * sizeof(l_uint32), 0, nullptr, nullptr, nullptr); - - if (memcopy) { - memcpy(pixGetData(pixd), pValues, elements * sizeof(l_uint32)); - } else { - pixSetData(pixd, pValues); - } - - clEnqueueUnmapMemObject(rEnv.mpkCmdQueue, clbuffer, pValues, 0, nullptr, nullptr); - - if (sync) { - clFinish(rEnv.mpkCmdQueue); - } - - return pixd; -} - -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; -} - -int OpenclDevice::initMorphCLAllocations(l_int32 wpl, l_int32 h, Image pixs) { - SetKernelEnv(&rEnv); - - 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); - } else { - // Get data from the source image - l_uint32 *srcdata = reinterpret_cast(malloc(wpl * h * sizeof(l_uint32))); - memcpy(srcdata, pixGetData(pixs), wpl * h * sizeof(l_uint32)); - - pixsCLBuffer = allocateZeroCopyBuffer(rEnv, srcdata, wpl * h, CL_MEM_USE_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); - - return (int)clStatus; -} - -int OpenclDevice::InitEnv() { -// tprintf("[OD] OpenclDevice::InitEnv()\n"); -# ifdef SAL_WIN32 - while (1) { - if (1 == LoadOpencl()) - break; - } -# endif - // sets up environment, compiles programs - - InitOpenclRunEnv_DeviceSelection(0); - return 1; -} - -int OpenclDevice::ReleaseOpenclRunEnv() { - ReleaseOpenclEnv(&gpuEnv); -# ifdef SAL_WIN32 - FreeOpenclDll(); -# endif - return 1; -} - -inline int OpenclDevice::AddKernelConfig(int kCount, const char *kName) { - ASSERT_HOST(kCount > 0); - ASSERT_HOST(strlen(kName) < sizeof(gpuEnv.mArrykernelNames[kCount - 1])); - strcpy(gpuEnv.mArrykernelNames[kCount - 1], kName); - gpuEnv.mnKernelCount++; - return 0; -} - -int OpenclDevice::RegistOpenclKernel() { - if (!gpuEnv.mnIsUserCreated) - memset(&gpuEnv, 0, sizeof(gpuEnv)); - - gpuEnv.mnFileCount = 0; // argc; - gpuEnv.mnKernelCount = 0UL; - - AddKernelConfig(1, "oclAverageSub1"); - return 0; -} - -int OpenclDevice::InitOpenclRunEnv_DeviceSelection(int argc) { - if (!isInited) { - // after programs compiled, selects best device - ds_device bestDevice_DS = getDeviceSelection(); - cl_device_id bestDevice = bestDevice_DS.oclDeviceID; - // overwrite global static GPUEnv with new device - if (selectedDeviceIsOpenCL()) { - // tprintf("[DS] InitOpenclRunEnv_DS::Calling populateGPUEnvFromDevice() - // for selected device\n"); - populateGPUEnvFromDevice(&gpuEnv, bestDevice); - gpuEnv.mnFileCount = 0; // argc; - gpuEnv.mnKernelCount = 0UL; - CompileKernelFile(&gpuEnv, ""); - } else { - // tprintf("[DS] InitOpenclRunEnv_DS::Skipping populateGPUEnvFromDevice() - // b/c native cpu selected\n"); - } - isInited = 1; - } - return 0; -} - -OpenclDevice::OpenclDevice() { - // InitEnv(); -} - -OpenclDevice::~OpenclDevice() { - // ReleaseOpenclRunEnv(); -} - -int OpenclDevice::ReleaseOpenclEnv(GPUEnv *gpuInfo) { - int i = 0; - int clStatus = 0; - - if (!isInited) { - return 1; - } - - for (i = 0; i < gpuEnv.mnFileCount; i++) { - if (gpuEnv.mpArryPrograms[i]) { - clStatus = clReleaseProgram(gpuEnv.mpArryPrograms[i]); - CHECK_OPENCL(clStatus, "clReleaseProgram"); - gpuEnv.mpArryPrograms[i] = nullptr; - } - } - if (gpuEnv.mpCmdQueue) { - clReleaseCommandQueue(gpuEnv.mpCmdQueue); - gpuEnv.mpCmdQueue = nullptr; - } - if (gpuEnv.mpContext) { - clReleaseContext(gpuEnv.mpContext); - gpuEnv.mpContext = nullptr; - } - isInited = 0; - gpuInfo->mnIsUserCreated = 0; - delete[] gpuInfo->mpArryDevsID; - return 1; -} -int OpenclDevice::BinaryGenerated(const char *clFileName, FILE **fhandle) { - unsigned int i = 0; - cl_int clStatus; - int status = 0; - FILE *fd = nullptr; - char fileName[256]; - char cl_name[128]; - char deviceName[1024]; - clStatus = clGetDeviceInfo(gpuEnv.mpArryDevsID[i], CL_DEVICE_NAME, sizeof(deviceName), deviceName, - nullptr); - CHECK_OPENCL(clStatus, "clGetDeviceInfo"); - const char *str = strstr(clFileName, ".cl"); - memcpy(cl_name, clFileName, str - clFileName); - cl_name[str - clFileName] = '\0'; - snprintf(fileName, sizeof(fileName), "%s-%s.bin", cl_name, deviceName); - legalizeFileName(fileName); - fd = fopen(fileName, "rb"); - status = (fd != nullptr) ? 1 : 0; - if (fd != nullptr) { - *fhandle = fd; - } - return status; -} -int OpenclDevice::CachedOfKernerPrg(const GPUEnv *gpuEnvCached, const char *clFileName) { - int i; - for (i = 0; i < gpuEnvCached->mnFileCount; i++) { - if (strcasecmp(gpuEnvCached->mArryKnelSrcFile[i], clFileName) == 0) { - if (gpuEnvCached->mpArryPrograms[i] != nullptr) { - return 1; - } - } - } - - return 0; -} -int OpenclDevice::WriteBinaryToFile(const char *fileName, const char *birary, size_t numBytes) { - FILE *output = nullptr; - output = fopen(fileName, "wb"); - if (output == nullptr) { - return 0; - } - - fwrite(birary, sizeof(char), numBytes, output); - fclose(output); - - return 1; -} - -int OpenclDevice::GeneratBinFromKernelSource(cl_program program, const char *clFileName) { - unsigned int i = 0; - cl_int clStatus; - cl_uint numDevices; - - clStatus = - clGetProgramInfo(program, CL_PROGRAM_NUM_DEVICES, sizeof(numDevices), &numDevices, nullptr); - CHECK_OPENCL(clStatus, "clGetProgramInfo"); - - std::vector mpArryDevsID(numDevices); - - /* grab the handles to all of the devices in the program. */ - clStatus = clGetProgramInfo(program, CL_PROGRAM_DEVICES, sizeof(cl_device_id) * numDevices, - &mpArryDevsID[0], nullptr); - CHECK_OPENCL(clStatus, "clGetProgramInfo"); - - /* figure out the sizes of each of the binaries. */ - std::vector binarySizes(numDevices); - - clStatus = clGetProgramInfo(program, CL_PROGRAM_BINARY_SIZES, sizeof(size_t) * numDevices, - &binarySizes[0], nullptr); - CHECK_OPENCL(clStatus, "clGetProgramInfo"); - - /* copy over all of the generated binaries. */ - std::vector binaries(numDevices); - - for (i = 0; i < numDevices; i++) { - if (binarySizes[i] != 0) { - binaries[i] = new char[binarySizes[i]]; - } else { - binaries[i] = nullptr; - } - } - - clStatus = clGetProgramInfo(program, CL_PROGRAM_BINARIES, sizeof(char *) * numDevices, - &binaries[0], nullptr); - CHECK_OPENCL(clStatus, "clGetProgramInfo"); - - /* dump out each binary into its own separate file. */ - for (i = 0; i < numDevices; i++) { - if (binarySizes[i] != 0) { - char fileName[256]; - char cl_name[128]; - char deviceName[1024]; - clStatus = - clGetDeviceInfo(mpArryDevsID[i], CL_DEVICE_NAME, sizeof(deviceName), deviceName, nullptr); - CHECK_OPENCL(clStatus, "clGetDeviceInfo"); - - const char *str = strstr(clFileName, ".cl"); - memcpy(cl_name, clFileName, str - clFileName); - cl_name[str - clFileName] = '\0'; - snprintf(fileName, sizeof(fileName), "%s-%s.bin", cl_name, deviceName); - legalizeFileName(fileName); - if (!WriteBinaryToFile(fileName, binaries[i], binarySizes[i])) { - tprintf("[OD] write binary[%s] failed\n", fileName); - return 0; - } // else - tprintf("[OD] write binary[%s] successfully\n", fileName); - } - } - - // Release all resources and memory - for (i = 0; i < numDevices; i++) { - delete[] binaries[i]; - } - - return 1; -} - -int OpenclDevice::CompileKernelFile(GPUEnv *gpuInfo, const char *buildOption) { - cl_int clStatus = 0; - const char *source; - size_t source_size[1]; - int binary_status, binaryExisted, idx; - cl_uint numDevices; - FILE *fd, *fd1; - const char *filename = "kernel.cl"; - // fprintf(stderr, "[OD] CompileKernelFile ... \n"); - if (CachedOfKernerPrg(gpuInfo, filename) == 1) { - return 1; - } - - idx = gpuInfo->mnFileCount; - - source = kernel_src; - - source_size[0] = strlen(source); - binaryExisted = 0; - binaryExisted = BinaryGenerated(filename, &fd); // don't check for binary during microbenchmark - if (binaryExisted == 1) { - clStatus = clGetContextInfo(gpuInfo->mpContext, CL_CONTEXT_NUM_DEVICES, sizeof(numDevices), - &numDevices, nullptr); - CHECK_OPENCL(clStatus, "clGetContextInfo"); - - std::vector mpArryDevsID(numDevices); - bool b_error = fseek(fd, 0, SEEK_END) < 0; - auto pos = std::ftell(fd); - b_error |= (pos <= 0); - size_t length = pos; - b_error |= fseek(fd, 0, SEEK_SET) < 0; - if (b_error) { - fclose(fd); - return 0; - } - - std::vector binary(length + 2); - - memset(&binary[0], 0, length + 2); - b_error |= fread(&binary[0], 1, length, fd) != length; - - fclose(fd); - 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[0], nullptr); - CHECK_OPENCL(clStatus, "clGetContextInfo"); - // fprintf(stderr, "[OD] Create kernel from binary\n"); - const uint8_t *c_binary = &binary[0]; - gpuInfo->mpArryPrograms[idx] = - clCreateProgramWithBinary(gpuInfo->mpContext, numDevices, &mpArryDevsID[0], &length, - &c_binary, &binary_status, &clStatus); - CHECK_OPENCL(clStatus, "clCreateProgramWithBinary"); - } else { - // create a CL program using the kernel source - // fprintf(stderr, "[OD] Create kernel from source\n"); - gpuInfo->mpArryPrograms[idx] = - clCreateProgramWithSource(gpuInfo->mpContext, 1, &source, source_size, &clStatus); - CHECK_OPENCL(clStatus, "clCreateProgramWithSource"); - } - - if (gpuInfo->mpArryPrograms[idx] == (cl_program) nullptr) { - return 0; - } - - // char options[512]; - // create a cl program executable for all the devices specified - // tprintf("[OD] BuildProgram.\n"); - if (!gpuInfo->mnIsUserCreated) { - clStatus = clBuildProgram(gpuInfo->mpArryPrograms[idx], 1, gpuInfo->mpArryDevsID, buildOption, - nullptr, nullptr); - } else { - clStatus = clBuildProgram(gpuInfo->mpArryPrograms[idx], 1, &(gpuInfo->mpDevID), buildOption, - nullptr, nullptr); - } - if (clStatus != CL_SUCCESS) { - tprintf("BuildProgram error!\n"); - size_t length; - if (!gpuInfo->mnIsUserCreated) { - 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); - } - if (clStatus != CL_SUCCESS) { - tprintf("opencl create build log fail\n"); - return 0; - } - std::vector buildLog(length); - if (!gpuInfo->mnIsUserCreated) { - clStatus = clGetProgramBuildInfo(gpuInfo->mpArryPrograms[idx], gpuInfo->mpArryDevsID[0], - CL_PROGRAM_BUILD_LOG, length, &buildLog[0], &length); - } else { - clStatus = clGetProgramBuildInfo(gpuInfo->mpArryPrograms[idx], gpuInfo->mpDevID, - CL_PROGRAM_BUILD_LOG, length, &buildLog[0], &length); - } - if (clStatus != CL_SUCCESS) { - tprintf("opencl program build info fail\n"); - return 0; - } - - fd1 = fopen("kernel-build.log", "w+"); - if (fd1 != nullptr) { - fwrite(&buildLog[0], sizeof(char), length, fd1); - fclose(fd1); - } - - return 0; - } - - strcpy(gpuInfo->mArryKnelSrcFile[idx], filename); - if (binaryExisted == 0) { - GeneratBinFromKernelSource(gpuInfo->mpArryPrograms[idx], filename); - } - - gpuInfo->mnFileCount += 1; - return 1; -} - -l_uint32 *OpenclDevice::pixReadFromTiffKernel(l_uint32 *tiffdata, l_int32 w, l_int32 h, l_int32 wpl, - l_uint32 *line) { - cl_int clStatus; - KernelEnv rEnv; - size_t globalThreads[2]; - size_t localThreads[2]; - int gsize; - cl_mem valuesCl; - cl_mem outputCl; - - // global and local work dimensions for Horizontal pass - gsize = (w + GROUPSIZE_X - 1) / GROUPSIZE_X * GROUPSIZE_X; - globalThreads[0] = gsize; - gsize = (h + GROUPSIZE_Y - 1) / GROUPSIZE_Y * GROUPSIZE_Y; - globalThreads[1] = gsize; - localThreads[0] = GROUPSIZE_X; - localThreads[1] = GROUPSIZE_Y; - - SetKernelEnv(&rEnv); - - l_uint32 *pResult = (l_uint32 *)malloc(w * h * sizeof(l_uint32)); - rEnv.mpkKernel = clCreateKernel(rEnv.mpkProgram, "composeRGBPixel", &clStatus); - 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); - CHECK_OPENCL(clStatus, "clSetKernelArg"); - clStatus = clSetKernelArg(rEnv.mpkKernel, 1, sizeof(w), &w); - CHECK_OPENCL(clStatus, "clSetKernelArg"); - clStatus = clSetKernelArg(rEnv.mpkKernel, 2, sizeof(h), &h); - CHECK_OPENCL(clStatus, "clSetKernelArg"); - clStatus = clSetKernelArg(rEnv.mpkKernel, 3, sizeof(wpl), &wpl); - CHECK_OPENCL(clStatus, "clSetKernelArg"); - clStatus = clSetKernelArg(rEnv.mpkKernel, 4, sizeof(cl_mem), &outputCl); - CHECK_OPENCL(clStatus, "clSetKernelArg"); - - // Kernel enqueue - 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); - return pResult; -} - -// Morphology Dilate operation for 5x5 structuring element. Invokes the relevant -// OpenCL kernels -static cl_int pixDilateCL_55(l_int32 wpl, l_int32 h) { - size_t globalThreads[2]; - cl_mem pixtemp; - cl_int status; - int gsize; - size_t localThreads[2]; - - // Horizontal pass - gsize = (wpl * h + GROUPSIZE_HMORX - 1) / GROUPSIZE_HMORX * GROUPSIZE_HMORX; - globalThreads[0] = gsize; - globalThreads[1] = GROUPSIZE_HMORY; - localThreads[0] = GROUPSIZE_HMORX; - localThreads[1] = GROUPSIZE_HMORY; - - rEnv.mpkKernel = clCreateKernel(rEnv.mpkProgram, "morphoDilateHor_5x5", &status); - CHECK_OPENCL(status, "clCreateKernel morphoDilateHor_5x5"); - - status = clSetKernelArg(rEnv.mpkKernel, 0, sizeof(cl_mem), &pixsCLBuffer); - status = clSetKernelArg(rEnv.mpkKernel, 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); - - // Swap source and dest buffers - pixtemp = pixsCLBuffer; - pixsCLBuffer = pixdCLBuffer; - pixdCLBuffer = pixtemp; - - // Vertical - gsize = (wpl + GROUPSIZE_X - 1) / GROUPSIZE_X * GROUPSIZE_X; - globalThreads[0] = gsize; - gsize = (h + GROUPSIZE_Y - 1) / GROUPSIZE_Y * GROUPSIZE_Y; - globalThreads[1] = gsize; - localThreads[0] = GROUPSIZE_X; - localThreads[1] = GROUPSIZE_Y; - - rEnv.mpkKernel = clCreateKernel(rEnv.mpkProgram, "morphoDilateVer_5x5", &status); - CHECK_OPENCL(status, "clCreateKernel morphoDilateVer_5x5"); - - status = clSetKernelArg(rEnv.mpkKernel, 0, sizeof(cl_mem), &pixsCLBuffer); - status = clSetKernelArg(rEnv.mpkKernel, 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); - - return status; -} - -// Morphology Erode operation for 5x5 structuring element. Invokes the relevant -// OpenCL kernels -static cl_int pixErodeCL_55(l_int32 wpl, l_int32 h) { - size_t globalThreads[2]; - cl_mem pixtemp; - cl_int status; - int gsize; - l_uint32 fwmask, lwmask; - size_t localThreads[2]; - - lwmask = lmask32[31 - 2]; - fwmask = rmask32[31 - 2]; - - // Horizontal pass - gsize = (wpl * h + GROUPSIZE_HMORX - 1) / GROUPSIZE_HMORX * GROUPSIZE_HMORX; - globalThreads[0] = gsize; - globalThreads[1] = GROUPSIZE_HMORY; - localThreads[0] = GROUPSIZE_HMORX; - localThreads[1] = GROUPSIZE_HMORY; - - rEnv.mpkKernel = clCreateKernel(rEnv.mpkProgram, "morphoErodeHor_5x5", &status); - CHECK_OPENCL(status, "clCreateKernel morphoErodeHor_5x5"); - - status = clSetKernelArg(rEnv.mpkKernel, 0, sizeof(cl_mem), &pixsCLBuffer); - status = clSetKernelArg(rEnv.mpkKernel, 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); - - // Swap source and dest buffers - pixtemp = pixsCLBuffer; - pixsCLBuffer = pixdCLBuffer; - pixdCLBuffer = pixtemp; - - // Vertical - gsize = (wpl + GROUPSIZE_X - 1) / GROUPSIZE_X * GROUPSIZE_X; - globalThreads[0] = gsize; - gsize = (h + GROUPSIZE_Y - 1) / GROUPSIZE_Y * GROUPSIZE_Y; - globalThreads[1] = gsize; - localThreads[0] = GROUPSIZE_X; - localThreads[1] = GROUPSIZE_Y; - - rEnv.mpkKernel = clCreateKernel(rEnv.mpkProgram, "morphoErodeVer_5x5", &status); - CHECK_OPENCL(status, "clCreateKernel morphoErodeVer_5x5"); - - status = clSetKernelArg(rEnv.mpkKernel, 0, sizeof(cl_mem), &pixsCLBuffer); - status = clSetKernelArg(rEnv.mpkKernel, 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); - - return status; -} - -// Morphology Dilate operation. Invokes the relevant OpenCL kernels -static cl_int pixDilateCL(l_int32 hsize, l_int32 vsize, l_int32 wpl, l_int32 h) { - l_int32 xp, yp, xn, yn; - SEL *sel; - size_t globalThreads[2]; - cl_mem pixtemp; - cl_int status = 0; - int gsize; - size_t localThreads[2]; - char isEven; - - OpenclDevice::SetKernelEnv(&rEnv); - - if (hsize == 5 && vsize == 5) { - // Specific case for 5x5 - status = pixDilateCL_55(wpl, h); - return status; - } - - sel = selCreateBrick(vsize, hsize, vsize / 2, hsize / 2, SEL_HIT); - - selFindMaxTranslations(sel, &xp, &yp, &xn, &yn); - selDestroy(&sel); - // global and local work dimensions for Horizontal pass - gsize = (wpl + GROUPSIZE_X - 1) / GROUPSIZE_X * GROUPSIZE_X; - globalThreads[0] = gsize; - gsize = (h + GROUPSIZE_Y - 1) / GROUPSIZE_Y * GROUPSIZE_Y; - globalThreads[1] = gsize; - localThreads[0] = GROUPSIZE_X; - localThreads[1] = GROUPSIZE_Y; - - if (xp > 31 || xn > 31) { - // Generic case. - rEnv.mpkKernel = clCreateKernel(rEnv.mpkProgram, "morphoDilateHor", &status); - CHECK_OPENCL(status, "clCreateKernel morphoDilateHor"); - - 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(xn), &xn); - 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); - - if (yp > 0 || yn > 0) { - pixtemp = pixsCLBuffer; - pixsCLBuffer = pixdCLBuffer; - pixdCLBuffer = pixtemp; - } - } else if (xp > 0 || xn > 0) { - // Specific Horizontal pass kernel for half width < 32 - rEnv.mpkKernel = clCreateKernel(rEnv.mpkProgram, "morphoDilateHor_32word", &status); - CHECK_OPENCL(status, "clCreateKernel morphoDilateHor_32word"); - isEven = (xp != xn); - - 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, 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); - - if (yp > 0 || yn > 0) { - pixtemp = pixsCLBuffer; - pixsCLBuffer = pixdCLBuffer; - pixdCLBuffer = pixtemp; - } - } - - if (yp > 0 || yn > 0) { - rEnv.mpkKernel = clCreateKernel(rEnv.mpkProgram, "morphoDilateVer", &status); - CHECK_OPENCL(status, "clCreateKernel morphoDilateVer"); - - status = clSetKernelArg(rEnv.mpkKernel, 0, sizeof(cl_mem), &pixsCLBuffer); - status = clSetKernelArg(rEnv.mpkKernel, 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); - } - - return status; -} - -// Morphology Erode operation. Invokes the relevant OpenCL kernels -static cl_int pixErodeCL(l_int32 hsize, l_int32 vsize, l_uint32 wpl, l_uint32 h) { - l_int32 xp, yp, xn, yn; - SEL *sel; - size_t globalThreads[2]; - size_t localThreads[2]; - cl_mem pixtemp; - cl_int status = 0; - int gsize; - char isAsymmetric = (MORPH_BC == ASYMMETRIC_MORPH_BC); - l_uint32 rwmask, lwmask; - char isEven; - - sel = selCreateBrick(vsize, hsize, vsize / 2, hsize / 2, SEL_HIT); - - selFindMaxTranslations(sel, &xp, &yp, &xn, &yn); - selDestroy(&sel); - OpenclDevice::SetKernelEnv(&rEnv); - - if (hsize == 5 && vsize == 5 && isAsymmetric) { - // Specific kernel for 5x5 - status = pixErodeCL_55(wpl, h); - return status; - } - - lwmask = lmask32[31 - (xn & 31)]; - rwmask = rmask32[31 - (xp & 31)]; - - // global and local work dimensions for Horizontal pass - gsize = (wpl + GROUPSIZE_X - 1) / GROUPSIZE_X * GROUPSIZE_X; - globalThreads[0] = gsize; - gsize = (h + GROUPSIZE_Y - 1) / GROUPSIZE_Y * GROUPSIZE_Y; - globalThreads[1] = gsize; - localThreads[0] = GROUPSIZE_X; - localThreads[1] = GROUPSIZE_Y; - - // Horizontal Pass - if (xp > 31 || xn > 31) { - // Generic case. - rEnv.mpkKernel = clCreateKernel(rEnv.mpkProgram, "morphoErodeHor", &status); - - 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(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); - - if (yp > 0 || yn > 0) { - pixtemp = pixsCLBuffer; - pixsCLBuffer = pixdCLBuffer; - pixdCLBuffer = pixtemp; - } - } else if (xp > 0 || xn > 0) { - rEnv.mpkKernel = clCreateKernel(rEnv.mpkProgram, "morphoErodeHor_32word", &status); - isEven = (xp != xn); - - 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, 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); - - if (yp > 0 || yn > 0) { - pixtemp = pixsCLBuffer; - pixsCLBuffer = pixdCLBuffer; - pixdCLBuffer = pixtemp; - } - } - - // Vertical Pass - if (yp > 0 || yn > 0) { - rEnv.mpkKernel = clCreateKernel(rEnv.mpkProgram, "morphoErodeVer", &status); - CHECK_OPENCL(status, "clCreateKernel morphoErodeVer"); - - status = clSetKernelArg(rEnv.mpkKernel, 0, sizeof(cl_mem), &pixsCLBuffer); - status = clSetKernelArg(rEnv.mpkKernel, 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(isAsymmetric), &isAsymmetric); - status = clSetKernelArg(rEnv.mpkKernel, 6, sizeof(yn), &yn); - status = clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, rEnv.mpkKernel, 2, nullptr, globalThreads, - localThreads, 0, nullptr, nullptr); - } - - return status; -} - -// Morphology Open operation. Invokes the relevant OpenCL kernels -static cl_int pixOpenCL(l_int32 hsize, l_int32 vsize, l_int32 wpl, l_int32 h) { - cl_int status; - cl_mem pixtemp; - - // Erode followed by Dilate - status = pixErodeCL(hsize, vsize, wpl, h); - - pixtemp = pixsCLBuffer; - pixsCLBuffer = pixdCLBuffer; - pixdCLBuffer = pixtemp; - - status = pixDilateCL(hsize, vsize, wpl, h); - - return status; -} - -// Morphology Close operation. Invokes the relevant OpenCL kernels -static cl_int pixCloseCL(l_int32 hsize, l_int32 vsize, l_int32 wpl, l_int32 h) { - cl_int status; - cl_mem pixtemp; - - // Dilate followed by Erode - status = pixDilateCL(hsize, vsize, wpl, h); - - pixtemp = pixsCLBuffer; - pixsCLBuffer = pixdCLBuffer; - pixdCLBuffer = pixtemp; - - status = pixErodeCL(hsize, vsize, wpl, h); - - return status; -} - -// output = buffer1 & ~(buffer2) -static cl_int pixSubtractCL_work(l_uint32 wpl, l_uint32 h, cl_mem buffer1, cl_mem buffer2) { - 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; - - 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); - status = clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, rEnv.mpkKernel, 2, nullptr, globalThreads, - localThreads, 0, nullptr, nullptr); - - return status; -} - -// OpenCL implementation of Get Lines from pix function -// Note: Assumes the source and dest opencl buffer are initialized. No check -// done -void OpenclDevice::pixGetLinesCL(Image pixd, Image pixs, Image *pix_vline, Image *pix_hline, - Image *pixClosed, bool getpixClosed, l_int32 close_hsize, - l_int32 close_vsize, l_int32 open_hsize, l_int32 open_vsize, - l_int32 line_hsize, l_int32 line_vsize) { - l_uint32 wpl, h; - cl_mem pixtemp; - - wpl = pixGetWpl(pixs); - h = pixGetHeight(pixs); - - // First step : Close Morph operation: Dilate followed by Erode - clStatus = pixCloseCL(close_hsize, close_vsize, wpl, h); - - // Copy the Close output to CPU buffer - if (getpixClosed) { - *pixClosed = - mapOutputCLBuffer(rEnv, pixdCLBuffer, *pixClosed, pixs, wpl * h, CL_MAP_READ, true, false); - } - - // Store the output of close operation in an intermediate buffer - // this will be later used for pixsubtract - clStatus = clEnqueueCopyBuffer(rEnv.mpkCmdQueue, pixdCLBuffer, pixdCLIntermediate, 0, 0, - sizeof(int) * wpl * h, 0, nullptr, nullptr); - - // Second step: Open Operation - Erode followed by Dilate - pixtemp = pixsCLBuffer; - pixsCLBuffer = pixdCLBuffer; - pixdCLBuffer = pixtemp; - - clStatus = pixOpenCL(open_hsize, open_vsize, wpl, h); - - // Third step: Subtract : (Close - Open) - pixtemp = pixsCLBuffer; - pixsCLBuffer = pixdCLBuffer; - pixdCLBuffer = pixdCLIntermediate; - pixdCLIntermediate = pixtemp; - - clStatus = pixSubtractCL_work(wpl, h, pixdCLBuffer, pixsCLBuffer); - - // Store the output of Hollow operation in an intermediate buffer - // this will be later used - clStatus = clEnqueueCopyBuffer(rEnv.mpkCmdQueue, pixdCLBuffer, pixdCLIntermediate, 0, 0, - sizeof(int) * wpl * h, 0, nullptr, nullptr); - - pixtemp = pixsCLBuffer; - pixsCLBuffer = pixdCLBuffer; - pixdCLBuffer = pixtemp; - - // Fourth step: Get vertical line - // pixOpenBrick(nullptr, pix_hollow, 1, min_line_length); - clStatus = pixOpenCL(1, line_vsize, wpl, h); - - // Copy the vertical line output to CPU buffer - *pix_vline = - mapOutputCLBuffer(rEnv, pixdCLBuffer, *pix_vline, pixs, wpl * h, CL_MAP_READ, true, false); - - pixtemp = pixsCLBuffer; - pixsCLBuffer = pixdCLIntermediate; - pixdCLIntermediate = pixtemp; - - // Fifth step: Get horizontal line - // pixOpenBrick(nullptr, pix_hollow, min_line_length, 1); - clStatus = pixOpenCL(line_hsize, 1, wpl, h); - - // Copy the horizontal line output to CPU buffer - *pix_hline = - mapOutputCLBuffer(rEnv, pixdCLBuffer, *pix_hline, pixs, wpl * h, CL_MAP_READ, true, true); - - return; -} - -/************************************************************************* - * HistogramRect - * Otsu Thresholding Operations - * histogramAllChannels is laid out as all channel 0, then all channel 1... - * only supports 1 or 4 channels (bytes_per_pixel) - ************************************************************************/ -int OpenclDevice::HistogramRectOCL(void *imageData, int bytes_per_pixel, int bytes_per_line, - int left, // always 0 - int top, // always 0 - int width, int height, int kHistogramSize, - int *histogramAllChannels) { - cl_int clStatus; - int retVal = 0; - KernelEnv histKern; - SetKernelEnv(&histKern); - KernelEnv histRedKern; - SetKernelEnv(&histRedKern); - /* map imagedata to device as read only */ - // USE_HOST_PTR uses onion+ bus which is slowest option; also happens to be - // 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(histKern.mpkContext, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, - width * height * bytes_per_pixel * sizeof(char), imageData, &clStatus); - CHECK_OPENCL(clStatus, "clCreateBuffer imageBuffer"); - - /* setup work group size parameters */ - int block_size = 256; - cl_uint numCUs; - clStatus = clGetDeviceInfo(gpuEnv.mpDevID, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(numCUs), &numCUs, - nullptr); - CHECK_OPENCL(clStatus, "clCreateBuffer imageBuffer"); - - int requestedOccupancy = 10; - int numWorkGroups = numCUs * requestedOccupancy; - int numThreads = block_size * numWorkGroups; - size_t local_work_size[] = {static_cast(block_size)}; - size_t global_work_size[] = {static_cast(numThreads)}; - size_t red_global_work_size[] = { - static_cast(block_size * kHistogramSize * bytes_per_pixel)}; - - /* map histogramAllChannels as write only */ - - cl_mem histogramBuffer = clCreateBuffer( - histKern.mpkContext, CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR, - kHistogramSize * bytes_per_pixel * sizeof(int), histogramAllChannels, &clStatus); - CHECK_OPENCL(clStatus, "clCreateBuffer histogramBuffer"); - - /* intermediate histogram buffer */ - int histRed = 256; - int tmpHistogramBins = kHistogramSize * bytes_per_pixel * histRed; - - cl_mem tmpHistogramBuffer = - clCreateBuffer(histKern.mpkContext, CL_MEM_READ_WRITE, tmpHistogramBins * sizeof(cl_uint), - nullptr, &clStatus); - CHECK_OPENCL(clStatus, "clCreateBuffer tmpHistogramBuffer"); - - /* atomic sync buffer */ - int *zeroBuffer = new int[1]; - zeroBuffer[0] = 0; - cl_mem atomicSyncBuffer = - clCreateBuffer(histKern.mpkContext, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(cl_int), - zeroBuffer, &clStatus); - CHECK_OPENCL(clStatus, "clCreateBuffer atomicSyncBuffer"); - delete[] zeroBuffer; - // Create kernel objects based on bytes_per_pixel - if (bytes_per_pixel == 1) { - histKern.mpkKernel = - clCreateKernel(histKern.mpkProgram, "kernel_HistogramRectOneChannel", &clStatus); - CHECK_OPENCL(clStatus, "clCreateKernel kernel_HistogramRectOneChannel"); - - histRedKern.mpkKernel = clCreateKernel(histRedKern.mpkProgram, - "kernel_HistogramRectOneChannelReduction", &clStatus); - CHECK_OPENCL(clStatus, "clCreateKernel kernel_HistogramRectOneChannelReduction"); - } else { - histKern.mpkKernel = - clCreateKernel(histKern.mpkProgram, "kernel_HistogramRectAllChannels", &clStatus); - CHECK_OPENCL(clStatus, "clCreateKernel kernel_HistogramRectAllChannels"); - - histRedKern.mpkKernel = clCreateKernel(histRedKern.mpkProgram, - "kernel_HistogramRectAllChannelsReduction", &clStatus); - CHECK_OPENCL(clStatus, "clCreateKernel kernel_HistogramRectAllChannelsReduction"); - } - - void *ptr; - - // Initialize tmpHistogramBuffer buffer - 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); - - /* set kernel 1 arguments */ - 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); - CHECK_OPENCL(clStatus, "clSetKernelArg numPixels"); - 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); - CHECK_OPENCL(clStatus, "clSetKernelArg imageBuffer"); - clStatus = clSetKernelArg(histRedKern.mpkKernel, 1, sizeof(cl_mem), &tmpHistogramBuffer); - CHECK_OPENCL(clStatus, "clSetKernelArg tmpHistogramBuffer"); - clStatus = clSetKernelArg(histRedKern.mpkKernel, 2, sizeof(cl_mem), &histogramBuffer); - CHECK_OPENCL(clStatus, "clSetKernelArg histogramBuffer"); - - /* launch histogram */ - 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); - CHECK_OPENCL(clStatus, "clEnqueueNDRangeKernel kernel_HistogramRectAllChannelsReduction"); - clFinish(histRedKern.mpkCmdQueue); - if (clStatus != 0) { - retVal = -1; - } - - /* 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); - CHECK_OPENCL(clStatus, "clEnqueueMapBuffer histogramBuffer"); - if (clStatus != 0) { - retVal = -1; - } - clEnqueueUnmapMemObject(histRedKern.mpkCmdQueue, histogramBuffer, ptr, 0, nullptr, nullptr); - - clReleaseMemObject(histogramBuffer); - clReleaseMemObject(imageBuffer); - return retVal; -} - -/************************************************************************* - * Threshold the rectangle, taking everything except the image buffer pointer - * from the class, using thresholds/hi_values to the output IMAGE. - * only supports 1 or 4 channels - ************************************************************************/ -int OpenclDevice::ThresholdRectToPixOCL(unsigned char *imageData, int bytes_per_pixel, - int bytes_per_line, int *thresholds, int *hi_values, - Image *pix, int height, int width, int top, int left) { - int retVal = 0; - /* create pix result buffer */ - *pix = pixCreate(width, height, 1); - uint32_t *pixData = pixGetData(*pix); - int wpl = pixGetWpl(*pix); - int pixSize = wpl * height * sizeof(uint32_t); // number of pixels - - cl_int clStatus; - KernelEnv rEnv; - SetKernelEnv(&rEnv); - - /* setup work group size parameters */ - int block_size = 256; - cl_uint numCUs = 6; - clStatus = clGetDeviceInfo(gpuEnv.mpDevID, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(numCUs), &numCUs, - nullptr); - CHECK_OPENCL(clStatus, "clCreateBuffer imageBuffer"); - - int requestedOccupancy = 10; - int numWorkGroups = numCUs * requestedOccupancy; - int numThreads = block_size * numWorkGroups; - size_t local_work_size[] = {(size_t)block_size}; - size_t global_work_size[] = {(size_t)numThreads}; - - /* map imagedata to device as read only */ - // USE_HOST_PTR uses onion+ bus which is slowest option; also happens to be - // 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); - CHECK_OPENCL(clStatus, "clCreateBuffer imageBuffer"); - - /* map pix as write only */ - pixThBuffer = clCreateBuffer(rEnv.mpkContext, CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR, pixSize, - pixData, &clStatus); - 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); - 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); - CHECK_OPENCL(clStatus, "clCreateBuffer hiValuesBuffer"); - - /* compile kernel */ - if (bytes_per_pixel == 4) { - rEnv.mpkKernel = clCreateKernel(rEnv.mpkProgram, "kernel_ThresholdRectToPix", &clStatus); - CHECK_OPENCL(clStatus, "clCreateKernel kernel_ThresholdRectToPix"); - } else { - rEnv.mpkKernel = - clCreateKernel(rEnv.mpkProgram, "kernel_ThresholdRectToPix_OneChan", &clStatus); - CHECK_OPENCL(clStatus, "clCreateKernel kernel_ThresholdRectToPix_OneChan"); - } - - /* set kernel arguments */ - clStatus = clSetKernelArg(rEnv.mpkKernel, 0, sizeof(cl_mem), &imageBuffer); - CHECK_OPENCL(clStatus, "clSetKernelArg imageBuffer"); - clStatus = clSetKernelArg(rEnv.mpkKernel, 1, sizeof(int), &height); - CHECK_OPENCL(clStatus, "clSetKernelArg height"); - clStatus = clSetKernelArg(rEnv.mpkKernel, 2, sizeof(int), &width); - 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); - CHECK_OPENCL(clStatus, "clSetKernelArg thresholdsBuffer"); - clStatus = clSetKernelArg(rEnv.mpkKernel, 5, sizeof(cl_mem), &hiValuesBuffer); - CHECK_OPENCL(clStatus, "clSetKernelArg hiValuesBuffer"); - clStatus = clSetKernelArg(rEnv.mpkKernel, 6, sizeof(cl_mem), &pixThBuffer); - CHECK_OPENCL(clStatus, "clSetKernelArg pixThBuffer"); - - /* launch kernel & wait */ - 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); - if (clStatus != 0) { - tprintf("Setting return value to -1\n"); - retVal = -1; - } - /* map results back from gpu */ - void *ptr = 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); - - clReleaseMemObject(imageBuffer); - clReleaseMemObject(thresholdsBuffer); - clReleaseMemObject(hiValuesBuffer); - - return retVal; -} - -/****************************************************************************** - * Data Types for Device Selection - *****************************************************************************/ - -struct TessScoreEvaluationInputData { - int height; - int width; - int numChannels; - unsigned char *imageData; - Image pix; -}; - -static void populateTessScoreEvaluationInputData(TessScoreEvaluationInputData *input) { - srand(1); - // 8.5x11 inches @ 300dpi rounded to clean multiples - int height = 3328; // %256 - int width = 2560; // %512 - int numChannels = 4; - input->height = height; - input->width = width; - input->numChannels = numChannels; - unsigned char(*imageData4)[4] = (unsigned char(*)[4])malloc( - height * width * numChannels * sizeof(unsigned char)); // new unsigned char[4][height*width]; - input->imageData = (unsigned char *)&imageData4[0]; - - // zero out image - unsigned char pixelWhite[4] = {0, 0, 0, 255}; - unsigned char pixelBlack[4] = {255, 255, 255, 255}; - for (int p = 0; p < height * width; p++) { - // unsigned char tmp[4] = imageData4[0]; - imageData4[p][0] = pixelWhite[0]; - imageData4[p][1] = pixelWhite[1]; - imageData4[p][2] = pixelWhite[2]; - imageData4[p][3] = pixelWhite[3]; - } - // random lines to be eliminated - int maxLineWidth = 64; // pixels wide - int numLines = 10; - // vertical lines - for (int i = 0; i < numLines; i++) { - int lineWidth = rand() % maxLineWidth; - int vertLinePos = lineWidth + rand() % (width - 2 * lineWidth); - // tprintf("[PI] VerticalLine @ %i (w=%i)\n", vertLinePos, lineWidth); - for (int row = vertLinePos - lineWidth / 2; row < vertLinePos + lineWidth / 2; row++) { - for (int col = 0; col < height; col++) { - // imageData4[row*width+col] = pixelBlack; - imageData4[row * width + col][0] = pixelBlack[0]; - imageData4[row * width + col][1] = pixelBlack[1]; - imageData4[row * width + col][2] = pixelBlack[2]; - imageData4[row * width + col][3] = pixelBlack[3]; - } - } - } - // horizontal lines - for (int i = 0; i < numLines; i++) { - int lineWidth = rand() % maxLineWidth; - int horLinePos = lineWidth + rand() % (height - 2 * lineWidth); - // tprintf("[PI] HorizontalLine @ %i (w=%i)\n", horLinePos, lineWidth); - for (int row = 0; row < width; row++) { - for (int col = horLinePos - lineWidth / 2; col < horLinePos + lineWidth / 2; - col++) { // for (int row = vertLinePos-lineWidth/2; row < - // vertLinePos+lineWidth/2; row++) { - // tprintf("[PI] HoizLine pix @ (%3i, %3i)\n", row, col); - // imageData4[row*width+col] = pixelBlack; - imageData4[row * width + col][0] = pixelBlack[0]; - imageData4[row * width + col][1] = pixelBlack[1]; - imageData4[row * width + col][2] = pixelBlack[2]; - imageData4[row * width + col][3] = pixelBlack[3]; - } - } - } - // spots (noise, squares) - float fractionBlack = 0.1; // how much of the image should be blackened - int numSpots = (height * width) * fractionBlack / (maxLineWidth * maxLineWidth / 2 / 2); - for (int i = 0; i < numSpots; i++) { - int lineWidth = rand() % maxLineWidth; - int col = lineWidth + rand() % (width - 2 * lineWidth); - int row = lineWidth + rand() % (height - 2 * lineWidth); - // tprintf("[PI] Spot[%i/%i] @ (%3i, %3i)\n", i, numSpots, row, col ); - for (int r = row - lineWidth / 2; r < row + lineWidth / 2; r++) { - for (int c = col - lineWidth / 2; c < col + lineWidth / 2; c++) { - // tprintf("[PI] \tSpot[%i/%i] @ (%3i, %3i)\n", i, numSpots, r, c ); - // imageData4[row*width+col] = pixelBlack; - imageData4[r * width + c][0] = pixelBlack[0]; - imageData4[r * width + c][1] = pixelBlack[1]; - imageData4[r * width + c][2] = pixelBlack[2]; - imageData4[r * width + c][3] = pixelBlack[3]; - } - } - } - - input->pix = pixCreate(input->width, input->height, 8 * input->numChannels); -} - -struct TessDeviceScore { - float time; // small time means faster device - bool clError; // were there any opencl errors - bool valid; // was the correct response generated -}; - -/****************************************************************************** - * Micro Benchmarks for Device Selection - *****************************************************************************/ - -static double composeRGBPixelMicroBench(GPUEnv *env, TessScoreEvaluationInputData input, - ds_device_type type) { - double time = 0; -# if ON_WINDOWS - LARGE_INTEGER freq, time_funct_start, time_funct_end; - QueryPerformanceFrequency(&freq); -# elif ON_APPLE - mach_timebase_info_data_t info = {0, 0}; - mach_timebase_info(&info); - long long start, stop; -# else - timespec time_funct_start, time_funct_end; -# endif - // input data - l_uint32 *tiffdata = (l_uint32 *)input.imageData; // same size and random data; data doesn't - // change workload - - // function call - if (type == DS_DEVICE_OPENCL_DEVICE) { -# if ON_WINDOWS - QueryPerformanceCounter(&time_funct_start); -# elif ON_APPLE - start = mach_absolute_time(); -# else - clock_gettime(CLOCK_MONOTONIC, &time_funct_start); -# endif - - OpenclDevice::gpuEnv = *env; - int wpl = pixGetWpl(input.pix); - 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); -# elif ON_APPLE - stop = mach_absolute_time(); - time = ((stop - start) * (double)info.numer / info.denom) / 1.0E9; -# else - clock_gettime(CLOCK_MONOTONIC, &time_funct_end); - time = (time_funct_end.tv_sec - time_funct_start.tv_sec) * 1.0 + - (time_funct_end.tv_nsec - time_funct_start.tv_nsec) / 1000000000.0; -# endif - - } else { -# if ON_WINDOWS - QueryPerformanceCounter(&time_funct_start); -# elif ON_APPLE - start = mach_absolute_time(); -# else - clock_gettime(CLOCK_MONOTONIC, &time_funct_start); -# endif - Image pix = pixCreate(input.width, input.height, 32); - l_uint32 *pixData = pixGetData(pix); - int i, j; - int idx = 0; - for (i = 0; i < input.height; i++) { - for (j = 0; j < input.width; j++) { - l_uint32 tiffword = tiffdata[i * input.width + j]; - l_int32 rval = ((tiffword)&0xff); - l_int32 gval = (((tiffword) >> 8) & 0xff); - l_int32 bval = (((tiffword) >> 16) & 0xff); - l_uint32 value = (rval << 24) | (gval << 16) | (bval << 8); - pixData[idx] = value; - idx++; - } - } -# if ON_WINDOWS - QueryPerformanceCounter(&time_funct_end); - time = (time_funct_end.QuadPart - time_funct_start.QuadPart) / (double)(freq.QuadPart); -# elif ON_APPLE - stop = mach_absolute_time(); - time = ((stop - start) * (double)info.numer / info.denom) / 1.0E9; -# else - clock_gettime(CLOCK_MONOTONIC, &time_funct_end); - time = (time_funct_end.tv_sec - time_funct_start.tv_sec) * 1.0 + - (time_funct_end.tv_nsec - time_funct_start.tv_nsec) / 1000000000.0; -# endif - pix.destroy(); - } - - return time; -} - -static double histogramRectMicroBench(GPUEnv *env, TessScoreEvaluationInputData input, - ds_device_type type) { - double time; -# if ON_WINDOWS - LARGE_INTEGER freq, time_funct_start, time_funct_end; - QueryPerformanceFrequency(&freq); -# elif ON_APPLE - mach_timebase_info_data_t info = {0, 0}; - mach_timebase_info(&info); - long long start, stop; -# else - timespec time_funct_start, time_funct_end; -# endif - - const int left = 0; - const int top = 0; - int kHistogramSize = 256; - int bytes_per_line = input.width * input.numChannels; - int *histogramAllChannels = new int[kHistogramSize * input.numChannels]; - // function call - if (type == DS_DEVICE_OPENCL_DEVICE) { -# if ON_WINDOWS - QueryPerformanceCounter(&time_funct_start); -# elif ON_APPLE - start = mach_absolute_time(); -# else - clock_gettime(CLOCK_MONOTONIC, &time_funct_start); -# endif - - OpenclDevice::gpuEnv = *env; - int retVal = OpenclDevice::HistogramRectOCL(input.imageData, input.numChannels, bytes_per_line, - left, top, input.width, input.height, - kHistogramSize, histogramAllChannels); - -# if ON_WINDOWS - QueryPerformanceCounter(&time_funct_end); - time = (time_funct_end.QuadPart - time_funct_start.QuadPart) / (double)(freq.QuadPart); -# elif ON_APPLE - stop = mach_absolute_time(); - if (retVal == 0) { - time = ((stop - start) * (double)info.numer / info.denom) / 1.0E9; - } else { - time = FLT_MAX; - } -# else - clock_gettime(CLOCK_MONOTONIC, &time_funct_end); - time = (time_funct_end.tv_sec - time_funct_start.tv_sec) * 1.0 + - (time_funct_end.tv_nsec - time_funct_start.tv_nsec) / 1000000000.0; -# endif - } else { - int *histogram = new int[kHistogramSize]; -# if ON_WINDOWS - QueryPerformanceCounter(&time_funct_start); -# elif ON_APPLE - start = mach_absolute_time(); -# else - clock_gettime(CLOCK_MONOTONIC, &time_funct_start); -# endif - for (int ch = 0; ch < input.numChannels; ++ch) { - tesseract::HistogramRect(input.pix, input.numChannels, left, top, input.width, input.height, - histogram); - } -# if ON_WINDOWS - QueryPerformanceCounter(&time_funct_end); - time = (time_funct_end.QuadPart - time_funct_start.QuadPart) / (double)(freq.QuadPart); -# elif ON_APPLE - stop = mach_absolute_time(); - time = ((stop - start) * (double)info.numer / info.denom) / 1.0E9; -# else - clock_gettime(CLOCK_MONOTONIC, &time_funct_end); - time = (time_funct_end.tv_sec - time_funct_start.tv_sec) * 1.0 + - (time_funct_end.tv_nsec - time_funct_start.tv_nsec) / 1000000000.0; -# endif - delete[] histogram; - } - - // cleanup - delete[] histogramAllChannels; - return time; -} - -// Reproducing the ThresholdRectToPix native version -static void ThresholdRectToPix_Native(const unsigned char *imagedata, int bytes_per_pixel, - int bytes_per_line, const int *thresholds, - const int *hi_values, Image *pix) { - int top = 0; - int left = 0; - int width = pixGetWidth(*pix); - int height = pixGetHeight(*pix); - - *pix = pixCreate(width, height, 1); - 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; - for (int x = 0; x < width; ++x, linedata += bytes_per_pixel) { - bool white_result = true; - for (int ch = 0; ch < bytes_per_pixel; ++ch) { - if (hi_values[ch] >= 0 && (linedata[ch] > thresholds[ch]) == (hi_values[ch] == 0)) { - white_result = false; - break; - } - } - if (white_result) - CLEAR_DATA_BIT(pixline, x); - else - SET_DATA_BIT(pixline, x); - } - srcdata += bytes_per_line; - } -} - -static double thresholdRectToPixMicroBench(GPUEnv *env, TessScoreEvaluationInputData input, - ds_device_type type) { - double time; -# if ON_WINDOWS - LARGE_INTEGER freq, time_funct_start, time_funct_end; - QueryPerformanceFrequency(&freq); -# elif ON_APPLE - mach_timebase_info_data_t info = {0, 0}; - mach_timebase_info(&info); - long long start, stop; -# else - timespec time_funct_start, time_funct_end; -# endif - - // input data - unsigned char pixelHi = (unsigned char)255; - int thresholds[4] = {pixelHi, pixelHi, pixelHi, pixelHi}; - - // Pix* pix = pixCreate(width, height, 1); - int top = 0; - int left = 0; - int bytes_per_line = input.width * input.numChannels; - - // function call - if (type == DS_DEVICE_OPENCL_DEVICE) { -# if ON_WINDOWS - QueryPerformanceCounter(&time_funct_start); -# elif ON_APPLE - start = mach_absolute_time(); -# else - clock_gettime(CLOCK_MONOTONIC, &time_funct_start); -# endif - - OpenclDevice::gpuEnv = *env; - int hi_values[4]; - int retVal = OpenclDevice::ThresholdRectToPixOCL( - input.imageData, input.numChannels, bytes_per_line, thresholds, hi_values, &input.pix, - input.height, input.width, top, left); - -# if ON_WINDOWS - QueryPerformanceCounter(&time_funct_end); - time = (time_funct_end.QuadPart - time_funct_start.QuadPart) / (double)(freq.QuadPart); -# elif ON_APPLE - stop = mach_absolute_time(); - if (retVal == 0) { - time = ((stop - start) * (double)info.numer / info.denom) / 1.0E9; - } else { - time = FLT_MAX; - } - -# else - clock_gettime(CLOCK_MONOTONIC, &time_funct_end); - time = (time_funct_end.tv_sec - time_funct_start.tv_sec) * 1.0 + - (time_funct_end.tv_nsec - time_funct_start.tv_nsec) / 1000000000.0; -# endif - } else { - tesseract::ImageThresholder thresholder; - thresholder.SetImage(input.pix); -# if ON_WINDOWS - QueryPerformanceCounter(&time_funct_start); -# elif ON_APPLE - start = mach_absolute_time(); -# else - clock_gettime(CLOCK_MONOTONIC, &time_funct_start); -# endif - int hi_values[4] = {}; - ThresholdRectToPix_Native(input.imageData, input.numChannels, bytes_per_line, thresholds, - hi_values, &input.pix); - -# if ON_WINDOWS - QueryPerformanceCounter(&time_funct_end); - time = (time_funct_end.QuadPart - time_funct_start.QuadPart) / (double)(freq.QuadPart); -# elif ON_APPLE - stop = mach_absolute_time(); - time = ((stop - start) * (double)info.numer / info.denom) / 1.0E9; -# else - clock_gettime(CLOCK_MONOTONIC, &time_funct_end); - time = (time_funct_end.tv_sec - time_funct_start.tv_sec) * 1.0 + - (time_funct_end.tv_nsec - time_funct_start.tv_nsec) / 1000000000.0; -# endif - } - - return time; -} - -static double getLineMasksMorphMicroBench(GPUEnv *env, TessScoreEvaluationInputData input, - ds_device_type type) { - double time = 0; -# if ON_WINDOWS - LARGE_INTEGER freq, time_funct_start, time_funct_end; - QueryPerformanceFrequency(&freq); -# elif ON_APPLE - mach_timebase_info_data_t info = {0, 0}; - mach_timebase_info(&info); - long long start, stop; -# else - timespec time_funct_start, time_funct_end; -# endif - - // input data - int resolution = 300; - int wpl = pixGetWpl(input.pix); - int kThinLineFraction = 20; // tess constant - int kMinLineLengthFraction = 4; // tess constant - int max_line_width = resolution / kThinLineFraction; - int min_line_length = resolution / kMinLineLengthFraction; - int closing_brick = max_line_width / 3; - - // function call - if (type == DS_DEVICE_OPENCL_DEVICE) { -# if ON_WINDOWS - QueryPerformanceCounter(&time_funct_start); -# elif ON_APPLE - start = mach_absolute_time(); -# else - clock_gettime(CLOCK_MONOTONIC, &time_funct_start); -# endif - OpenclDevice::gpuEnv = *env; - OpenclDevice::initMorphCLAllocations(wpl, input.height, input.pix); - Image 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::releaseMorphCLBuffers(); - -# if ON_WINDOWS - QueryPerformanceCounter(&time_funct_end); - time = (time_funct_end.QuadPart - time_funct_start.QuadPart) / (double)(freq.QuadPart); -# elif ON_APPLE - stop = mach_absolute_time(); - time = ((stop - start) * (double)info.numer / info.denom) / 1.0E9; -# else - clock_gettime(CLOCK_MONOTONIC, &time_funct_end); - time = (time_funct_end.tv_sec - time_funct_start.tv_sec) * 1.0 + - (time_funct_end.tv_nsec - time_funct_start.tv_nsec) / 1000000000.0; -# endif - } else { -# if ON_WINDOWS - QueryPerformanceCounter(&time_funct_start); -# elif ON_APPLE - start = mach_absolute_time(); -# else - clock_gettime(CLOCK_MONOTONIC, &time_funct_start); -# endif - - // native serial code - Image src_pix = input.pix; - Image pix_closed = pixCloseBrick(nullptr, src_pix, closing_brick, closing_brick); - Image pix_solid = pixOpenBrick(nullptr, pix_closed, max_line_width, max_line_width); - Image pix_hollow = pixSubtract(nullptr, pix_closed, pix_solid); - pix_solid.destroy(); - Image pix_vline = pixOpenBrick(nullptr, pix_hollow, 1, min_line_length); - Image pix_hline = pixOpenBrick(nullptr, pix_hollow, min_line_length, 1); - pix_hline.destroy(); - pix_vline.destroy(); - pix_hollow.destroy(); - -# if ON_WINDOWS - QueryPerformanceCounter(&time_funct_end); - time = (time_funct_end.QuadPart - time_funct_start.QuadPart) / (double)(freq.QuadPart); -# elif ON_APPLE - stop = mach_absolute_time(); - time = ((stop - start) * (double)info.numer / info.denom) / 1.0E9; -# else - clock_gettime(CLOCK_MONOTONIC, &time_funct_end); - time = (time_funct_end.tv_sec - time_funct_start.tv_sec) * 1.0 + - (time_funct_end.tv_nsec - time_funct_start.tv_nsec) / 1000000000.0; -# endif - } - - return time; -} - -/****************************************************************************** - * Device Selection - *****************************************************************************/ - -// encode score object as byte string -static ds_status serializeScore(ds_device *device, uint8_t **serializedScore, - unsigned int *serializedScoreSize) { - *serializedScoreSize = sizeof(TessDeviceScore); - *serializedScore = new uint8_t[*serializedScoreSize]; - memcpy(*serializedScore, device->score, *serializedScoreSize); - return DS_SUCCESS; -} - -// parses byte string and stores in score object -static ds_status deserializeScore(ds_device *device, const uint8_t *serializedScore, - unsigned int serializedScoreSize) { - // check that serializedScoreSize == sizeof(TessDeviceScore); - device->score = new TessDeviceScore; - memcpy(device->score, serializedScore, serializedScoreSize); - return DS_SUCCESS; -} - -static ds_status releaseScore(TessDeviceScore *score) { - delete score; - return DS_SUCCESS; -} - -// evaluate devices -static ds_status evaluateScoreForDevice(ds_device *device, void *inputData) { - // overwrite statuc gpuEnv w/ current device - // so native opencl calls can be used; they use static gpuEnv - tprintf("\n[DS] Device: \"%s\" (%s) evaluation...\n", device->oclDeviceName, - device->type == DS_DEVICE_OPENCL_DEVICE ? "OpenCL" : "Native"); - GPUEnv *env = nullptr; - if (device->type == DS_DEVICE_OPENCL_DEVICE) { - env = &OpenclDevice::gpuEnv; - memset(env, 0, sizeof(*env)); - // tprintf("[DS] populating tmp GPUEnv from device\n"); - populateGPUEnvFromDevice(env, device->oclDeviceID); - env->mnFileCount = 0; // argc; - env->mnKernelCount = 0UL; - // tprintf("[DS] compiling kernels for tmp GPUEnv\n"); - OpenclDevice::CompileKernelFile(env, ""); - } - - TessScoreEvaluationInputData *input = static_cast(inputData); - - // pixReadTiff - double composeRGBPixelTime = composeRGBPixelMicroBench(env, *input, device->type); - - // HistogramRect - double histogramRectTime = histogramRectMicroBench(env, *input, device->type); - - // ThresholdRectToPix - double thresholdRectToPixTime = thresholdRectToPixMicroBench(env, *input, device->type); - - // getLineMasks - double getLineMasksMorphTime = getLineMasksMorphMicroBench(env, *input, device->type); - - // weigh times (% of cpu time) - // these weights should be the % execution time that the native cpu code took - float composeRGBPixelWeight = 1.2f; - float histogramRectWeight = 2.4f; - float thresholdRectToPixWeight = 4.5f; - float getLineMasksMorphWeight = 5.0f; - - float weightedTime = composeRGBPixelWeight * composeRGBPixelTime + - histogramRectWeight * histogramRectTime + - thresholdRectToPixWeight * thresholdRectToPixTime + - getLineMasksMorphWeight * getLineMasksMorphTime; - device->score = new TessDeviceScore; - device->score->time = weightedTime; - - tprintf("[DS] Device: \"%s\" (%s) evaluated\n", device->oclDeviceName, - device->type == DS_DEVICE_OPENCL_DEVICE ? "OpenCL" : "Native"); - tprintf("[DS]%25s: %f (w=%.1f)\n", "composeRGBPixel", composeRGBPixelTime, composeRGBPixelWeight); - tprintf("[DS]%25s: %f (w=%.1f)\n", "HistogramRect", histogramRectTime, histogramRectWeight); - tprintf("[DS]%25s: %f (w=%.1f)\n", "ThresholdRectToPix", thresholdRectToPixTime, - thresholdRectToPixWeight); - tprintf("[DS]%25s: %f (w=%.1f)\n", "getLineMasksMorph", getLineMasksMorphTime, - getLineMasksMorphWeight); - tprintf("[DS]%25s: %f\n", "Score", device->score->time); - return DS_SUCCESS; -} - -// initial call to select device -ds_device OpenclDevice::getDeviceSelection() { - if (!deviceIsSelected) { - // check if opencl is available at runtime - if (1 == LoadOpencl()) { - // opencl is available - // setup devices - ds_status status; - ds_profile *profile; - status = initDSProfile(&profile, "v0.1"); - // try reading scores from file - const char *fileName = "tesseract_opencl_profile_devices.dat"; - status = readProfileFromFile(profile, deserializeScore, fileName); - if (status != DS_SUCCESS) { - // need to run evaluation - tprintf("[DS] Profile file not available (%s); performing profiling.\n", fileName); - - // create input data - TessScoreEvaluationInputData input; - populateTessScoreEvaluationInputData(&input); - // perform evaluations - unsigned int numUpdates; - status = - profileDevices(profile, DS_EVALUATE_ALL, evaluateScoreForDevice, &input, &numUpdates); - // write scores to file - if (status == DS_SUCCESS) { - status = writeProfileToFile(profile, serializeScore, fileName); - if (status == DS_SUCCESS) { - tprintf("[DS] Scores written to file (%s).\n", fileName); - } else { - tprintf( - "[DS] Error saving scores to file (%s); scores not written to " - "file.\n", - fileName); - } - } else { - tprintf( - "[DS] Unable to evaluate performance; scores not written to " - "file.\n"); - } - } else { - tprintf("[DS] Profile read from file (%s).\n", fileName); - } - - // we now have device scores either from file or evaluation - // select fastest using custom Tesseract selection algorithm - float bestTime = FLT_MAX; // begin search with worst possible time - int bestDeviceIdx = -1; - for (unsigned d = 0; d < profile->numDevices; d++) { - ds_device device = profile->devices[d]; - if (device.score == nullptr) - continue; - TessDeviceScore score = *device.score; - - float time = score.time; - tprintf("[DS] Device[%u] %i:%s score is %f\n", d + 1, device.type, device.oclDeviceName, - time); - if (time < bestTime) { - bestTime = time; - bestDeviceIdx = d; - } - } - if (bestDeviceIdx >= 0) { - tprintf( - "[DS] Selected Device[%i]: \"%s\" (%s)\n", bestDeviceIdx + 1, - profile->devices[bestDeviceIdx].oclDeviceName, - profile->devices[bestDeviceIdx].type == DS_DEVICE_OPENCL_DEVICE ? "OpenCL" : "Native"); - } - // cleanup - // TODO: call destructor for profile object? - - bool overridden = false; - char *overrideDeviceStr = getenv("TESSERACT_OPENCL_DEVICE"); - if (overrideDeviceStr != nullptr) { - int overrideDeviceIdx = atoi(overrideDeviceStr); - if (overrideDeviceIdx > 0 && overrideDeviceIdx <= profile->numDevices) { - tprintf( - "[DS] Overriding Device Selection (TESSERACT_OPENCL_DEVICE=%s, " - "%i)\n", - overrideDeviceStr, overrideDeviceIdx); - bestDeviceIdx = overrideDeviceIdx - 1; - overridden = true; - } else { - tprintf( - "[DS] Ignoring invalid TESSERACT_OPENCL_DEVICE=%s ([1,%i] are " - "valid devices).\n", - overrideDeviceStr, profile->numDevices); - } - } - - if (overridden) { - tprintf( - "[DS] Overridden Device[%i]: \"%s\" (%s)\n", bestDeviceIdx + 1, - profile->devices[bestDeviceIdx].oclDeviceName, - profile->devices[bestDeviceIdx].type == DS_DEVICE_OPENCL_DEVICE ? "OpenCL" : "Native"); - } - selectedDevice = profile->devices[bestDeviceIdx]; - // cleanup - releaseDSProfile(profile, releaseScore); - } else { - // opencl isn't available at runtime, select native cpu device - tprintf("[DS] OpenCL runtime not available.\n"); - selectedDevice.type = DS_DEVICE_NATIVE_CPU; - selectedDevice.oclDeviceName = "(null)"; - selectedDevice.score = nullptr; - selectedDevice.oclDeviceID = nullptr; - selectedDevice.oclDriverVersion = nullptr; - } - deviceIsSelected = true; - } - return selectedDevice; -} - -bool OpenclDevice::selectedDeviceIsOpenCL() { - ds_device device = getDeviceSelection(); - return (device.type == DS_DEVICE_OPENCL_DEVICE); -} - -} // namespace - -#endif diff --git a/src/opencl/openclwrapper.h b/src/opencl/openclwrapper.h deleted file mode 100644 index fcf535ad..00000000 --- a/src/opencl/openclwrapper.h +++ /dev/null @@ -1,179 +0,0 @@ -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// http://www.apache.org/licenses/LICENSE-2.0 -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. - -#ifndef TESSERACT_OPENCL_OPENCLWRAPPER_H_ -#define TESSERACT_OPENCL_OPENCLWRAPPER_H_ - -#include -#include -#include "pix.h" -#include "tprintf.h" - -// including CL/cl.h doesn't occur until USE_OPENCL defined below - -/************************************************************************** - * enable/disable use of OpenCL - **************************************************************************/ - -#ifdef USE_OPENCL - -# ifdef __APPLE__ -# include -# else -# include -# endif - -namespace tesseract { - -class Image; -struct TessDeviceScore; - -// device type -enum ds_device_type { DS_DEVICE_NATIVE_CPU = 0, DS_DEVICE_OPENCL_DEVICE }; - -struct ds_device { - ds_device_type type; - cl_device_id oclDeviceID; - char *oclDeviceName; - char *oclDriverVersion; - // a pointer to the score data, the content/format is application defined. - TessDeviceScore *score; -}; - -# ifndef strcasecmp -# define strcasecmp strcmp -# endif - -# define MAX_KERNEL_STRING_LEN 64 -# define MAX_CLFILE_NUM 50 -# define MAX_CLKERNEL_NUM 200 -# define MAX_KERNEL_NAME_LEN 64 -# define CL_QUEUE_THREAD_HANDLE_AMD 0x403E -# define GROUPSIZE_X 16 -# define GROUPSIZE_Y 16 -# define GROUPSIZE_HMORX 256 -# define GROUPSIZE_HMORY 1 - -struct KernelEnv { - cl_context mpkContext; - cl_command_queue mpkCmdQueue; - cl_program mpkProgram; - cl_kernel mpkKernel; - char mckKernelName[150]; -}; - -struct OpenCLEnv { - cl_platform_id mpOclPlatformID; - cl_context mpOclContext; - cl_device_id mpOclDevsID; - cl_command_queue mpOclCmdQueue; -}; -typedef int (*cl_kernel_function)(void **userdata, KernelEnv *kenv); - -# define CHECK_OPENCL(status, name) \ - if (status != CL_SUCCESS) { \ - tprintf("OpenCL error code is %d at when %s .\n", status, name); \ - } - -struct GPUEnv { - // share vb in all modules in hb library - cl_platform_id mpPlatformID; - cl_device_type mDevType; - cl_context mpContext; - cl_device_id *mpArryDevsID; - cl_device_id mpDevID; - cl_command_queue mpCmdQueue; - cl_kernel mpArryKernels[MAX_CLFILE_NUM]; - cl_program mpArryPrograms[MAX_CLFILE_NUM]; // one program object maps one - // kernel source file - char mArryKnelSrcFile[MAX_CLFILE_NUM][256], // the max len of kernel file name is 256 - mArrykernelNames[MAX_CLKERNEL_NUM][MAX_KERNEL_STRING_LEN + 1]; - cl_kernel_function mpArryKnelFuncs[MAX_CLKERNEL_NUM]; - int mnKernelCount, mnFileCount, // only one kernel file - mnIsUserCreated; // 1: created , 0:no create and needed to create by - // opencl wrapper - int mnKhrFp64Flag; - int mnAmdFp64Flag; -}; - -class OpenclDevice { -public: - static GPUEnv gpuEnv; - static int isInited; - OpenclDevice(); - ~OpenclDevice(); - static int InitEnv(); // load dll, call InitOpenclRunEnv(0) - static int InitOpenclRunEnv(int argc); // RegistOpenclKernel, double flags, compile kernels - static int InitOpenclRunEnv_DeviceSelection( - int argc); // RegistOpenclKernel, double flags, compile kernels - static int RegistOpenclKernel(); - static int ReleaseOpenclRunEnv(); - static int ReleaseOpenclEnv(GPUEnv *gpuInfo); - static int CompileKernelFile(GPUEnv *gpuInfo, const char *buildOption); - static int CachedOfKernerPrg(const GPUEnv *gpuEnvCached, const char *clFileName); - static int GeneratBinFromKernelSource(cl_program program, const char *clFileName); - static int WriteBinaryToFile(const char *fileName, const char *birary, size_t numBytes); - static int BinaryGenerated(const char *clFileName, FILE **fhandle); - // static int CompileKernelFile( const char *filename, GPUEnv *gpuInfo, const - // char *buildOption ); - static l_uint32 *pixReadFromTiffKernel(l_uint32 *tiffdata, l_int32 w, l_int32 h, l_int32 wpl, - l_uint32 *line); - static int composeRGBPixelCl(int *tiffdata, int *line, int h, int w); - - /* OpenCL implementations of Morphological operations*/ - - // Initialization of OCL buffers used in Morph operations - static int initMorphCLAllocations(l_int32 wpl, l_int32 h, Image pixs); - static void releaseMorphCLBuffers(); - - static void pixGetLinesCL(Image pixd, Image pixs, Image *pix_vline, Image *pix_hline, Image *pixClosed, - bool getpixClosed, l_int32 close_hsize, l_int32 close_vsize, - l_int32 open_hsize, l_int32 open_vsize, l_int32 line_hsize, - l_int32 line_vsize); - - // int InitOpenclAttr( OpenCLEnv * env ); - // int ReleaseKernel( KernelEnv * env ); - static int SetKernelEnv(KernelEnv *envInfo); - // int CreateKernel( char * kernelname, KernelEnv * env ); - // int RunKernel( const char *kernelName, void **userdata ); - // int ConvertToString( const char *filename, char **source ); - // int CheckKernelName( KernelEnv *envInfo, const char *kernelName ); - // int RegisterKernelWrapper( const char *kernelName, cl_kernel_function - // function ); int RunKernelWrapper( cl_kernel_function function, const char * - // kernelName, void **usrdata ); int GetKernelEnvAndFunc( const char - // *kernelName, KernelEnv *env, cl_kernel_function *function ); - - static int LoadOpencl(); -# ifdef WIN32 - // static int OpenclInite(); - static void FreeOpenclDll(); -# endif - - inline static int AddKernelConfig(int kCount, const char *kName); - - /* for binarization */ - static int HistogramRectOCL(void *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, Image *pix, - int rect_height, int rect_width, int rect_top, int rect_left); - - static ds_device getDeviceSelection(); - static ds_device selectedDevice; - static bool deviceIsSelected; - static bool selectedDeviceIsOpenCL(); -}; - -} - -#endif // USE_OPENCL -#endif // TESSERACT_OPENCL_OPENCLWRAPPER_H_ diff --git a/src/tesseract.cpp b/src/tesseract.cpp index 48081556..cf19f668 100644 --- a/src/tesseract.cpp +++ b/src/tesseract.cpp @@ -34,9 +34,6 @@ #include #include #include "dict.h" -#if defined(USE_OPENCL) -# include "openclwrapper.h" // for OpenclDevice -#endif #include #include "simddetect.h" #include "tesseractclass.h" // for AnyTessLang @@ -112,34 +109,6 @@ static void PrintVersionInfo() { printf(" %s\n", versionStrP); lept_free(versionStrP); -#ifdef USE_OPENCL - cl_platform_id platform[4]; - cl_uint num_platforms; - - printf(" OpenCL info:\n"); - if (clGetPlatformIDs(4, platform, &num_platforms) == CL_SUCCESS) { - printf(" Found %u platform(s).\n", num_platforms); - for (unsigned n = 0; n < num_platforms; n++) { - char info[256]; - if (clGetPlatformInfo(platform[n], CL_PLATFORM_NAME, 256, info, 0) == CL_SUCCESS) { - printf(" Platform %u name: %s.\n", n + 1, info); - } - if (clGetPlatformInfo(platform[n], CL_PLATFORM_VERSION, 256, info, 0) == CL_SUCCESS) { - printf(" Version: %s.\n", info); - } - cl_device_id devices[2]; - cl_uint num_devices; - if (clGetDeviceIDs(platform[n], CL_DEVICE_TYPE_ALL, 2, devices, &num_devices) == CL_SUCCESS) { - printf(" Found %u device(s).\n", num_devices); - for (unsigned i = 0; i < num_devices; ++i) { - if (clGetDeviceInfo(devices[i], CL_DEVICE_NAME, 256, info, 0) == CL_SUCCESS) { - printf(" Device %u name: %s.\n", i + 1, info); - } - } - } - } - } -#endif #if defined(HAVE_NEON) || defined(__aarch64__) if (tesseract::SIMDDetect::IsNEONAvailable()) printf(" Found NEON\n"); diff --git a/src/textord/linefind.cpp b/src/textord/linefind.cpp index 886e1989..19de5835 100644 --- a/src/textord/linefind.cpp +++ b/src/textord/linefind.cpp @@ -27,9 +27,6 @@ #include "edgblob.h" #include "linefind.h" #include "tabvector.h" -#if defined(USE_OPENCL) -# include "openclwrapper.h" // for OpenclDevice -#endif #include @@ -469,48 +466,33 @@ static void GetLineMasks(int resolution, Image src_pix, Image *pix_vline, Image } int closing_brick = max_line_width / 3; -// only use opencl if compiled w/ OpenCL and selected device is opencl -#ifdef USE_OPENCL - if (OpenclDevice::selectedDeviceIsOpenCL()) { - // OpenCL pixGetLines Operation - int clStatus = - OpenclDevice::initMorphCLAllocations(pixGetWpl(src_pix), pixGetHeight(src_pix), src_pix); - bool getpixclosed = pix_music_mask != nullptr; - OpenclDevice::pixGetLinesCL(nullptr, src_pix, pix_vline, pix_hline, &pix_closed, getpixclosed, - closing_brick, closing_brick, max_line_width, max_line_width, - min_line_length, min_line_length); - } else { -#endif - // Close up small holes, making it less likely that false alarms are found - // in thickened text (as it will become more solid) and also smoothing over - // some line breaks and nicks in the edges of the lines. - pix_closed = pixCloseBrick(nullptr, src_pix, closing_brick, closing_brick); - if (pixa_display != nullptr) { - pixaAddPix(pixa_display, pix_closed, L_CLONE); - } - // Open up with a big box to detect solid areas, which can then be - // subtracted. This is very generous and will leave in even quite wide - // lines. - Image pix_solid = pixOpenBrick(nullptr, pix_closed, max_line_width, max_line_width); - if (pixa_display != nullptr) { - pixaAddPix(pixa_display, pix_solid, L_CLONE); - } - pix_hollow = pixSubtract(nullptr, pix_closed, pix_solid); - - pix_solid.destroy(); - - // Now open up in both directions independently to find lines of at least - // 1 inch/kMinLineLengthFraction in length. - if (pixa_display != nullptr) { - pixaAddPix(pixa_display, pix_hollow, L_CLONE); - } - *pix_vline = pixOpenBrick(nullptr, pix_hollow, 1, min_line_length); - *pix_hline = pixOpenBrick(nullptr, pix_hollow, min_line_length, 1); - - pix_hollow.destroy(); -#ifdef USE_OPENCL + // Close up small holes, making it less likely that false alarms are found + // in thickened text (as it will become more solid) and also smoothing over + // some line breaks and nicks in the edges of the lines. + pix_closed = pixCloseBrick(nullptr, src_pix, closing_brick, closing_brick); + if (pixa_display != nullptr) { + pixaAddPix(pixa_display, pix_closed, L_CLONE); } -#endif + // Open up with a big box to detect solid areas, which can then be + // subtracted. This is very generous and will leave in even quite wide + // lines. + Image pix_solid = pixOpenBrick(nullptr, pix_closed, max_line_width, max_line_width); + if (pixa_display != nullptr) { + pixaAddPix(pixa_display, pix_solid, L_CLONE); + } + pix_hollow = pixSubtract(nullptr, pix_closed, pix_solid); + + pix_solid.destroy(); + + // Now open up in both directions independently to find lines of at least + // 1 inch/kMinLineLengthFraction in length. + if (pixa_display != nullptr) { + pixaAddPix(pixa_display, pix_hollow, L_CLONE); + } + *pix_vline = pixOpenBrick(nullptr, pix_hollow, 1, min_line_length); + *pix_hline = pixOpenBrick(nullptr, pix_hollow, min_line_length, 1); + + pix_hollow.destroy(); // Lines are sufficiently rare, that it is worth checking for a zero image. bool v_empty = pix_vline->isZero(); diff --git a/sw.cpp b/sw.cpp index 22cba4a6..dfa26a27 100644 --- a/sw.cpp +++ b/sw.cpp @@ -21,7 +21,6 @@ void build(Solution &s) libtesseract.Public += "include"_idir; libtesseract.Protected += - "src/opencl"_id, "src/ccmain"_id, "src/api"_id, "src/dict"_id, diff --git a/tesseract.pc.in b/tesseract.pc.in index f907f07c..81b1531f 100644 --- a/tesseract.pc.in +++ b/tesseract.pc.in @@ -12,5 +12,5 @@ URL: https://github.com/tesseract-ocr/tesseract Version: @VERSION@ Requires.private: lept Libs: -L${libdir} -ltesseract @libarchive_LIBS@ @libcurl_LIBS@ @TENSORFLOW_LIBS@ -Libs.private: -lpthread @OPENCL_LDFLAGS@ +Libs.private: -lpthread Cflags: -I${includedir}