Merge remote-tracking branch 'origin/2.4' into merge-2.4

Conflicts:
	CMakeLists.txt
	modules/calib3d/src/calibration.cpp
	modules/ocl/src/cl_programcache.cpp
	modules/ocl/src/filtering.cpp
	modules/ocl/src/imgproc.cpp
	samples/ocl/adaptive_bilateral_filter.cpp
	samples/ocl/bgfg_segm.cpp
	samples/ocl/clahe.cpp
	samples/ocl/facedetect.cpp
	samples/ocl/pyrlk_optical_flow.cpp
	samples/ocl/squares.cpp
	samples/ocl/surf_matcher.cpp
	samples/ocl/tvl1_optical_flow.cpp
pull/1704/head
Roman Donchenko 11 years ago
commit 78be4f66f7
  1. 5
      CMakeLists.txt
  2. 14
      cmake/OpenCVFindLibsPerf.cmake
  3. 19
      cmake/OpenCVGenConfig.cmake
  4. 13
      cmake/OpenCVGenPkgconfig.cmake
  5. 5
      cmake/OpenCVModule.cmake
  6. 39
      cmake/OpenCVUtils.cmake
  7. 7
      cmake/templates/OpenCVConfig.cmake.in
  8. 9
      doc/tutorials/ios/image_manipulation/image_manipulation.rst
  9. 4
      modules/calib3d/src/calibration.cpp
  10. 7
      modules/java/CMakeLists.txt
  11. 2
      modules/nonfree/perf/perf_surf.ocl.cpp
  12. 2
      modules/nonfree/src/opencl/surf.cl
  13. 2
      modules/nonfree/src/surf.ocl.cpp
  14. 2
      modules/nonfree/test/test_surf.ocl.cpp
  15. 2
      modules/ocl/include/opencv2/ocl/matrix_operations.hpp
  16. 2
      modules/ocl/perf/main.cpp
  17. 2
      modules/ocl/perf/perf_arithm.cpp
  18. 2
      modules/ocl/perf/perf_bgfg.cpp
  19. 2
      modules/ocl/perf/perf_blend.cpp
  20. 2
      modules/ocl/perf/perf_brute_force_matcher.cpp
  21. 2
      modules/ocl/perf/perf_calib3d.cpp
  22. 2
      modules/ocl/perf/perf_canny.cpp
  23. 2
      modules/ocl/perf/perf_color.cpp
  24. 2
      modules/ocl/perf/perf_fft.cpp
  25. 2
      modules/ocl/perf/perf_filters.cpp
  26. 2
      modules/ocl/perf/perf_gemm.cpp
  27. 2
      modules/ocl/perf/perf_gftt.cpp
  28. 2
      modules/ocl/perf/perf_haar.cpp
  29. 2
      modules/ocl/perf/perf_hog.cpp
  30. 2
      modules/ocl/perf/perf_imgproc.cpp
  31. 2
      modules/ocl/perf/perf_kalman.cpp
  32. 2
      modules/ocl/perf/perf_match_template.cpp
  33. 2
      modules/ocl/perf/perf_matrix_operation.cpp
  34. 2
      modules/ocl/perf/perf_ml.cpp
  35. 2
      modules/ocl/perf/perf_moments.cpp
  36. 2
      modules/ocl/perf/perf_norm.cpp
  37. 2
      modules/ocl/perf/perf_opticalflow.cpp
  38. 2
      modules/ocl/perf/perf_precomp.hpp
  39. 2
      modules/ocl/perf/perf_pyramid.cpp
  40. 2
      modules/ocl/perf/perf_split_merge.cpp
  41. 2
      modules/ocl/src/arithm.cpp
  42. 2
      modules/ocl/src/bgfg_mog.cpp
  43. 2
      modules/ocl/src/blend.cpp
  44. 2
      modules/ocl/src/brute_force_matcher.cpp
  45. 2
      modules/ocl/src/build_warps.cpp
  46. 2
      modules/ocl/src/canny.cpp
  47. 5
      modules/ocl/src/cl_context.cpp
  48. 2
      modules/ocl/src/cl_operations.cpp
  49. 43
      modules/ocl/src/cl_programcache.cpp
  50. 3
      modules/ocl/src/cl_programcache.hpp
  51. 2
      modules/ocl/src/color.cpp
  52. 2
      modules/ocl/src/columnsum.cpp
  53. 2
      modules/ocl/src/error.cpp
  54. 5
      modules/ocl/src/fft.cpp
  55. 70
      modules/ocl/src/filtering.cpp
  56. 2
      modules/ocl/src/gemm.cpp
  57. 2
      modules/ocl/src/gftt.cpp
  58. 2
      modules/ocl/src/haar.cpp
  59. 2
      modules/ocl/src/hog.cpp
  60. 43
      modules/ocl/src/imgproc.cpp
  61. 2
      modules/ocl/src/interpolate_frames.cpp
  62. 2
      modules/ocl/src/kalman.cpp
  63. 2
      modules/ocl/src/kmeans.cpp
  64. 2
      modules/ocl/src/knearest.cpp
  65. 2
      modules/ocl/src/match_template.cpp
  66. 2
      modules/ocl/src/matrix_operations.cpp
  67. 2
      modules/ocl/src/mcwutil.cpp
  68. 2
      modules/ocl/src/moments.cpp
  69. 2
      modules/ocl/src/mssegmentation.cpp
  70. 2
      modules/ocl/src/opencl/arithm_LUT.cl
  71. 2
      modules/ocl/src/opencl/arithm_absdiff_nonsaturate.cl
  72. 2
      modules/ocl/src/opencl/arithm_add.cl
  73. 2
      modules/ocl/src/opencl/arithm_addWeighted.cl
  74. 2
      modules/ocl/src/opencl/arithm_add_mask.cl
  75. 2
      modules/ocl/src/opencl/arithm_add_scalar.cl
  76. 2
      modules/ocl/src/opencl/arithm_add_scalar_mask.cl
  77. 2
      modules/ocl/src/opencl/arithm_bitwise_binary.cl
  78. 2
      modules/ocl/src/opencl/arithm_bitwise_binary_mask.cl
  79. 2
      modules/ocl/src/opencl/arithm_bitwise_binary_scalar.cl
  80. 2
      modules/ocl/src/opencl/arithm_bitwise_not.cl
  81. 2
      modules/ocl/src/opencl/arithm_cartToPolar.cl
  82. 2
      modules/ocl/src/opencl/arithm_compare.cl
  83. 2
      modules/ocl/src/opencl/arithm_exp.cl
  84. 2
      modules/ocl/src/opencl/arithm_flip.cl
  85. 2
      modules/ocl/src/opencl/arithm_flip_rc.cl
  86. 2
      modules/ocl/src/opencl/arithm_log.cl
  87. 2
      modules/ocl/src/opencl/arithm_nonzero.cl
  88. 2
      modules/ocl/src/opencl/arithm_phase.cl
  89. 2
      modules/ocl/src/opencl/arithm_pow.cl
  90. 2
      modules/ocl/src/opencl/arithm_setidentity.cl
  91. 2
      modules/ocl/src/opencl/arithm_sum.cl
  92. 2
      modules/ocl/src/opencl/arithm_transpose.cl
  93. 2
      modules/ocl/src/opencl/bgfg_mog.cl
  94. 2
      modules/ocl/src/opencl/brute_force_match.cl
  95. 2
      modules/ocl/src/opencl/build_warps.cl
  96. 2
      modules/ocl/src/opencl/convertC3C4.cl
  97. 2
      modules/ocl/src/opencl/cvt_color.cl
  98. 43
      modules/ocl/src/opencl/filter_sep_col.cl
  99. 360
      modules/ocl/src/opencl/filter_sep_row.cl
  100. 114
      modules/ocl/src/opencl/filtering_adaptive_bilateral.cl
  101. Some files were not shown because too many files have changed in this diff Show More

@ -32,6 +32,10 @@ else(NOT CMAKE_TOOLCHAIN_FILE)
endif(NOT CMAKE_TOOLCHAIN_FILE)
if(POLICY CMP0022)
cmake_policy(SET CMP0022 OLD)
endif()
# must go before the project command
set(CMAKE_CONFIGURATION_TYPES "Debug;Release" CACHE STRING "Configs" FORCE)
if(DEFINED CMAKE_BUILD_TYPE)
@ -136,6 +140,7 @@ OCV_OPTION(WITH_QT "Build with Qt Backend support" OFF
OCV_OPTION(WITH_WIN32UI "Build with Win32 UI Backend support" ON IF WIN32 )
OCV_OPTION(WITH_QUICKTIME "Use QuickTime for Video I/O insted of QTKit" OFF IF APPLE )
OCV_OPTION(WITH_TBB "Include Intel TBB support" OFF IF (NOT IOS) )
OCV_OPTION(WITH_OPENMP "Include OpenMP support" OFF)
OCV_OPTION(WITH_CSTRIPES "Include C= support" OFF IF WIN32 )
OCV_OPTION(WITH_TIFF "Include TIFF support" ON IF (NOT IOS) )
OCV_OPTION(WITH_UNICAP "Include Unicap support (GPL)" OFF IF (UNIX AND NOT APPLE AND NOT ANDROID) )

@ -86,13 +86,13 @@ else()
endif()
# --- OpenMP ---
if(NOT HAVE_TBB AND NOT HAVE_CSTRIPES)
set(_fname "${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/CMakeTmp/omptest.cpp")
file(WRITE "${_fname}" "#ifndef _OPENMP\n#error\n#endif\nint main() { return 0; }\n")
try_compile(HAVE_OPENMP "${CMAKE_BINARY_DIR}" "${_fname}")
file(REMOVE "${_fname}")
else()
set(HAVE_OPENMP 0)
if(WITH_OPENMP AND NOT HAVE_TBB AND NOT HAVE_CSTRIPES)
find_package(OpenMP)
if(OPENMP_FOUND)
set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${OpenMP_C_FLAGS}")
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${OpenMP_CXX_FLAGS}")
endif()
set(HAVE_OPENMP "${OPENMP_FOUND}")
endif()
# --- GCD ---

@ -74,7 +74,14 @@ if(ANDROID AND NOT BUILD_SHARED_LIBS AND HAVE_TBB)
list(APPEND OpenCV2_INCLUDE_DIRS_CONFIGCMAKE ${TBB_INCLUDE_DIRS})
endif()
export(TARGETS ${OpenCVModules_TARGETS} FILE "${CMAKE_BINARY_DIR}/OpenCVModules.cmake")
set(modules_file_suffix "")
if(ANDROID)
# the REPLACE here is needed, because OpenCVModules_armeabi.cmake includes
# OpenCVModules_armeabi-*.cmake, which would match OpenCVModules_armeabi-v7a*.cmake.
string(REPLACE - _ modules_file_suffix "_${ANDROID_NDK_ABI_NAME}")
endif()
export(TARGETS ${OpenCVModules_TARGETS} FILE "${CMAKE_BINARY_DIR}/OpenCVModules${modules_file_suffix}.cmake")
configure_file("${OpenCV_SOURCE_DIR}/cmake/templates/OpenCVConfig.cmake.in" "${CMAKE_BINARY_DIR}/OpenCVConfig.cmake" IMMEDIATE @ONLY)
#support for version checking when finding opencv. find_package(OpenCV 2.3.1 EXACT) should now work.
@ -94,7 +101,7 @@ endif()
configure_file("${OpenCV_SOURCE_DIR}/cmake/templates/OpenCVConfig.cmake.in" "${CMAKE_BINARY_DIR}/unix-install/OpenCVConfig.cmake" IMMEDIATE @ONLY)
configure_file("${OpenCV_SOURCE_DIR}/cmake/templates/OpenCVConfig-version.cmake.in" "${CMAKE_BINARY_DIR}/unix-install/OpenCVConfig-version.cmake" IMMEDIATE @ONLY)
if(UNIX)
if(UNIX) # ANDROID configuration is created here also
#http://www.vtk.org/Wiki/CMake/Tutorials/Packaging reference
# For a command "find_package(<name> [major[.minor]] [EXACT] [REQUIRED|QUIET])"
# cmake will look in the following dir on unix:
@ -104,11 +111,11 @@ if(UNIX)
if(INSTALL_TO_MANGLED_PATHS)
install(FILES ${CMAKE_BINARY_DIR}/unix-install/OpenCVConfig.cmake DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}-${OPENCV_VERSION}/)
install(FILES ${CMAKE_BINARY_DIR}/unix-install/OpenCVConfig-version.cmake DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}-${OPENCV_VERSION}/)
install(EXPORT OpenCVModules DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}-${OPENCV_VERSION}/)
install(EXPORT OpenCVModules DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}-${OPENCV_VERSION}/ FILE OpenCVModules${modules_file_suffix}.cmake)
else()
install(FILES "${CMAKE_BINARY_DIR}/unix-install/OpenCVConfig.cmake" DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}/)
install(FILES ${CMAKE_BINARY_DIR}/unix-install/OpenCVConfig-version.cmake DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}/)
install(EXPORT OpenCVModules DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}/)
install(EXPORT OpenCVModules DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}/ FILE OpenCVModules${modules_file_suffix}.cmake)
endif()
endif()
@ -128,10 +135,10 @@ if(WIN32)
configure_file("${OpenCV_SOURCE_DIR}/cmake/templates/OpenCVConfig-version.cmake.in" "${CMAKE_BINARY_DIR}/win-install/OpenCVConfig-version.cmake" IMMEDIATE @ONLY)
if(BUILD_SHARED_LIBS)
install(FILES "${CMAKE_BINARY_DIR}/win-install/OpenCVConfig.cmake" DESTINATION "${OpenCV_INSTALL_BINARIES_PREFIX}/lib")
install(EXPORT OpenCVModules DESTINATION "${OpenCV_INSTALL_BINARIES_PREFIX}/lib")
install(EXPORT OpenCVModules DESTINATION "${OpenCV_INSTALL_BINARIES_PREFIX}/lib" FILE OpenCVModules${modules_file_suffix}.cmake)
else()
install(FILES "${CMAKE_BINARY_DIR}/win-install/OpenCVConfig.cmake" DESTINATION "${OpenCV_INSTALL_BINARIES_PREFIX}/staticlib")
install(EXPORT OpenCVModules DESTINATION "${OpenCV_INSTALL_BINARIES_PREFIX}/staticlib")
install(EXPORT OpenCVModules DESTINATION "${OpenCV_INSTALL_BINARIES_PREFIX}/staticlib" FILE OpenCVModules${modules_file_suffix}.cmake)
endif()
install(FILES "${CMAKE_BINARY_DIR}/win-install/OpenCVConfig-version.cmake" DESTINATION "${CMAKE_INSTALL_PREFIX}")
install(FILES "${OpenCV_SOURCE_DIR}/cmake/OpenCVConfig.cmake" DESTINATION "${CMAKE_INSTALL_PREFIX}/")

@ -57,8 +57,17 @@ endforeach()
# add extra dependencies required for OpenCV
set(OpenCV_LIB_COMPONENTS ${OpenCV_LIB_COMPONENTS_})
if(OpenCV_EXTRA_COMPONENTS)
string(REPLACE ";" " " OpenCV_EXTRA_COMPONENTS "${OpenCV_EXTRA_COMPONENTS}")
set(OpenCV_LIB_COMPONENTS "${OpenCV_LIB_COMPONENTS} ${OpenCV_EXTRA_COMPONENTS}")
foreach(extra_component ${OpenCV_EXTRA_COMPONENTS})
if(extra_component MATCHES "^-[lL]" OR extra_component MATCHES "[\\/]")
set(maybe_l_prefix "")
else()
set(maybe_l_prefix "-l")
endif()
set(OpenCV_LIB_COMPONENTS "${OpenCV_LIB_COMPONENTS} ${maybe_l_prefix}${extra_component}")
endforeach()
endif()
#generate the .pc file

@ -538,9 +538,10 @@ macro(ocv_create_module)
if(NOT "${ARGN}" STREQUAL "SKIP_LINK")
target_link_libraries(${the_module} ${OPENCV_MODULE_${the_module}_DEPS})
target_link_libraries(${the_module} LINK_PRIVATE ${OPENCV_MODULE_${the_module}_DEPS_EXT} ${OPENCV_LINKER_LIBS} ${IPP_LIBS} ${ARGN})
target_link_libraries(${the_module} LINK_INTERFACE_LIBRARIES ${OPENCV_MODULE_${the_module}_DEPS})
target_link_libraries(${the_module} ${OPENCV_MODULE_${the_module}_DEPS_EXT} ${OPENCV_LINKER_LIBS} ${IPP_LIBS} ${ARGN})
if (HAVE_CUDA)
target_link_libraries(${the_module} LINK_PRIVATE ${CUDA_LIBRARIES} ${CUDA_npp_LIBRARY})
target_link_libraries(${the_module} ${CUDA_LIBRARIES} ${CUDA_npp_LIBRARY})
endif()
endif()

@ -11,6 +11,18 @@ if(NOT COMMAND find_host_program)
endmacro()
endif()
# assert macro
# Note: it doesn't support lists in arguments
# Usage samples:
# ocv_assert(MyLib_FOUND)
# ocv_assert(DEFINED MyLib_INCLUDE_DIRS)
macro(ocv_assert)
if(NOT (${ARGN}))
string(REPLACE ";" " " __assert_msg "${ARGN}")
message(AUTHOR_WARNING "Assertion failed: ${__assert_msg}")
endif()
endmacro()
macro(ocv_check_environment_variables)
foreach(_var ${ARGN})
if(NOT DEFINED ${_var} AND DEFINED ENV{${_var}})
@ -467,9 +479,10 @@ function(ocv_install_target)
set(isArchive 0)
set(isDst 0)
unset(__dst)
foreach(e ${ARGN})
if(isDst EQUAL 1)
set(DST "${e}")
set(__dst "${e}")
break()
endif()
if(isArchive EQUAL 1 AND e STREQUAL "DESTINATION")
@ -482,18 +495,20 @@ function(ocv_install_target)
endif()
endforeach()
# message(STATUS "Process ${__target} dst=${DST}...")
if(NOT DEFINED DST)
set(DST "OPENCV_LIB_INSTALL_PATH")
endif()
get_target_property(fname ${__target} LOCATION_DEBUG)
string(REPLACE ".lib" ".pdb" fname "${fname}")
install(FILES ${fname} DESTINATION ${DST} CONFIGURATIONS Debug)
# message(STATUS "Process ${__target} dst=${__dst}...")
if(DEFINED __dst)
get_target_property(fname ${__target} LOCATION_DEBUG)
if(fname MATCHES "\\.lib$")
string(REGEX REPLACE "\\.lib$" ".pdb" fname "${fname}")
install(FILES ${fname} DESTINATION ${__dst} CONFIGURATIONS Debug)
endif()
get_target_property(fname ${__target} LOCATION_RELEASE)
string(REPLACE ".lib" ".pdb" fname "${fname}")
install(FILES ${fname} DESTINATION ${DST} CONFIGURATIONS Release)
get_target_property(fname ${__target} LOCATION_RELEASE)
if(fname MATCHES "\\.lib$")
string(REGEX REPLACE "\\.lib$" ".pdb" fname "${fname}")
install(FILES ${fname} DESTINATION ${__dst} CONFIGURATIONS Release)
endif()
endif()
endif()
endif()
endfunction()

@ -37,7 +37,12 @@
#
# ===================================================================================
include(${CMAKE_CURRENT_LIST_DIR}/OpenCVModules.cmake)
set(modules_file_suffix "")
if(ANDROID)
string(REPLACE - _ modules_file_suffix "_${ANDROID_NDK_ABI_NAME}")
endif()
include(${CMAKE_CURRENT_LIST_DIR}/OpenCVModules${modules_file_suffix}.cmake)
# TODO All things below should be reviewed. What is about of moving this code into related modules (special vars/hooks/files)

@ -12,7 +12,8 @@ In this tutorial we will learn how to do basic image processing using OpenCV in
*Introduction*
==============
In *OpenCV* all the image processing operations are done on *Mat*. iOS uses UIImage object to display image. One of the thing is to convert UIImage object to Mat object. Below is the code to convert UIImage to Mat.
In *OpenCV* all the image processing operations are usually carried out on the *Mat* structure. In iOS however, to render an image on screen it have to be an instance of the *UIImage* class. To convert an *OpenCV Mat* to an *UIImage* we use the *Core Graphics* framework available in iOS. Below is the code needed to covert back and forth between Mat's and UIImage's.
.. code-block:: cpp
@ -22,7 +23,7 @@ In *OpenCV* all the image processing operations are done on *Mat*. iOS uses UIIm
CGFloat cols = image.size.width;
CGFloat rows = image.size.height;
cv::Mat cvMat(rows, cols, CV_8UC4); // 8 bits per component, 4 channels
cv::Mat cvMat(rows, cols, CV_8UC4); // 8 bits per component, 4 channels (color channels + alpha)
CGContextRef contextRef = CGBitmapContextCreate(cvMat.data, // Pointer to data
cols, // Width of bitmap
@ -35,7 +36,6 @@ In *OpenCV* all the image processing operations are done on *Mat*. iOS uses UIIm
CGContextDrawImage(contextRef, CGRectMake(0, 0, cols, rows), image.CGImage);
CGContextRelease(contextRef);
CGColorSpaceRelease(colorSpace);
return cvMat;
}
@ -61,12 +61,11 @@ In *OpenCV* all the image processing operations are done on *Mat*. iOS uses UIIm
CGContextDrawImage(contextRef, CGRectMake(0, 0, cols, rows), image.CGImage);
CGContextRelease(contextRef);
CGColorSpaceRelease(colorSpace);
return cvMat;
}
Once we obtain the Mat Object. We can do all our processing on Mat object, similar to cpp. For example if we want to convert image to gray, we can do it via below code.
After the processing we need to convert it back to UIImage. The code below can handle both gray-scale and color image conversions (determined by the number of channels in the *if* statement).
.. code-block:: cpp

@ -1164,8 +1164,8 @@ CV_IMPL void cvInitIntrinsicParams2D( const CvMat* objectPoints,
matA.reset(cvCreateMat( 2*nimages, 2, CV_64F ));
_b.reset(cvCreateMat( 2*nimages, 1, CV_64F ));
a[2] = (imageSize.width - 1)*0.5;
a[5] = (imageSize.height - 1)*0.5;
a[2] = (!imageSize.width) ? 0.5 : (imageSize.width - 1)*0.5;
a[5] = (!imageSize.height) ? 0.5 : (imageSize.height - 1)*0.5;
_allH.reset(cvCreateMat( nimages, 9, CV_64F ));
// extract vanishing points in order to obtain initial value for the focal length

@ -280,7 +280,12 @@ set_target_properties(${the_module} PROPERTIES COMPILE_DEFINITIONS OPENCV_NOSTL)
if(BUILD_FAT_JAVA_LIB)
set(__deps ${OPENCV_MODULE_${the_module}_DEPS} ${OPENCV_MODULES_BUILD})
list(REMOVE_ITEM __deps ${the_module} opencv_ts)
foreach(m ${OPENCV_MODULES_BUILD}) # filterout INTERNAL (like opencv_ts) and BINDINGS (like opencv_python) modules
ocv_assert(DEFINED OPENCV_MODULE_${m}_CLASS)
if(NOT OPENCV_MODULE_${m}_CLASS STREQUAL "PUBLIC")
list(REMOVE_ITEM __deps ${m})
endif()
endforeach()
ocv_list_unique(__deps)
set(__extradeps ${__deps})
ocv_list_filterout(__extradeps "^opencv_")

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -26,7 +26,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -23,7 +23,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -22,7 +22,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -26,7 +26,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -26,7 +26,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -26,7 +26,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -26,7 +26,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -26,7 +26,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -26,7 +26,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -26,7 +26,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -26,7 +26,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -26,7 +26,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -26,7 +26,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -26,7 +26,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -26,7 +26,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -26,7 +26,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -26,7 +26,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -26,7 +26,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -26,7 +26,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -26,7 +26,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -26,7 +26,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -26,7 +26,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -22,7 +22,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -26,7 +26,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -26,7 +26,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -32,7 +32,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -26,7 +26,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -27,7 +27,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
@ -757,6 +757,9 @@ __Module::~__Module()
#if defined(WIN32) && defined(CVAPI_EXPORTS)
extern "C"
BOOL WINAPI DllMain(HINSTANCE /*hInst*/, DWORD fdwReason, LPVOID lpReserved);
extern "C"
BOOL WINAPI DllMain(HINSTANCE /*hInst*/, DWORD fdwReason, LPVOID lpReserved)
{

@ -27,7 +27,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -27,7 +27,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
@ -61,12 +61,16 @@ namespace cv { namespace ocl {
cv::Mutex ProgramCache::mutexFiles;
cv::Mutex ProgramCache::mutexCache;
std::auto_ptr<ProgramCache> _programCache;
ProgramCache* _programCache = NULL;
ProgramCache* ProgramCache::getProgramCache()
{
if (NULL == _programCache.get())
_programCache.reset(new ProgramCache());
return _programCache.get();
if (NULL == _programCache)
{
cv::AutoLock lock(getInitializationMutex());
if (NULL == _programCache)
_programCache = new ProgramCache();
}
return _programCache;
}
ProgramCache::ProgramCache()
@ -78,6 +82,12 @@ ProgramCache::ProgramCache()
ProgramCache::~ProgramCache()
{
releaseProgram();
if (this == _programCache)
{
cv::AutoLock lock(getInitializationMutex());
if (this == _programCache)
_programCache = NULL;
}
}
cl_program ProgramCache::progLookup(const String& srcsign)
@ -420,22 +430,17 @@ struct ProgramFileCache
{
if(status == CL_BUILD_PROGRAM_FAILURE)
{
cl_int logStatus;
char *buildLog = NULL;
size_t buildLogSize = 0;
logStatus = clGetProgramBuildInfo(program,
getClDeviceID(ctx), CL_PROGRAM_BUILD_LOG, buildLogSize,
buildLog, &buildLogSize);
if(logStatus != CL_SUCCESS)
std::cout << "Failed to build the program and get the build info." << std::endl;
buildLog = new char[buildLogSize];
CV_DbgAssert(!!buildLog);
memset(buildLog, 0, buildLogSize);
openCLSafeCall(clGetProgramBuildInfo(program, getClDeviceID(ctx),
CL_PROGRAM_BUILD_LOG, buildLogSize, buildLog, NULL));
std::cout << "\nBUILD LOG: " << options << "\n";
std::cout << buildLog << std::endl;
delete [] buildLog;
CL_PROGRAM_BUILD_LOG, 0, NULL, &buildLogSize));
std::vector<char> buildLog; buildLog.resize(buildLogSize);
memset(&buildLog[0], 0, buildLogSize);
openCLSafeCall(clGetProgramBuildInfo(program, getClDeviceID(ctx),
CL_PROGRAM_BUILD_LOG, buildLogSize, &buildLog[0], NULL));
std::cout << std::endl << "BUILD LOG: "
<< (source->name ? source->name : "dynamic program") << ": "
<< options << "\n";
std::cout << &buildLog[0] << std::endl;
}
openCLVerifyCall(status);
}

@ -24,7 +24,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
@ -52,7 +52,6 @@ class ProgramCache
protected:
ProgramCache();
~ProgramCache();
friend class std::auto_ptr<ProgramCache>;
public:
static ProgramCache *getProgramCache();

@ -26,7 +26,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -24,7 +24,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
@ -126,7 +126,8 @@ void cv::ocl::fft_setup()
{
return;
}
pCache.setupData = new clAmdFftSetupData;
if (pCache.setupData == NULL)
pCache.setupData = new clAmdFftSetupData;
openCLSafeCall(clAmdFftInitSetupData( pCache.setupData ));
pCache.started = true;
}

@ -29,7 +29,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
@ -1059,74 +1059,39 @@ template <> struct index_and_sizeof<float>
template <typename T>
void linearRowFilter_gpu(const oclMat &src, const oclMat &dst, oclMat mat_kernel, int ksize, int anchor, int bordertype)
{
Context *clCxt = src.clCxt;
CV_Assert(bordertype <= BORDER_REFLECT_101);
CV_Assert(ksize == (anchor << 1) + 1);
int channels = src.oclchannels();
size_t localThreads[3] = {16, 16, 1};
String kernelName = "row_filter";
char btype[30];
size_t localThreads[3] = { 16, 16, 1 };
size_t globalThreads[3] = { dst.cols, dst.rows, 1 };
switch (bordertype)
{
case 0:
sprintf(btype, "BORDER_CONSTANT");
break;
case 1:
sprintf(btype, "BORDER_REPLICATE");
break;
case 2:
sprintf(btype, "BORDER_REFLECT");
break;
case 3:
sprintf(btype, "BORDER_WRAP");
break;
case 4:
sprintf(btype, "BORDER_REFLECT_101");
break;
}
char compile_option[128];
sprintf(compile_option, "-D RADIUSX=%d -D LSIZE0=%d -D LSIZE1=%d -D CN=%d -D %s", anchor, (int)localThreads[0], (int)localThreads[1], channels, btype);
size_t globalThreads[3];
globalThreads[1] = (dst.rows + localThreads[1] - 1) / localThreads[1] * localThreads[1];
globalThreads[2] = (1 + localThreads[2] - 1) / localThreads[2] * localThreads[2];
const char * const borderMap[] = { "BORDER_CONSTANT", "BORDER_REPLICATE", "BORDER_REFLECT", "BORDER_WRAP", "BORDER_REFLECT_101" };
std::string buildOptions = format("-D RADIUSX=%d -D LSIZE0=%d -D LSIZE1=%d -D CN=%d -D %s",
anchor, (int)localThreads[0], (int)localThreads[1], channels, borderMap[bordertype]);
if (src.depth() == CV_8U)
{
switch (channels)
{
case 1:
case 3:
globalThreads[0] = ((dst.cols + 4) / 4 + localThreads[0] - 1) / localThreads[0] * localThreads[0];
globalThreads[0] = (dst.cols + 3) >> 2;
break;
case 2:
globalThreads[0] = ((dst.cols + 1) / 2 + localThreads[0] - 1) / localThreads[0] * localThreads[0];
globalThreads[0] = (dst.cols + 1) >> 1;
break;
case 4:
globalThreads[0] = (dst.cols + localThreads[0] - 1) / localThreads[0] * localThreads[0];
globalThreads[0] = dst.cols;
break;
}
}
else
{
globalThreads[0] = (dst.cols + localThreads[0] - 1) / localThreads[0] * localThreads[0];
}
//sanity checks
CV_Assert(clCxt == dst.clCxt);
CV_Assert(src.cols == dst.cols);
CV_Assert(src.oclchannels() == dst.oclchannels());
CV_Assert(ksize == (anchor << 1) + 1);
int src_pix_per_row, dst_pix_per_row;
int src_offset_x, src_offset_y;//, dst_offset_in_pixel;
src_pix_per_row = src.step / src.elemSize();
src_offset_x = (src.offset % src.step) / src.elemSize();
src_offset_y = src.offset / src.step;
dst_pix_per_row = dst.step / dst.elemSize();
//dst_offset_in_pixel = dst.offset / dst.elemSize();
int src_pix_per_row = src.step / src.elemSize();
int src_offset_x = (src.offset % src.step) / src.elemSize();
int src_offset_y = src.offset / src.step;
int dst_pix_per_row = dst.step / dst.elemSize();
int ridusy = (dst.rows - src.rows) >> 1;
std::vector<std::pair<size_t , const void *> > args;
args.push_back(std::make_pair(sizeof(cl_mem), &src.data));
args.push_back(std::make_pair(sizeof(cl_mem), &dst.data));
@ -1141,7 +1106,8 @@ void linearRowFilter_gpu(const oclMat &src, const oclMat &dst, oclMat mat_kernel
args.push_back(std::make_pair(sizeof(cl_int), (void *)&ridusy));
args.push_back(std::make_pair(sizeof(cl_mem), (void *)&mat_kernel.data));
openCLExecuteKernel(clCxt, &filter_sep_row, kernelName, globalThreads, localThreads, args, channels, src.depth(), compile_option);
openCLExecuteKernel(src.clCxt, &filter_sep_row, "row_filter", globalThreads, localThreads,
args, channels, src.depth(), buildOptions.c_str());
}
Ptr<BaseRowFilter_GPU> cv::ocl::getLinearRowFilter_GPU(int srcType, int /*bufType*/, const Mat &rowKernel, int anchor, int bordertype)

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -29,7 +29,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -36,7 +36,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
@ -437,7 +437,7 @@ namespace cv
CV_Assert(top >= 0 && bottom >= 0 && left >= 0 && right >= 0);
if( _src.offset != 0 && (bordertype & BORDER_ISOLATED) == 0 )
if( (_src.wholecols != _src.cols || _src.wholerows != _src.rows) && (bordertype & BORDER_ISOLATED) == 0 )
{
Size wholeSize;
Point ofs;
@ -454,34 +454,25 @@ namespace cv
}
bordertype &= ~cv::BORDER_ISOLATED;
// TODO need to remove this conditions and fix the code
if (bordertype == cv::BORDER_REFLECT || bordertype == cv::BORDER_WRAP)
{
CV_Assert((_src.cols >= left) && (_src.cols >= right) && (_src.rows >= top) && (_src.rows >= bottom));
}
else if (bordertype == cv::BORDER_REFLECT_101)
{
CV_Assert((_src.cols > left) && (_src.cols > right) && (_src.rows > top) && (_src.rows > bottom));
}
dst.create(_src.rows + top + bottom, _src.cols + left + right, _src.type());
int srcStep = _src.step1() / _src.oclchannels(), dstStep = dst.step1() / dst.oclchannels();
int srcStep = _src.step / _src.elemSize(), dstStep = dst.step / dst.elemSize();
int srcOffset = _src.offset / _src.elemSize(), dstOffset = dst.offset / dst.elemSize();
int depth = _src.depth(), ochannels = _src.oclchannels();
int __bordertype[] = {cv::BORDER_CONSTANT, cv::BORDER_REPLICATE, BORDER_REFLECT, BORDER_WRAP, BORDER_REFLECT_101};
const char *borderstr[] = {"BORDER_CONSTANT", "BORDER_REPLICATE", "BORDER_REFLECT", "BORDER_WRAP", "BORDER_REFLECT_101"};
size_t bordertype_index;
int __bordertype[] = { BORDER_CONSTANT, BORDER_REPLICATE, BORDER_REFLECT, BORDER_WRAP, BORDER_REFLECT_101 };
const char *borderstr[] = { "BORDER_CONSTANT", "BORDER_REPLICATE", "BORDER_REFLECT", "BORDER_WRAP", "BORDER_REFLECT_101" };
for(bordertype_index = 0; bordertype_index < sizeof(__bordertype) / sizeof(int); bordertype_index++)
if (__bordertype[bordertype_index] == bordertype)
int bordertype_index = -1;
for (int i = 0, end = sizeof(__bordertype) / sizeof(int); i < end; i++)
if (__bordertype[i] == bordertype)
{
bordertype_index = i;
break;
if (bordertype_index == sizeof(__bordertype) / sizeof(int))
}
if (bordertype_index < 0)
CV_Error(Error::StsBadArg, "Unsupported border type");
String kernelName = "copymakeborder";
size_t localThreads[3] = {16, 16, 1};
size_t localThreads[3] = { 16, 16, 1 };
size_t globalThreads[3] = { dst.cols, dst.rows, 1 };
std::vector< std::pair<size_t, const void *> > args;
@ -504,12 +495,6 @@ namespace cv
typeMap[depth], channelMap[ochannels],
borderstr[bordertype_index]);
if (src.type() == CV_8UC1 && (dst.offset & 3) == 0 && (dst.cols & 3) == 0)
{
kernelName = "copymakeborder_C1_D0";
globalThreads[0] = dst.cols >> 2;
}
int cn = src.channels(), ocn = src.oclchannels();
int bufSize = src.elemSize1() * ocn;
AutoBuffer<uchar> _buf(bufSize);
@ -519,7 +504,7 @@ namespace cv
args.push_back( std::make_pair( bufSize , (void *)buf ));
openCLExecuteKernel(src.clCxt, &imgproc_copymakeboder, kernelName, globalThreads,
openCLExecuteKernel(src.clCxt, &imgproc_copymakeboder, "copymakeborder", globalThreads,
localThreads, args, -1, -1, buildOptions.c_str());
}

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -27,7 +27,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -26,7 +26,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -24,7 +24,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -16,7 +16,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -26,7 +26,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -26,7 +26,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -26,7 +26,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -26,7 +26,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -26,7 +26,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -26,7 +26,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -26,7 +26,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -25,7 +25,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -15,7 +15,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -26,7 +26,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.

@ -47,36 +47,6 @@
#define READ_TIMES_ROW ((2*(RADIUS+LSIZE0)-1)/LSIZE0)
#endif
#ifdef BORDER_CONSTANT
//BORDER_CONSTANT: iiiiii|abcdefgh|iiiiiii
#define ELEM(i,l_edge,r_edge,elem1,elem2) (i)<(l_edge) | (i) >= (r_edge) ? (elem1) : (elem2)
#endif
#ifdef BORDER_REPLICATE
//BORDER_REPLICATE: aaaaaa|abcdefgh|hhhhhhh
#define ADDR_L(i,l_edge,r_edge) (i) < (l_edge) ? (l_edge) : (i)
#define ADDR_R(i,r_edge,addr) (i) >= (r_edge) ? (r_edge)-1 : (addr)
#endif
#ifdef BORDER_REFLECT
//BORDER_REFLECT: fedcba|abcdefgh|hgfedcb
#define ADDR_L(i,l_edge,r_edge) (i) < (l_edge) ? -(i)-1 : (i)
#define ADDR_R(i,r_edge,addr) (i) >= (r_edge) ? -(i)-1+((r_edge)<<1) : (addr)
#endif
#ifdef BORDER_REFLECT_101
//BORDER_REFLECT_101: gfedcb|abcdefgh|gfedcba
#define ADDR_L(i,l_edge,r_edge) (i) < (l_edge) ? -(i) : (i)
#define ADDR_R(i,r_edge,addr) (i) >= (r_edge) ? -(i)-2+((r_edge)<<1) : (addr)
#endif
#ifdef BORDER_WRAP
//BORDER_WRAP: cdefgh|abcdefgh|abcdefg
#define ADDR_L(i,l_edge,r_edge) (i) < (l_edge) ? (i)+(r_edge) : (i)
#define ADDR_R(i,r_edge,addr) (i) >= (r_edge) ? (i)-(r_edge) : (addr)
#endif
/**********************************************************************************
These kernels are written for separable filters such as Sobel, Scharr, GaussianBlur.
Now(6/29/2011) the kernels only support 8U data type and the anchor of the convovle
@ -107,15 +77,16 @@ __kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void col_filter
{
int x = get_global_id(0);
int y = get_global_id(1);
int l_x = get_local_id(0);
int l_y = get_local_id(1);
int start_addr = mad24(y,src_step_in_pixel,x);
int end_addr = mad24(src_whole_rows - 1,src_step_in_pixel,src_whole_cols);
int i;
GENTYPE_SRC sum;
GENTYPE_SRC temp[READ_TIMES_COL];
__local GENTYPE_SRC LDS_DAT[LSIZE1*READ_TIMES_COL][LSIZE0+1];
int start_addr = mad24(y, src_step_in_pixel, x);
int end_addr = mad24(src_whole_rows - 1, src_step_in_pixel, src_whole_cols);
int i;
GENTYPE_SRC sum, temp[READ_TIMES_COL];
__local GENTYPE_SRC LDS_DAT[LSIZE1 * READ_TIMES_COL][LSIZE0 + 1];
//read pixels from src
for(i = 0;i<READ_TIMES_COL;i++)

@ -16,7 +16,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
@ -48,34 +48,43 @@
#define ALIGN (RADIUS)
#endif
#ifdef BORDER_CONSTANT
//BORDER_CONSTANT: iiiiii|abcdefgh|iiiiiii
#define ELEM(i,l_edge,r_edge,elem1,elem2) (i)<(l_edge) | (i) >= (r_edge) ? (elem1) : (elem2)
#endif
#ifdef BORDER_REPLICATE
//BORDER_REPLICATE: aaaaaa|abcdefgh|hhhhhhh
#define ADDR_L(i,l_edge,r_edge,addr) (i) < (l_edge) ? (l_edge) : (addr)
#define ADDR_R(i,r_edge,addr) (i) >= (r_edge) ? (r_edge)-1 : (addr)
#endif
#elif defined BORDER_REPLICATE
#define EXTRAPOLATE(x, maxV) \
{ \
x = max(min(x, maxV - 1), 0); \
}
#elif defined BORDER_WRAP
#define EXTRAPOLATE(x, maxV) \
{ \
if (x < 0) \
x -= ((x - maxV + 1) / maxV) * maxV; \
if (x >= maxV) \
x %= maxV; \
}
#elif defined(BORDER_REFLECT) || defined(BORDER_REFLECT_101)
#define EXTRAPOLATE_(x, maxV, delta) \
{ \
if (maxV == 1) \
x = 0; \
else \
do \
{ \
if ( x < 0 ) \
x = -x - 1 + delta; \
else \
x = maxV - 1 - (x - maxV) - delta; \
} \
while (x >= maxV || x < 0); \
}
#ifdef BORDER_REFLECT
//BORDER_REFLECT: fedcba|abcdefgh|hgfedcb
#define ADDR_L(i,l_edge,r_edge,addr) (i) < (l_edge) ? -(i)-1 : (addr)
#define ADDR_R(i,r_edge,addr) (i) >= (r_edge) ? -(i)-1+((r_edge)<<1) : (addr)
#endif
#ifdef BORDER_REFLECT_101
//BORDER_REFLECT_101: gfedcb|abcdefgh|gfedcba
#define ADDR_L(i,l_edge,r_edge,addr) (i) < (l_edge) ? -(i) : (addr)
#define ADDR_R(i,r_edge,addr) (i) >= (r_edge) ? -(i)-2+((r_edge)<<1) : (addr)
#define EXTRAPOLATE(x, maxV) EXTRAPOLATE_(x, maxV, 0)
#else
#define EXTRAPOLATE(x, maxV) EXTRAPOLATE_(x, maxV, 1)
#endif
#ifdef BORDER_WRAP
//BORDER_WRAP: cdefgh|abcdefgh|abcdefg
#define ADDR_L(i,l_edge,r_edge,addr) (i) < (l_edge) ? (i)+(r_edge) : (addr)
#define ADDR_R(i,r_edge,addr) (i) >= (r_edge) ? (i)-(r_edge) : (addr)
#else
#error No extrapolation method
#endif
/**********************************************************************************
@ -96,73 +105,71 @@ The info above maybe obsolete.
***********************************************************************************/
__kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_C1_D0
(__global const uchar * restrict src,
__global float * dst,
const int dst_cols,
const int dst_rows,
const int src_whole_cols,
const int src_whole_rows,
const int src_step_in_pixel,
const int src_offset_x,
const int src_offset_y,
const int dst_step_in_pixel,
const int radiusy,
__constant float * mat_kernel __attribute__((max_constant_size(4*(2*RADIUSX+1)))))
(__global uchar * restrict src,
__global float * dst,
int dst_cols, int dst_rows,
int src_whole_cols, int src_whole_rows,
int src_step_in_pixel,
int src_offset_x, int src_offset_y,
int dst_step_in_pixel, int radiusy,
__constant float * mat_kernel __attribute__((max_constant_size(4*(2*RADIUSX+1)))))
{
int x = get_global_id(0)<<2;
int y = get_global_id(1);
int l_x = get_local_id(0);
int l_y = get_local_id(1);
int start_x = x+src_offset_x-RADIUSX & 0xfffffffc;
int offset = src_offset_x-RADIUSX & 3;
int start_y = y+src_offset_y-radiusy;
int start_addr = mad24(start_y,src_step_in_pixel,start_x);
int start_x = x+src_offset_x - RADIUSX & 0xfffffffc;
int offset = src_offset_x - RADIUSX & 3;
int start_y = y + src_offset_y - radiusy;
int start_addr = mad24(start_y, src_step_in_pixel, start_x);
int i;
float4 sum;
uchar4 temp[READ_TIMES_ROW];
__local uchar4 LDS_DAT[LSIZE1][READ_TIMES_ROW*LSIZE0+1];
#ifdef BORDER_CONSTANT
int end_addr = mad24(src_whole_rows - 1,src_step_in_pixel,src_whole_cols);
//read pixels from src
for(i = 0; i<READ_TIMES_ROW; i++)
int end_addr = mad24(src_whole_rows - 1, src_step_in_pixel, src_whole_cols);
// read pixels from src
for (i = 0; i < READ_TIMES_ROW; i++)
{
int current_addr = start_addr+i*LSIZE0*4;
current_addr = ((current_addr < end_addr) && (current_addr > 0)) ? current_addr : 0;
temp[i] = *(__global uchar4*)&src[current_addr];
}
//judge if read out of boundary
for(i = 0; i<READ_TIMES_ROW; i++)
// judge if read out of boundary
for (i = 0; i<READ_TIMES_ROW; i++)
{
temp[i].x= ELEM(start_x+i*LSIZE0*4,0,src_whole_cols,0,temp[i].x);
temp[i].y= ELEM(start_x+i*LSIZE0*4+1,0,src_whole_cols,0,temp[i].y);
temp[i].z= ELEM(start_x+i*LSIZE0*4+2,0,src_whole_cols,0,temp[i].z);
temp[i].w= ELEM(start_x+i*LSIZE0*4+3,0,src_whole_cols,0,temp[i].w);
temp[i]= ELEM(start_y,0,src_whole_rows,(uchar4)0,temp[i]);
temp[i].x = ELEM(start_x+i*LSIZE0*4,0,src_whole_cols,0,temp[i].x);
temp[i].y = ELEM(start_x+i*LSIZE0*4+1,0,src_whole_cols,0,temp[i].y);
temp[i].z = ELEM(start_x+i*LSIZE0*4+2,0,src_whole_cols,0,temp[i].z);
temp[i].w = ELEM(start_x+i*LSIZE0*4+3,0,src_whole_cols,0,temp[i].w);
temp[i] = ELEM(start_y,0,src_whole_rows,(uchar4)0,temp[i]);
}
#else
int not_all_in_range = (start_x<0) | (start_x + READ_TIMES_ROW*LSIZE0*4+4>src_whole_cols)| (start_y<0) | (start_y >= src_whole_rows);
int4 index[READ_TIMES_ROW];
int4 addr;
int s_y;
if(not_all_in_range)
if (not_all_in_range)
{
//judge if read out of boundary
for(i = 0; i<READ_TIMES_ROW; i++)
// judge if read out of boundary
for (i = 0; i < READ_TIMES_ROW; i++)
{
index[i].x= ADDR_L(start_x+i*LSIZE0*4,0,src_whole_cols,start_x+i*LSIZE0*4);
index[i].x= ADDR_R(start_x+i*LSIZE0*4,src_whole_cols,index[i].x);
index[i].y= ADDR_L(start_x+i*LSIZE0*4+1,0,src_whole_cols,start_x+i*LSIZE0*4+1);
index[i].y= ADDR_R(start_x+i*LSIZE0*4+1,src_whole_cols,index[i].y);
index[i].z= ADDR_L(start_x+i*LSIZE0*4+2,0,src_whole_cols,start_x+i*LSIZE0*4+2);
index[i].z= ADDR_R(start_x+i*LSIZE0*4+2,src_whole_cols,index[i].z);
index[i].w= ADDR_L(start_x+i*LSIZE0*4+3,0,src_whole_cols,start_x+i*LSIZE0*4+3);
index[i].w= ADDR_R(start_x+i*LSIZE0*4+3,src_whole_cols,index[i].w);
index[i] = (int4)(start_x+i*LSIZE0*4) + (int4)(0, 1, 2, 3);
EXTRAPOLATE(index[i].x, src_whole_cols);
EXTRAPOLATE(index[i].y, src_whole_cols);
EXTRAPOLATE(index[i].z, src_whole_cols);
EXTRAPOLATE(index[i].w, src_whole_cols);
}
s_y= ADDR_L(start_y,0,src_whole_rows,start_y);
s_y= ADDR_R(start_y,src_whole_rows,s_y);
//read pixels from src
for(i = 0; i<READ_TIMES_ROW; i++)
s_y = start_y;
EXTRAPOLATE(s_y, src_whole_rows);
// read pixels from src
for (i = 0; i<READ_TIMES_ROW; i++)
{
addr = mad24((int4)s_y,(int4)src_step_in_pixel,index[i]);
temp[i].x = src[addr.x];
@ -173,64 +180,55 @@ __kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_
}
else
{
//read pixels from src
for(i = 0; i<READ_TIMES_ROW; i++)
{
// read pixels from src
for (i = 0; i<READ_TIMES_ROW; i++)
temp[i] = *(__global uchar4*)&src[start_addr+i*LSIZE0*4];
}
}
#endif
//save pixels to lds
for(i = 0; i<READ_TIMES_ROW; i++)
{
// save pixels to lds
for (i = 0; i<READ_TIMES_ROW; i++)
LDS_DAT[l_y][l_x+i*LSIZE0]=temp[i];
}
barrier(CLK_LOCAL_MEM_FENCE);
//read pixels from lds and calculate the result
// read pixels from lds and calculate the result
sum =convert_float4(vload4(0,(__local uchar*)&LDS_DAT[l_y][l_x]+RADIUSX+offset))*mat_kernel[RADIUSX];
for(i=1; i<=RADIUSX; i++)
for (i=1; i<=RADIUSX; i++)
{
temp[0]=vload4(0,(__local uchar*)&LDS_DAT[l_y][l_x]+RADIUSX+offset-i);
temp[1]=vload4(0,(__local uchar*)&LDS_DAT[l_y][l_x]+RADIUSX+offset+i);
sum += convert_float4(temp[0])*mat_kernel[RADIUSX-i]+convert_float4(temp[1])*mat_kernel[RADIUSX+i];
temp[0] = vload4(0, (__local uchar*)&LDS_DAT[l_y][l_x] + RADIUSX + offset - i);
temp[1] = vload4(0, (__local uchar*)&LDS_DAT[l_y][l_x] + RADIUSX + offset + i);
sum += convert_float4(temp[0]) * mat_kernel[RADIUSX-i] + convert_float4(temp[1]) * mat_kernel[RADIUSX+i];
}
start_addr = mad24(y,dst_step_in_pixel,x);
//write the result to dst
if((x+3<dst_cols) & (y<dst_rows))
{
// write the result to dst
if ((x+3<dst_cols) & (y<dst_rows))
*(__global float4*)&dst[start_addr] = sum;
}
else if((x+2<dst_cols) & (y<dst_rows))
else if ((x+2<dst_cols) && (y<dst_rows))
{
dst[start_addr] = sum.x;
dst[start_addr+1] = sum.y;
dst[start_addr+2] = sum.z;
}
else if((x+1<dst_cols) & (y<dst_rows))
else if ((x+1<dst_cols) && (y<dst_rows))
{
dst[start_addr] = sum.x;
dst[start_addr+1] = sum.y;
}
else if((x<dst_cols) & (y<dst_rows))
{
else if (x<dst_cols && y<dst_rows)
dst[start_addr] = sum.x;
}
}
__kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_C4_D0
(__global const uchar4 * restrict src,
__global float4 * dst,
const int dst_cols,
const int dst_rows,
const int src_whole_cols,
const int src_whole_rows,
const int src_step_in_pixel,
const int src_offset_x,
const int src_offset_y,
const int dst_step_in_pixel,
const int radiusy,
__constant float * mat_kernel __attribute__((max_constant_size(4*(2*RADIUSX+1)))))
(__global uchar4 * restrict src,
__global float4 * dst,
int dst_cols, int dst_rows,
int src_whole_cols, int src_whole_rows,
int src_step_in_pixel,
int src_offset_x, int src_offset_y,
int dst_step_in_pixel, int radiusy,
__constant float * mat_kernel __attribute__((max_constant_size(4*(2*RADIUSX+1)))))
{
int x = get_global_id(0);
int y = get_global_id(1);
@ -246,15 +244,17 @@ __kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_
__local uchar4 LDS_DAT[LSIZE1][READ_TIMES_ROW*LSIZE0+1];
#ifdef BORDER_CONSTANT
int end_addr = mad24(src_whole_rows - 1,src_step_in_pixel,src_whole_cols);
//read pixels from src
for(i = 0; i<READ_TIMES_ROW; i++)
// read pixels from src
for (i = 0; i<READ_TIMES_ROW; i++)
{
int current_addr = start_addr+i*LSIZE0;
current_addr = ((current_addr < end_addr) && (current_addr > 0)) ? current_addr : 0;
temp[i] = src[current_addr];
}
//judge if read out of boundary
for(i = 0; i<READ_TIMES_ROW; i++)
for (i = 0; i<READ_TIMES_ROW; i++)
{
temp[i]= ELEM(start_x+i*LSIZE0,0,src_whole_cols,(uchar4)0,temp[i]);
temp[i]= ELEM(start_y,0,src_whole_rows,(uchar4)0,temp[i]);
@ -262,39 +262,37 @@ __kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_
#else
int index[READ_TIMES_ROW];
int s_x,s_y;
//judge if read out of boundary
for(i = 0; i<READ_TIMES_ROW; i++)
// judge if read out of boundary
for (i = 0; i<READ_TIMES_ROW; i++)
{
s_x= ADDR_L(start_x+i*LSIZE0,0,src_whole_cols,start_x+i*LSIZE0);
s_x= ADDR_R(start_x+i*LSIZE0,src_whole_cols,s_x);
s_y= ADDR_L(start_y,0,src_whole_rows,start_y);
s_y= ADDR_R(start_y,src_whole_rows,s_y);
s_x = start_x+i*LSIZE0;
EXTRAPOLATE(s_x, src_whole_cols);
s_y = start_y;
EXTRAPOLATE(s_y, src_whole_rows);
index[i]=mad24(s_y,src_step_in_pixel,s_x);
}
//read pixels from src
for(i = 0; i<READ_TIMES_ROW; i++)
{
for (i = 0; i<READ_TIMES_ROW; i++)
temp[i] = src[index[i]];
}
#endif
//save pixels to lds
for(i = 0; i<READ_TIMES_ROW; i++)
{
for (i = 0; i<READ_TIMES_ROW; i++)
LDS_DAT[l_y][l_x+i*LSIZE0]=temp[i];
}
barrier(CLK_LOCAL_MEM_FENCE);
//read pixels from lds and calculate the result
sum =convert_float4(LDS_DAT[l_y][l_x+RADIUSX])*mat_kernel[RADIUSX];
for(i=1; i<=RADIUSX; i++)
for (i=1; i<=RADIUSX; i++)
{
temp[0]=LDS_DAT[l_y][l_x+RADIUSX-i];
temp[1]=LDS_DAT[l_y][l_x+RADIUSX+i];
sum += convert_float4(temp[0])*mat_kernel[RADIUSX-i]+convert_float4(temp[1])*mat_kernel[RADIUSX+i];
}
//write the result to dst
if((x<dst_cols) & (y<dst_rows))
if (x<dst_cols && y<dst_rows)
{
start_addr = mad24(y,dst_step_in_pixel,x);
dst[start_addr] = sum;
@ -302,18 +300,14 @@ __kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_
}
__kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_C1_D5
(__global const float * restrict src,
__global float * dst,
const int dst_cols,
const int dst_rows,
const int src_whole_cols,
const int src_whole_rows,
const int src_step_in_pixel,
const int src_offset_x,
const int src_offset_y,
const int dst_step_in_pixel,
const int radiusy,
__constant float * mat_kernel __attribute__((max_constant_size(4*(2*RADIUSX+1)))))
(__global float * restrict src,
__global float * dst,
int dst_cols, int dst_rows,
int src_whole_cols, int src_whole_rows,
int src_step_in_pixel,
int src_offset_x, int src_offset_y,
int dst_step_in_pixel, int radiusy,
__constant float * mat_kernel __attribute__((max_constant_size(4*(2*RADIUSX+1)))))
{
int x = get_global_id(0);
int y = get_global_id(1);
@ -329,15 +323,17 @@ __kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_
__local float LDS_DAT[LSIZE1][READ_TIMES_ROW*LSIZE0+1];
#ifdef BORDER_CONSTANT
int end_addr = mad24(src_whole_rows - 1,src_step_in_pixel,src_whole_cols);
//read pixels from src
for(i = 0; i<READ_TIMES_ROW; i++)
// read pixels from src
for (i = 0; i<READ_TIMES_ROW; i++)
{
int current_addr = start_addr+i*LSIZE0;
current_addr = ((current_addr < end_addr) && (current_addr > 0)) ? current_addr : 0;
temp[i] = src[current_addr];
}
//judge if read out of boundary
for(i = 0; i<READ_TIMES_ROW; i++)
// judge if read out of boundary
for (i = 0; i<READ_TIMES_ROW; i++)
{
temp[i]= ELEM(start_x+i*LSIZE0,0,src_whole_cols,(float)0,temp[i]);
temp[i]= ELEM(start_y,0,src_whole_rows,(float)0,temp[i]);
@ -345,39 +341,36 @@ __kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_
#else
int index[READ_TIMES_ROW];
int s_x,s_y;
//judge if read out of boundary
for(i = 0; i<READ_TIMES_ROW; i++)
// judge if read out of boundary
for (i = 0; i<READ_TIMES_ROW; i++)
{
s_x= ADDR_L(start_x+i*LSIZE0,0,src_whole_cols,start_x+i*LSIZE0);
s_x= ADDR_R(start_x+i*LSIZE0,src_whole_cols,s_x);
s_y= ADDR_L(start_y,0,src_whole_rows,start_y);
s_y= ADDR_R(start_y,src_whole_rows,s_y);
index[i]=mad24(s_y,src_step_in_pixel,s_x);
s_x = start_x + i*LSIZE0, s_y = start_y;
EXTRAPOLATE(s_x, src_whole_cols);
EXTRAPOLATE(s_y, src_whole_rows);
index[i]=mad24(s_y, src_step_in_pixel, s_x);
}
//read pixels from src
for(i = 0; i<READ_TIMES_ROW; i++)
{
// read pixels from src
for (i = 0; i<READ_TIMES_ROW; i++)
temp[i] = src[index[i]];
}
#endif
//save pixels to lds
for(i = 0; i<READ_TIMES_ROW; i++)
{
for (i = 0; i<READ_TIMES_ROW; i++)
LDS_DAT[l_y][l_x+i*LSIZE0]=temp[i];
}
barrier(CLK_LOCAL_MEM_FENCE);
//read pixels from lds and calculate the result
// read pixels from lds and calculate the result
sum =LDS_DAT[l_y][l_x+RADIUSX]*mat_kernel[RADIUSX];
for(i=1; i<=RADIUSX; i++)
for (i=1; i<=RADIUSX; i++)
{
temp[0]=LDS_DAT[l_y][l_x+RADIUSX-i];
temp[1]=LDS_DAT[l_y][l_x+RADIUSX+i];
sum += temp[0]*mat_kernel[RADIUSX-i]+temp[1]*mat_kernel[RADIUSX+i];
}
//write the result to dst
if((x<dst_cols) & (y<dst_rows))
// write the result to dst
if (x<dst_cols && y<dst_rows)
{
start_addr = mad24(y,dst_step_in_pixel,x);
dst[start_addr] = sum;
@ -385,18 +378,14 @@ __kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_
}
__kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_C4_D5
(__global const float4 * restrict src,
__global float4 * dst,
const int dst_cols,
const int dst_rows,
const int src_whole_cols,
const int src_whole_rows,
const int src_step_in_pixel,
const int src_offset_x,
const int src_offset_y,
const int dst_step_in_pixel,
const int radiusy,
__constant float * mat_kernel __attribute__((max_constant_size(4*(2*RADIUSX+1)))))
(__global float4 * restrict src,
__global float4 * dst,
int dst_cols, int dst_rows,
int src_whole_cols, int src_whole_rows,
int src_step_in_pixel,
int src_offset_x, int src_offset_y,
int dst_step_in_pixel, int radiusy,
__constant float * mat_kernel __attribute__((max_constant_size(4*(2*RADIUSX+1)))))
{
int x = get_global_id(0);
int y = get_global_id(1);
@ -412,15 +401,17 @@ __kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_
__local float4 LDS_DAT[LSIZE1][READ_TIMES_ROW*LSIZE0+1];
#ifdef BORDER_CONSTANT
int end_addr = mad24(src_whole_rows - 1,src_step_in_pixel,src_whole_cols);
//read pixels from src
for(i = 0; i<READ_TIMES_ROW; i++)
// read pixels from src
for (i = 0; i<READ_TIMES_ROW; i++)
{
int current_addr = start_addr+i*LSIZE0;
current_addr = ((current_addr < end_addr) && (current_addr > 0)) ? current_addr : 0;
temp[i] = src[current_addr];
}
//judge if read out of boundary
for(i = 0; i<READ_TIMES_ROW; i++)
// judge if read out of boundary
for (i = 0; i<READ_TIMES_ROW; i++)
{
temp[i]= ELEM(start_x+i*LSIZE0,0,src_whole_cols,(float4)0,temp[i]);
temp[i]= ELEM(start_y,0,src_whole_rows,(float4)0,temp[i]);
@ -428,42 +419,39 @@ __kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_
#else
int index[READ_TIMES_ROW];
int s_x,s_y;
//judge if read out of boundary
for(i = 0; i<READ_TIMES_ROW; i++)
// judge if read out of boundary
for (i = 0; i<READ_TIMES_ROW; i++)
{
s_x= ADDR_L(start_x+i*LSIZE0,0,src_whole_cols,start_x+i*LSIZE0);
s_x= ADDR_R(start_x+i*LSIZE0,src_whole_cols,s_x);
s_y= ADDR_L(start_y,0,src_whole_rows,start_y);
s_y= ADDR_R(start_y,src_whole_rows,s_y);
s_x = start_x + i*LSIZE0, s_y = start_y;
EXTRAPOLATE(s_x, src_whole_cols);
EXTRAPOLATE(s_y, src_whole_rows);
index[i]=mad24(s_y,src_step_in_pixel,s_x);
}
//read pixels from src
for(i = 0; i<READ_TIMES_ROW; i++)
{
// read pixels from src
for (i = 0; i<READ_TIMES_ROW; i++)
temp[i] = src[index[i]];
}
#endif
//save pixels to lds
for(i = 0; i<READ_TIMES_ROW; i++)
{
// save pixels to lds
for (i = 0; i<READ_TIMES_ROW; i++)
LDS_DAT[l_y][l_x+i*LSIZE0]=temp[i];
}
barrier(CLK_LOCAL_MEM_FENCE);
//read pixels from lds and calculate the result
// read pixels from lds and calculate the result
sum =LDS_DAT[l_y][l_x+RADIUSX]*mat_kernel[RADIUSX];
for(i=1; i<=RADIUSX; i++)
for (i=1; i<=RADIUSX; i++)
{
temp[0]=LDS_DAT[l_y][l_x+RADIUSX-i];
temp[1]=LDS_DAT[l_y][l_x+RADIUSX+i];
sum += temp[0]*mat_kernel[RADIUSX-i]+temp[1]*mat_kernel[RADIUSX+i];
}
//write the result to dst
if((x<dst_cols) & (y<dst_rows))
// write the result to dst
if (x<dst_cols && y<dst_rows)
{
start_addr = mad24(y,dst_step_in_pixel,x);
dst[start_addr] = sum;
}
}

@ -27,7 +27,7 @@
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other oclMaterials provided with the distribution.
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
@ -45,38 +45,43 @@
//
//M*/
#ifdef BORDER_REPLICATE
//BORDER_REPLICATE: aaaaaa|abcdefgh|hhhhhhh
#define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? (l_edge) : (i))
#define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? (r_edge)-1 : (addr))
#define ADDR_H(i, t_edge, b_edge) ((i) < (t_edge) ? (t_edge) :(i))
#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? (b_edge)-1 :(addr))
#endif
#ifdef BORDER_CONSTANT
#define ELEM(i,l_edge,r_edge,elem1,elem2) (i)<(l_edge) | (i) >= (r_edge) ? (elem1) : (elem2)
#elif defined BORDER_REPLICATE
#define EXTRAPOLATE(x, maxV) \
{ \
x = max(min(x, maxV - 1), 0); \
}
#elif defined BORDER_WRAP
#define EXTRAPOLATE(x, maxV) \
{ \
if (x < 0) \
x -= ((x - maxV + 1) / maxV) * maxV; \
if (x >= maxV) \
x %= maxV; \
}
#elif defined(BORDER_REFLECT) || defined(BORDER_REFLECT_101)
#define EXTRAPOLATE_(x, maxV, delta) \
{ \
if (maxV == 1) \
x = 0; \
else \
do \
{ \
if ( x < 0 ) \
x = -x - 1 + delta; \
else \
x = maxV - 1 - (x - maxV) - delta; \
} \
while (x >= maxV || x < 0); \
}
#ifdef BORDER_REFLECT
//BORDER_REFLECT: fedcba|abcdefgh|hgfedcb
#define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? -(i)-1 : (i))
#define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? -(i)-1+((r_edge)<<1) : (addr))
#define ADDR_H(i, t_edge, b_edge) ((i) < (t_edge) ? -(i)-1 : (i))
#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? -(i)-1+((b_edge)<<1) : (addr))
#endif
#ifdef BORDER_REFLECT_101
//BORDER_REFLECT_101: gfedcb|abcdefgh|gfedcba
#define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? -(i) : (i))
#define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? -(i)-2+((r_edge)<<1) : (addr))
#define ADDR_H(i, t_edge, b_edge) ((i) < (t_edge) ? -(i) : (i))
#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? -(i)-2+((b_edge)<<1) : (addr))
#define EXTRAPOLATE(x, maxV) EXTRAPOLATE_(x, maxV, 0)
#else
#define EXTRAPOLATE(x, maxV) EXTRAPOLATE_(x, maxV, 1)
#endif
//blur function does not support BORDER_WRAP
#ifdef BORDER_WRAP
//BORDER_WRAP: cdefgh|abcdefgh|abcdefg
#define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? (i)+(r_edge) : (i))
#define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? (i)-(r_edge) : (addr))
#define ADDR_H(i, t_edge, b_edge) ((i) < (t_edge) ? (i)+(b_edge) : (i))
#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? (i)-(b_edge) : (addr))
#else
#error No extrapolation method
#endif
__kernel void
@ -117,9 +122,7 @@ edgeEnhancingFilter_C4_D0(
float4 tmp_sum[1+EXTRA];
for(int tmpint = 0; tmpint < 1+EXTRA; tmpint++)
{
tmp_sum[tmpint] = (float4)(0,0,0,0);
}
#ifdef BORDER_CONSTANT
bool con;
@ -127,25 +130,18 @@ edgeEnhancingFilter_C4_D0(
for(int j = 0; j < ksY+EXTRA; j++)
{
con = (startX+col >= 0 && startX+col < src_whole_cols && startY+j >= 0 && startY+j < src_whole_rows);
int cur_col = clamp(startX + col, 0, src_whole_cols);
if(con)
{
if (con)
ss = src[(startY+j)*(src_step>>2) + cur_col];
}
data[j][col] = con ? ss : (uchar4)0;
}
#else
for(int j= 0; j < ksY+EXTRA; j++)
{
int selected_row;
int selected_col;
selected_row = ADDR_H(startY+j, 0, src_whole_rows);
selected_row = ADDR_B(startY+j, src_whole_rows, selected_row);
selected_col = ADDR_L(startX+col, 0, src_whole_cols);
selected_col = ADDR_R(startX+col, src_whole_cols, selected_col);
int selected_row = startY+j, selected_col = startX+col;
EXTRAPOLATE(selected_row, src_whole_rows)
EXTRAPOLATE(selected_col, src_whole_cols)
data[j][col] = src[selected_row * (src_step>>2) + selected_col];
}
@ -172,7 +168,6 @@ edgeEnhancingFilter_C4_D0(
if(col < (THREADS-(ksX-1)))
{
int4 currVal;
int howManyAll = (2*anX+1)*(ksY);
//find variance of all data
@ -187,15 +182,14 @@ edgeEnhancingFilter_C4_D0(
sumVal =0;
sumValSqr=0;
for(int j = startLMj; j < endLMj; j++)
{
for(int i=-anX; i<=anX; i++)
{
currVal = convert_int4(data[j][col+anX+i]) ;
currVal = convert_int4(data[j][col+anX+i]);
sumVal += currVal;
sumValSqr += mul24(currVal, currVal);
}
}
var[extraCnt] = convert_float4( ( (sumValSqr * howManyAll)- mul24(sumVal , sumVal) ) ) / ( (float)(howManyAll*howManyAll) ) ;
#else
var[extraCnt] = (float4)(900.0, 900.0, 900.0, 0.0);
@ -228,17 +222,15 @@ edgeEnhancingFilter_C4_D0(
weight = 1.0f;
#endif
#else
currVal = convert_int4(data[j][col+anX+i]) ;
currVal = convert_int4(data[j][col+anX+i]);
currWRTCenter = currVal-currValCenter;
#if VAR_PER_CHANNEL
weight = var[extraCnt] / (var[extraCnt] + convert_float4(currWRTCenter * currWRTCenter)) * (float4)(lut[lut_j*lut_step+anX+i]);
//weight.x = var[extraCnt].x / ( var[extraCnt].x + (float) mul24(currWRTCenter.x , currWRTCenter.x) ) ;
//weight.y = var[extraCnt].y / ( var[extraCnt].y + (float) mul24(currWRTCenter.y , currWRTCenter.y) ) ;
//weight.z = var[extraCnt].z / ( var[extraCnt].z + (float) mul24(currWRTCenter.z , currWRTCenter.z) ) ;
//weight.w = 0;
weight = var[extraCnt] / (var[extraCnt] + convert_float4(currWRTCenter * currWRTCenter)) *
(float4)(lut[lut_j*lut_step+anX+i]);
#else
weight = 1.0f/(1.0f+( mul24(currWRTCenter.x, currWRTCenter.x) + mul24(currWRTCenter.y, currWRTCenter.y) + mul24(currWRTCenter.z, currWRTCenter.z))/(var.x+var.y+var.z));
weight = 1.0f/(1.0f+( mul24(currWRTCenter.x, currWRTCenter.x) + mul24(currWRTCenter.y, currWRTCenter.y) +
mul24(currWRTCenter.z, currWRTCenter.z))/(var.x+var.y+var.z));
#endif
#endif
tmp_sum[extraCnt] += convert_float4(data[j][col+anX+i]) * weight;
@ -249,9 +241,7 @@ edgeEnhancingFilter_C4_D0(
tmp_sum[extraCnt] /= totalWeight;
if(posX >= 0 && posX < dst_cols && (posY+extraCnt) >= 0 && (posY+extraCnt) < dst_rows)
{
dst[(dst_startY+extraCnt) * (dst_step>>2)+ dst_startX + col] = convert_uchar4(tmp_sum[extraCnt]);
}
#if VAR_PER_CHANNEL
totalWeight = (float4)(0,0,0,0);
@ -323,13 +313,9 @@ edgeEnhancingFilter_C1_D0(
#else
for(int j= 0; j < ksY+EXTRA; j++)
{
int selected_row;
int selected_col;
selected_row = ADDR_H(startY+j, 0, src_whole_rows);
selected_row = ADDR_B(startY+j, src_whole_rows, selected_row);
selected_col = ADDR_L(startX+col, 0, src_whole_cols);
selected_col = ADDR_R(startX+col, src_whole_cols, selected_col);
int selected_row = startY+j, selected_col = startX+col;
EXTRAPOLATE(selected_row, src_whole_rows)
EXTRAPOLATE(selected_col, src_whole_cols)
data[j][col] = src[selected_row * (src_step) + selected_col];
}

Some files were not shown because too many files have changed in this diff Show More

Loading…
Cancel
Save