From 08d8faf9daf2647d3701ac2807ded394d6308cb0 Mon Sep 17 00:00:00 2001 From: GregoryMorse Date: Mon, 23 Dec 2013 00:21:51 +0800 Subject: [PATCH 1/3] Update system.cpp Add native C++ support --- modules/core/src/system.cpp | 31 +++++++++++++++++++++++++++++++ 1 file changed, 31 insertions(+) diff --git a/modules/core/src/system.cpp b/modules/core/src/system.cpp index b301d95db..09daceed5 100644 --- a/modules/core/src/system.cpp +++ b/modules/core/src/system.cpp @@ -87,10 +87,41 @@ #ifdef HAVE_WINRT #include +#ifndef __cplusplus_winrt +#include +#pragma comment(lib, "runtimeobject.lib") +#endif std::wstring GetTempPathWinRT() { +#ifdef __cplusplus_winrt return std::wstring(Windows::Storage::ApplicationData::Current->TemporaryFolder->Path->Data()); +#else + Microsoft::WRL::ComPtr appdataFactory; + Microsoft::WRL::ComPtr appdataRef; + Microsoft::WRL::ComPtr storagefolderRef; + Microsoft::WRL::ComPtr storageitemRef; + HSTRING str; + HSTRING_HEADER hstrHead; + std::wstring wstr; + if (FAILED(WindowsCreateStringReference(RuntimeClass_Windows_Storage_ApplicationData, + (UINT32)wcslen(RuntimeClass_Windows_Storage_ApplicationData), &hstrHead, &str))) + return wstr; + if (FAILED(RoGetActivationFactory(str, IID_PPV_ARGS(appdataFactory.ReleaseAndGetAddressOf())))) + return wstr; + if (FAILED(appdataFactory->get_Current(appdataRef.ReleaseAndGetAddressOf()))) + return wstr; + if (FAILED(appdataRef->get_TemporaryFolder(storagefolderRef.ReleaseAndGetAddressOf()))) + return wstr; + if (FAILED(storagefolderRef.As(&storageitemRef))) + return wstr; + str = NULL; + if (FAILED(storageitemRef->get_Path(&str))) + return wstr; + wstr = WindowsGetStringRawBuffer(str, NULL); + WindowsDeleteString(str); + return wstr; +#endif } std::wstring GetTempFileNameWinRT(std::wstring prefix) From 0ccc903647955d632b9a9091d8ad989a2cd9b038 Mon Sep 17 00:00:00 2001 From: Peng Xiao Date: Fri, 27 Dec 2013 11:54:08 +0800 Subject: [PATCH 2/3] fixed a buffer overrun of ocl canny the `map` buffer does not have the same size with CUDA and index starts at [1, 1] instead of [0, 0]. --- modules/ocl/src/opencl/imgproc_canny.cl | 16 ++++++++-------- 1 file changed, 8 insertions(+), 8 deletions(-) diff --git a/modules/ocl/src/opencl/imgproc_canny.cl b/modules/ocl/src/opencl/imgproc_canny.cl index 0a54f1468..2ddfdae5f 100644 --- a/modules/ocl/src/opencl/imgproc_canny.cl +++ b/modules/ocl/src/opencl/imgproc_canny.cl @@ -381,8 +381,8 @@ struct PtrStepSz { int step; int rows, cols; }; -inline int get(struct PtrStepSz data, int y, int x) { return *((__global int *)((__global char*)data.ptr + data.step * y + sizeof(int) * x)); } -inline void set(struct PtrStepSz data, int y, int x, int value) { *((__global int *)((__global char*)data.ptr + data.step * y + sizeof(int) * x)) = value; } +inline int get(struct PtrStepSz data, int y, int x) { return *((__global int *)((__global char*)data.ptr + data.step * (y + 1) + sizeof(int) * (x + 1))); } +inline void set(struct PtrStepSz data, int y, int x, int value) { *((__global int *)((__global char*)data.ptr + data.step * (y + 1) + sizeof(int) * (x + 1))) = value; } ////////////////////////////////////////////////////////////////////////////////////////// // do Hysteresis for pixel whose edge type is 1 @@ -494,7 +494,7 @@ edgesHysteresisLocal } } #else - struct PtrStepSz map = {((__global int *)((__global char*)map_ptr + map_offset)), map_step, rows, cols}; + struct PtrStepSz map = {((__global int *)((__global char*)map_ptr + map_offset)), map_step, rows + 1, cols + 1}; __local int smem[18][18]; @@ -507,13 +507,13 @@ edgesHysteresisLocal smem[threadIdx.y + 1][threadIdx.x + 1] = x < map.cols && y < map.rows ? get(map, y, x) : 0; if (threadIdx.y == 0) - smem[0][threadIdx.x + 1] = y > 0 ? get(map, y - 1, x) : 0; + smem[0][threadIdx.x + 1] = x < map.cols ? get(map, y - 1, x) : 0; if (threadIdx.y == blockDim.y - 1) smem[blockDim.y + 1][threadIdx.x + 1] = y + 1 < map.rows ? get(map, y + 1, x) : 0; if (threadIdx.x == 0) - smem[threadIdx.y + 1][0] = x > 0 ? get(map, y, x - 1) : 0; + smem[threadIdx.y + 1][0] = y < map.rows ? get(map, y, x - 1) : 0; if (threadIdx.x == blockDim.x - 1) - smem[threadIdx.y + 1][blockDim.x + 1] = x + 1 < map.cols ? get(map, y, x + 1) : 0; + smem[threadIdx.y + 1][blockDim.x + 1] = x + 1 < map.cols && y < map.rows ? get(map, y, x + 1) : 0; if (threadIdx.x == 0 && threadIdx.y == 0) smem[0][0] = y > 0 && x > 0 ? get(map, y - 1, x - 1) : 0; if (threadIdx.x == blockDim.x - 1 && threadIdx.y == 0) @@ -525,7 +525,7 @@ edgesHysteresisLocal barrier(CLK_LOCAL_MEM_FENCE); - if (x >= map.cols || y >= map.rows) + if (x >= cols || y >= rows) return; int n; @@ -576,7 +576,7 @@ edgesHysteresisLocal if (n > 0) { const int ind = atomic_inc(counter); - st[ind] = (ushort2)(x, y); + st[ind] = (ushort2)(x + 1, y + 1); } #endif } From d014cb8fb48982ffec87dad36a40a455896ca88f Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Fri, 27 Dec 2013 14:44:58 +0400 Subject: [PATCH 3/3] fixed warning [-Wempty-body] --- modules/ocl/src/gftt.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/modules/ocl/src/gftt.cpp b/modules/ocl/src/gftt.cpp index a82196d78..4f24d1358 100644 --- a/modules/ocl/src/gftt.cpp +++ b/modules/ocl/src/gftt.cpp @@ -208,7 +208,7 @@ void cv::ocl::GoodFeaturesToTrackDetector_OCL::operator ()(const oclMat& image, if(!use_cpu_sorter) { // round to 2^n unsigned int n=1; - for(n=1;n<(unsigned int)corner_array_size;n<<=1); + for(n=1;n<(unsigned int)corner_array_size;n<<=1) ; corner_array_size = (int)n; ensureSizeIsEnough(1, corner_array_size , CV_32FC2, tmpCorners_);