From b0606b0557e65487f9ec0408554500a008d90439 Mon Sep 17 00:00:00 2001
From: Marina Kolpakova <no@email>
Date: Tue, 10 Jul 2012 11:58:10 +0000
Subject: [PATCH] LBP classifer moved to ptr from DevMem2D

---
 modules/gpu/src/cuda/lbp.cu                |  2 +-
 modules/gpu/src/opencv2/gpu/device/lbp.hpp | 26 +++++++++++-----------
 2 files changed, 14 insertions(+), 14 deletions(-)

diff --git a/modules/gpu/src/cuda/lbp.cu b/modules/gpu/src/cuda/lbp.cu
index aa3cdb3b1..769430ed8 100644
--- a/modules/gpu/src/cuda/lbp.cu
+++ b/modules/gpu/src/cuda/lbp.cu
@@ -69,7 +69,7 @@ namespace cv { namespace gpu { namespace device
                     ClNode node = nodes[current_node];
 
                     uchar4 feature = features[node.featureIdx];
-                    int c = evaluator( (y + feature.y) * istep + x + feature.x , feature, integral, istep);
+                    int c = evaluator( (y + feature.y) * istep + x + feature.x , feature.w * istep, feature.z, integral, istep);
                     const int* subsetIdx = subsets + (current_node * subsetSize);
 
                     int idx =  (subsetIdx[c >> 5] & ( 1 << (c & 31))) ? current_leave : current_leave + 1;
diff --git a/modules/gpu/src/opencv2/gpu/device/lbp.hpp b/modules/gpu/src/opencv2/gpu/device/lbp.hpp
index 8508fa41b..b5ef36584 100644
--- a/modules/gpu/src/opencv2/gpu/device/lbp.hpp
+++ b/modules/gpu/src/opencv2/gpu/device/lbp.hpp
@@ -160,38 +160,38 @@ __device__ __forceinline__ T __atomicMin(T* address, T val)
         __device__ __forceinline__ LBP() {}
 
         //feature as uchar x, y - left top, z,w - right bottom
-        __device__ __forceinline__ int operator() (unsigned int y, uchar4 feature, const int* integral, int step) const
+        __device__ __forceinline__ int operator() (unsigned int y, int featurew, int featurez, const int* integral, int step) const
         {
-            int x_off = 2 * feature.z;
+            int x_off = 2 * featurez;
             int anchors[9];
 
             anchors[0]  = integral[y];
-            anchors[1]  = integral[y + feature.z];
+            anchors[1]  = integral[y + featurez];
             anchors[0] -= anchors[1];
             anchors[2]  = integral[y + x_off];
             anchors[1] -= anchors[2];
-            anchors[2] -= integral[y + feature.z + x_off];
-            y+=feature.w * step;
+            anchors[2] -= integral[y + featurez + x_off];
+            y += featurew;
 
             anchors[3]  = integral[y];
-            anchors[4]  = integral[y + feature.z];
+            anchors[4]  = integral[y + featurez];
             anchors[3] -= anchors[4];
             anchors[5]  = integral[y + x_off];
             anchors[4] -= anchors[5];
-            anchors[5] -= integral[y + feature.z + x_off];
+            anchors[5] -= integral[y + featurez + x_off];
 
             anchors[0] -= anchors[3];
             anchors[1] -= anchors[4];
             anchors[2] -= anchors[5];
             // 0 - 2 contains s0 - s2
 
-            y+=feature.w * step;
+            y += featurew;
             anchors[6]  = integral[y];
-            anchors[7]  = integral[y + feature.z];
+            anchors[7]  = integral[y + featurez];
             anchors[6] -= anchors[7];
             anchors[8]  = integral[y + x_off];
             anchors[7] -= anchors[8];
-            anchors[8] -= integral[y + x_off + feature.z];
+            anchors[8] -= integral[y + x_off + featurez];
 
             anchors[3] -= anchors[6];
             anchors[4] -= anchors[7];
@@ -210,13 +210,13 @@ __device__ __forceinline__ T __atomicMin(T* address, T val)
             response |= (~(anchors[5] >> 31)) & 16;
             response |= (~(anchors[3] >> 31)) & 1;
 
-            y+=feature.w * step;
+            y += featurew;
             anchors[0]  = integral[y];
-            anchors[1]  = integral[y + feature.z];
+            anchors[1]  = integral[y + featurez];
             anchors[0] -= anchors[1];
             anchors[2]  = integral[y + x_off];
             anchors[1] -= anchors[2];
-            anchors[2] -= integral[y + x_off + feature.z];
+            anchors[2] -= integral[y + x_off + featurez];
 
             anchors[6] -= anchors[0];
             anchors[7] -= anchors[1];