     enum { BIN_OP_OR, BIN_OP_AND, BIN_OP_XOR };
 
     template <typename T, int opid>
-    struct BinOp { __device__ T operator()(T lhs, T rhs); };
-
-    template <typename T>
-    struct BinOp<T, BIN_OP_OR>{ __device__ T operator()(T lhs, T rhs) { return lhs | rhs; } };
-
-    template <typename T>
-    struct BinOp<T, BIN_OP_AND>{ __device__ T operator()(T lhs, T rhs) { return lhs & rhs; } };
-
-    template <typename T>
-    struct BinOp<T, BIN_OP_XOR>{ __device__ T operator()(T lhs, T rhs) { return lhs ^ rhs; } };
+    struct BinOp;
+
+    template <typename T>
+    struct BinOp<T, BIN_OP_OR>
+    { 
+        static __device__ T call(T lhs, T rhs) 
+        { 
+            return lhs | rhs; 
+        } 
+    };
+
+    template <typename T>
+    struct BinOp<T, BIN_OP_AND>
+    { 
+        static __device__ T call(T lhs, T rhs) 
+        { 
+            return lhs & rhs; 
+        } 
+    };
+
+    template <typename T>
+    struct BinOp<T, BIN_OP_XOR>
+    { 
+        static __device__ T call(T lhs, T rhs) 
+        { 
+            return lhs ^ rhs; 
+        } 
+    };
 
     template <typename T, int cn, typename BinOp, typename Mask>
-    __global__ void bitwise_bin_op(int rows, int cols, const PtrStep src1, const PtrStep src2, PtrStep dst, BinOp op, Mask mask)
+    __global__ void bitwise_bin_op(int rows, int cols, const PtrStep src1, const PtrStep src2, PtrStep dst, Mask mask)
     {
         const int x = blockDim.x * blockIdx.x + threadIdx.x;
         const int y = blockDim.y * blockIdx.y + threadIdx.y;


            #pragma unroll
            for (int i = 0; i < cn; ++i)
                dsty[cn * x + i] = op(src1y[cn * x + i], src2y[cn * x + i]);
        }
    }



            #pragma unroll
            for (int i = 0; i < cn; ++i)
                dsty[cn * x + i] = UnOp::call(srcy[cn * x + i]);
        }
    }

