From aeaf1a6f6d724552f3b7dd4c5dbcd1d0652a6e78 Mon Sep 17 00:00:00 2001 From: "marina.kolpakova" Date: Mon, 20 Aug 2012 02:26:23 +0400 Subject: [PATCH] refactoring in Emulation --- modules/gpu/src/cuda/ccomponetns.cu | 4 ++-- modules/gpu/src/cuda/lbp.cu | 2 +- modules/gpu/src/opencv2/gpu/device/emulation.hpp | 4 ++-- 3 files changed, 5 insertions(+), 5 deletions(-) diff --git a/modules/gpu/src/cuda/ccomponetns.cu b/modules/gpu/src/cuda/ccomponetns.cu index 275066a..07f2410 100644 --- a/modules/gpu/src/cuda/ccomponetns.cu +++ b/modules/gpu/src/cuda/ccomponetns.cu @@ -316,7 +316,7 @@ namespace cv { namespace gpu { namespace device } } - changed = Emulation::sycthOr(changed); + changed = Emulation::syncthreadsOr(changed); if (!changed) break; @@ -474,7 +474,7 @@ namespace cv { namespace gpu { namespace device } } } - } while (Emulation::sycthOr(changed)); + } while (Emulation::syncthreadsOr(changed)); } __global__ void flatten(const DevMem2D edges, DevMem2Di comps) diff --git a/modules/gpu/src/cuda/lbp.cu b/modules/gpu/src/cuda/lbp.cu index e96692c..edfbf6e 100644 --- a/modules/gpu/src/cuda/lbp.cu +++ b/modules/gpu/src/cuda/lbp.cu @@ -279,7 +279,7 @@ namespace cv { namespace gpu { namespace device rect.z = __float2int_rn(windowW * scale); rect.w = __float2int_rn(windowH * scale); - int res = Emulation::smem::atomicInc(classified, (unsigned int)objects.cols); + int res = atomicInc(classified, (unsigned int)objects.cols); objects(0, res) = rect; } } diff --git a/modules/gpu/src/opencv2/gpu/device/emulation.hpp b/modules/gpu/src/opencv2/gpu/device/emulation.hpp index 1a6f579..074e911 100644 --- a/modules/gpu/src/opencv2/gpu/device/emulation.hpp +++ b/modules/gpu/src/opencv2/gpu/device/emulation.hpp @@ -51,11 +51,11 @@ namespace cv { namespace gpu { namespace device struct Emulation { - static __device__ __forceinline__ int sycthOr(int pred) + static __device__ __forceinline__ int syncthreadsOr(int pred) { #if defined (__CUDA_ARCH__) && (__CUDA_ARCH__ < 200) // just campilation stab - return false; + return 0; #else return __syncthreads_or(pred); #endif -- 2.7.4