From 48808581190d3076b579c65498337a1fcfb97b20 Mon Sep 17 00:00:00 2001 From: GregoryMorse Date: Mon, 23 Dec 2013 00:28:50 +0800 Subject: [PATCH 1/2] Update CMakeLists.txt WinRT native C++ support allowing building of static libraries Update CMakeLists.txt Update OpenCVCRTLinkage.cmake Update OpenCVCRTLinkage.cmake --- CMakeLists.txt | 3 ++- cmake/OpenCVCRTLinkage.cmake | 12 ++++++++---- modules/core/CMakeLists.txt | 5 ++++- 3 files changed, 14 insertions(+), 6 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 3f793f107..daf185fba 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -219,6 +219,7 @@ OCV_OPTION(ENABLE_VFPV3 "Enable VFPv3-D32 instructions" OCV_OPTION(ENABLE_NOISY_WARNINGS "Show all warnings even if they are too noisy" OFF ) OCV_OPTION(OPENCV_WARNINGS_ARE_ERRORS "Treat warnings as errors" OFF ) OCV_OPTION(ENABLE_WINRT_MODE "Build with Windows Runtime support" OFF IF WIN32 ) +OCV_OPTION(ENABLE_WINRT_MODE_NATIVE "Build with Windows Runtime native C++ support" OFF IF WIN32 ) # uncategorized options # =================================================== @@ -660,7 +661,7 @@ endif() if(WIN32) status("") status(" Windows RT support:" HAVE_WINRT THEN YES ELSE NO) - if (ENABLE_WINRT_MODE) + if (ENABLE_WINRT_MODE OR ENABLE_WINRT_MODE_NATIVE) status(" Windows SDK v8.0:" ${WINDOWS_SDK_PATH}) status(" Visual Studio 2012:" ${VISUAL_STUDIO_PATH}) endif() diff --git a/cmake/OpenCVCRTLinkage.cmake b/cmake/OpenCVCRTLinkage.cmake index 8a297c685..5265e3e8a 100644 --- a/cmake/OpenCVCRTLinkage.cmake +++ b/cmake/OpenCVCRTLinkage.cmake @@ -9,7 +9,7 @@ set(HAVE_WINRT FALSE) # search Windows Platform SDK message(STATUS "Checking for Windows Platform SDK") GET_FILENAME_COMPONENT(WINDOWS_SDK_PATH "[HKEY_LOCAL_MACHINE\\SOFTWARE\\Microsoft\\Microsoft SDKs\\Windows\\v8.0;InstallationFolder]" ABSOLUTE CACHE) -if (WINDOWS_SDK_PATH STREQUAL "") +if(WINDOWS_SDK_PATH STREQUAL "") set(HAVE_MSPDK FALSE) message(STATUS "Windows Platform SDK 8.0 was not found") else() @@ -19,7 +19,7 @@ endif() #search for Visual Studio 11.0 install directory message(STATUS "Checking for Visual Studio 2012") GET_FILENAME_COMPONENT(VISUAL_STUDIO_PATH [HKEY_LOCAL_MACHINE\\SOFTWARE\\Microsoft\\VisualStudio\\11.0\\Setup\\VS;ProductDir] REALPATH CACHE) -if (VISUAL_STUDIO_PATH STREQUAL "") +if(VISUAL_STUDIO_PATH STREQUAL "") set(HAVE_MSVC2012 FALSE) message(STATUS "Visual Studio 2012 was not found") else() @@ -30,11 +30,15 @@ try_compile(HAVE_WINRT_SDK "${OpenCV_BINARY_DIR}" "${OpenCV_SOURCE_DIR}/cmake/checks/winrttest.cpp") -if (ENABLE_WINRT_MODE AND HAVE_WINRT_SDK AND HAVE_MSVC2012 AND HAVE_MSPDK) +if(ENABLE_WINRT_MODE AND HAVE_WINRT_SDK AND HAVE_MSVC2012 AND HAVE_MSPDK) set(HAVE_WINRT TRUE) + set(HAVE_WINRT_CX TRUE) +elseif(ENABLE_WINRT_MODE_NATIVE AND HAVE_WINRT_SDK AND HAVE_MSVC2012 AND HAVE_MSPDK) + set(HAVE_WINRT TRUE) + set(HAVE_WINRT_CX FALSE) endif() -if (HAVE_WINRT) +if(HAVE_WINRT) add_definitions(/DWINVER=0x0602 /DNTDDI_VERSION=NTDDI_WIN8 /D_WIN32_WINNT=0x0602) set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} /appcontainer") set(CMAKE_MODULE_LINKER_FLAGS "${CMAKE_MODULE_LINKER_FLAGS} /appcontainer") diff --git a/modules/core/CMakeLists.txt b/modules/core/CMakeLists.txt index 66b8ae0d2..2adf5dbbd 100644 --- a/modules/core/CMakeLists.txt +++ b/modules/core/CMakeLists.txt @@ -2,8 +2,11 @@ set(the_description "The Core Functionality") ocv_add_module(core PRIVATE_REQUIRED ${ZLIB_LIBRARIES}) ocv_module_include_directories(${ZLIB_INCLUDE_DIR}) +if(HAVE_WINRT_CX) + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /ZW") +endif() if(HAVE_WINRT) - set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /ZW /GS /Gm- /AI\"${WINDOWS_SDK_PATH}/References/CommonConfiguration/Neutral\" /AI\"${VISUAL_STUDIO_PATH}/vcpackages\"") + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /GS /Gm- /AI\"${WINDOWS_SDK_PATH}/References/CommonConfiguration/Neutral\" /AI\"${VISUAL_STUDIO_PATH}/vcpackages\"") endif() if(HAVE_CUDA) From 4f6f6e8cacfec0cfac430a63a41a4ed62ee70492 Mon Sep 17 00:00:00 2001 From: Alexander Smorkalov Date: Thu, 26 Dec 2013 21:20:32 +0400 Subject: [PATCH 2/2] static function qualifier replaced on inline to enable kernel compilation with OpenCL 1.1 embedded profile. --- modules/ocl/src/opencl/bgfg_mog.cl | 8 ++++---- modules/ocl/src/opencl/kmeans_kernel.cl | 2 +- modules/ocl/src/opencl/meanShift.cl | 2 +- modules/ocl/src/opencl/objdetect_hog.cl | 2 +- modules/ocl/src/opencl/pyrlk.cl | 20 ++++++++++---------- modules/ocl/src/opencl/stereobp.cl | 4 ++-- modules/ocl/src/opencl/tvl1flow.cl | 6 +++--- 7 files changed, 22 insertions(+), 22 deletions(-) diff --git a/modules/ocl/src/opencl/bgfg_mog.cl b/modules/ocl/src/opencl/bgfg_mog.cl index 06e18c213..6a95316f0 100644 --- a/modules/ocl/src/opencl/bgfg_mog.cl +++ b/modules/ocl/src/opencl/bgfg_mog.cl @@ -63,7 +63,7 @@ inline float sum(float val) return val; } -static float clamp1(float var, float learningRate, float diff, float minVar) +inline float clamp1(float var, float learningRate, float diff, float minVar) { return fmax(var + learningRate * (diff * diff - var), minVar); } @@ -96,7 +96,7 @@ inline float sum(const float4 val) return (val.x + val.y + val.z); } -static void swap4(__global float4* ptr, int x, int y, int k, int rows, int ptr_step) +inline void swap4(__global float4* ptr, int x, int y, int k, int rows, int ptr_step) { float4 val = ptr[(k * rows + y) * ptr_step + x]; ptr[(k * rows + y) * ptr_step + x] = ptr[((k + 1) * rows + y) * ptr_step + x]; @@ -104,7 +104,7 @@ static void swap4(__global float4* ptr, int x, int y, int k, int rows, int ptr_s } -static float4 clamp1(const float4 var, float learningRate, const float4 diff, float minVar) +inline float4 clamp1(const float4 var, float learningRate, const float4 diff, float minVar) { float4 result; result.x = fmax(var.x + learningRate * (diff.x * diff.x - var.x), minVar); @@ -128,7 +128,7 @@ typedef struct uchar c_shadowVal; } con_srtuct_t; -static void swap(__global float* ptr, int x, int y, int k, int rows, int ptr_step) +inline void swap(__global float* ptr, int x, int y, int k, int rows, int ptr_step) { float val = ptr[(k * rows + y) * ptr_step + x]; ptr[(k * rows + y) * ptr_step + x] = ptr[((k + 1) * rows + y) * ptr_step + x]; diff --git a/modules/ocl/src/opencl/kmeans_kernel.cl b/modules/ocl/src/opencl/kmeans_kernel.cl index 244d52ca3..bb0e9c9a4 100644 --- a/modules/ocl/src/opencl/kmeans_kernel.cl +++ b/modules/ocl/src/opencl/kmeans_kernel.cl @@ -44,7 +44,7 @@ // //M*/ -static float distance_(__global const float * center, __global const float * src, int feature_length) +inline float distance_(__global const float * center, __global const float * src, int feature_length) { float res = 0; float4 v0, v1, v2; diff --git a/modules/ocl/src/opencl/meanShift.cl b/modules/ocl/src/opencl/meanShift.cl index ea5060e46..3fff473a8 100644 --- a/modules/ocl/src/opencl/meanShift.cl +++ b/modules/ocl/src/opencl/meanShift.cl @@ -46,7 +46,7 @@ // //M*/ -static short2 do_mean_shift(int x0, int y0, __global uchar4* out,int out_step, +inline short2 do_mean_shift(int x0, int y0, __global uchar4* out,int out_step, __global uchar4* in, int in_step, int dst_off, int src_off, int cols, int rows, int sp, int sr, int maxIter, float eps) { diff --git a/modules/ocl/src/opencl/objdetect_hog.cl b/modules/ocl/src/opencl/objdetect_hog.cl index 60d7346e5..e931e82b5 100644 --- a/modules/ocl/src/opencl/objdetect_hog.cl +++ b/modules/ocl/src/opencl/objdetect_hog.cl @@ -208,7 +208,7 @@ __kernel void normalize_hists_36_kernel(__global float* block_hists, //------------------------------------------------------------- // Normalization of histograms via L2Hys_norm // -static float reduce_smem(volatile __local float* smem, int size) +inline float reduce_smem(volatile __local float* smem, int size) { unsigned int tid = get_local_id(0); float sum = smem[tid]; diff --git a/modules/ocl/src/opencl/pyrlk.cl b/modules/ocl/src/opencl/pyrlk.cl index 303d26892..f34aee900 100644 --- a/modules/ocl/src/opencl/pyrlk.cl +++ b/modules/ocl/src/opencl/pyrlk.cl @@ -52,7 +52,7 @@ #endif #ifdef CPU -static void reduce3(float val1, float val2, float val3, __local float* smem1, __local float* smem2, __local float* smem3, int tid) +inline void reduce3(float val1, float val2, float val3, __local float* smem1, __local float* smem2, __local float* smem3, int tid) { smem1[tid] = val1; smem2[tid] = val2; @@ -71,7 +71,7 @@ static void reduce3(float val1, float val2, float val3, __local float* smem1, } } -static void reduce2(float val1, float val2, volatile __local float* smem1, volatile __local float* smem2, int tid) +inline void reduce2(float val1, float val2, volatile __local float* smem1, volatile __local float* smem2, int tid) { smem1[tid] = val1; smem2[tid] = val2; @@ -88,7 +88,7 @@ static void reduce2(float val1, float val2, volatile __local float* smem1, volat } } -static void reduce1(float val1, volatile __local float* smem1, int tid) +inline void reduce1(float val1, volatile __local float* smem1, int tid) { smem1[tid] = val1; barrier(CLK_LOCAL_MEM_FENCE); @@ -103,7 +103,7 @@ static void reduce1(float val1, volatile __local float* smem1, int tid) } } #else -static void reduce3(float val1, float val2, float val3, +inline void reduce3(float val1, float val2, float val3, __local volatile float* smem1, __local volatile float* smem2, __local volatile float* smem3, int tid) { smem1[tid] = val1; @@ -150,7 +150,7 @@ static void reduce3(float val1, float val2, float val3, barrier(CLK_LOCAL_MEM_FENCE); } -static void reduce2(float val1, float val2, __local volatile float* smem1, __local volatile float* smem2, int tid) +inline void reduce2(float val1, float val2, __local volatile float* smem1, __local volatile float* smem2, int tid) { smem1[tid] = val1; smem2[tid] = val2; @@ -189,7 +189,7 @@ static void reduce2(float val1, float val2, __local volatile float* smem1, __loc barrier(CLK_LOCAL_MEM_FENCE); } -static void reduce1(float val1, __local volatile float* smem1, int tid) +inline void reduce1(float val1, __local volatile float* smem1, int tid) { smem1[tid] = val1; barrier(CLK_LOCAL_MEM_FENCE); @@ -225,7 +225,7 @@ static void reduce1(float val1, __local volatile float* smem1, int tid) // Image read mode __constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR; -static void SetPatch(image2d_t I, float x, float y, +inline void SetPatch(image2d_t I, float x, float y, float* Pch, float* Dx, float* Dy, float* A11, float* A12, float* A22) { @@ -262,7 +262,7 @@ inline void GetError(image2d_t J, const float x, const float y, const float* Pch *errval += fabs(diff); } -static void SetPatch4(image2d_t I, const float x, const float y, +inline void SetPatch4(image2d_t I, const float x, const float y, float4* Pch, float4* Dx, float4* Dy, float* A11, float* A12, float* A22) { @@ -285,7 +285,7 @@ static void SetPatch4(image2d_t I, const float x, const float y, *A22 += sqIdx.x + sqIdx.y + sqIdx.z; } -static void GetPatch4(image2d_t J, const float x, const float y, +inline void GetPatch4(image2d_t J, const float x, const float y, const float4* Pch, const float4* Dx, const float4* Dy, float* b1, float* b2) { @@ -297,7 +297,7 @@ static void GetPatch4(image2d_t J, const float x, const float y, *b2 += xdiff.x + xdiff.y + xdiff.z; } -static void GetError4(image2d_t J, const float x, const float y, const float4* Pch, float* errval) +inline void GetError4(image2d_t J, const float x, const float y, const float4* Pch, float* errval) { float4 diff = read_imagef(J, sampler, (float2)(x,y))-*Pch; *errval += fabs(diff.x) + fabs(diff.y) + fabs(diff.z); diff --git a/modules/ocl/src/opencl/stereobp.cl b/modules/ocl/src/opencl/stereobp.cl index 4b5864f4c..5a1bf088c 100644 --- a/modules/ocl/src/opencl/stereobp.cl +++ b/modules/ocl/src/opencl/stereobp.cl @@ -97,7 +97,7 @@ inline float pix_diff_1(const uchar4 l, __global const uchar *rs) return abs((int)(l.x) - *rs); } -static float pix_diff_4(const uchar4 l, __global const uchar *rs) +inline float pix_diff_4(const uchar4 l, __global const uchar *rs) { uchar4 r; r = *((__global uchar4 *)rs); @@ -233,7 +233,7 @@ __kernel void level_up_message(__global T *src, int src_rows, int src_step, /////////////////////////////////////////////////////////////// //////////////////// calc all iterations ///////////////////// /////////////////////////////////////////////////////////////// -static void message(__global T *us_, __global T *ds_, __global T *ls_, __global T *rs_, +inline void message(__global T *us_, __global T *ds_, __global T *ls_, __global T *rs_, const __global T *dt, int u_step, int msg_disp_step, int data_disp_step, float4 cmax_disc_term, float4 cdisc_single_jump) diff --git a/modules/ocl/src/opencl/tvl1flow.cl b/modules/ocl/src/opencl/tvl1flow.cl index 6111a4a38..b488e8969 100644 --- a/modules/ocl/src/opencl/tvl1flow.cl +++ b/modules/ocl/src/opencl/tvl1flow.cl @@ -62,7 +62,7 @@ __kernel void centeredGradientKernel(__global const float* src, int src_col, int } -static float bicubicCoeff(float x_) +inline float bicubicCoeff(float x_) { float x = fabs(x_); @@ -156,7 +156,7 @@ __kernel void warpBackwardKernel(__global const float* I0, int I0_step, int I0_c } -static float readImage(__global float *image, int x, int y, int rows, int cols, int elemCntPerRow) +inline float readImage(__global float *image, int x, int y, int rows, int cols, int elemCntPerRow) { int i0 = clamp(x, 0, cols - 1); int j0 = clamp(y, 0, rows - 1); @@ -284,7 +284,7 @@ __kernel void estimateDualVariablesKernel(__global const float* u1, int u1_col, } -static float divergence(__global const float* v1, __global const float* v2, int y, int x, int v1_step, int v2_step) +inline float divergence(__global const float* v1, __global const float* v2, int y, int x, int v1_step, int v2_step) { if (x > 0 && y > 0)