Some small fixes

modules/core/src/matmul.cpp

@@ -788,7 +788,7 @@ static bool ocl_gemm( InputArray matA, InputArray matB, double alpha,
     const ocl::Device & dev = ocl::Device::getDefault();
     bool doubleSupport = dev.doubleFPConfig() > 0;
 
-    if ((!doubleSupport && depth == CV_64F))
+    if (!doubleSupport && depth == CV_64F)
         return false;
 
     bool haveC = matC.kind() != cv::_InputArray::NONE;
@@ -804,7 +804,7 @@ static bool ocl_gemm( InputArray matA, InputArray matB, double alpha,
 
     Size sizeD(sizeB.width, sizeA.height);
 
-    CV_Assert( matB.type() == type && (!haveC || matC.type() == type) );
+    CV_Assert( !haveC || matC.type() == type );
     CV_Assert( sizeA.width == sizeB.height && (!haveC || sizeC == sizeD) );
 
     int max_wg_size = (int)dev.maxWorkGroupSize();
@@ -822,11 +822,8 @@ static bool ocl_gemm( InputArray matA, InputArray matB, double alpha,
 
     if (haveC)
         ctrans ? transpose(matC, D) : matC.copyTo(D);
-    else
-        D.setTo(Scalar::all(0));
 
     int vectorWidths[] = { 4, 4, 2, 2, 1, 4, cn, -1 };
-
     int kercn = ocl::checkOptimalVectorWidth(vectorWidths, B, D);
 
     String opts = format("-D T=%s -D T1=%s -D WT=%s -D cn=%d -D kercn=%d -D LOCAL_SIZE=%d %s %s %s",

modules/core/src/opencl/gemm.cl

@@ -22,13 +22,13 @@
 
 #if cn==2
 #if kercn==2
-#define MUL(i, a, b)\
+#define MUL(a, b)\
     {\
     sum.x += fma(a.x, b.x, - a.y * b.y);\
     sum.y += fma(a.x, b.y, a.y * b.x);\
     }
 #else
-#define MUL(i, a, b)\
+#define MUL(a, b)\
     {\
     sum.x += fma(a.x, b.x, - a.y * b.y);\
     sum.y += fma(a.x, b.y, a.y * b.x);\
@@ -37,7 +37,7 @@
     }
 #endif
 #else
-#define MUL(i, a, b) sum = fma(a, b, sum);
+#define MUL(a, b) sum = fma(a, b, sum);
 #endif
 
 
@@ -62,7 +62,7 @@ __kernel void gemm(__global const uchar * A_ptr, int A_step, int A_offset,
     if (x < D_cols && y < D_rows)
     {
         for (int i = 0; i < n; ++i)
-            MUL(i, A[i], B[i*STEP_B]);
+            MUL(A[i], B[i*STEP_B]);
 #else
 
     __local T  a_local[LOCAL_SIZE*LOCAL_SIZE];
@@ -86,14 +86,14 @@ __kernel void gemm(__global const uchar * A_ptr, int A_step, int A_offset,
 
         if (x < D_cols && y < D_rows)
         {
-            for (int i = 0; i < LOCAL_SIZE
 #if NO_MULT
-                && p * LOCAL_SIZE + i < n
+            int ie = min(LOCAL_SIZE, n - p * LOCAL_SIZE);
+            for (int i = 0; i < ie; ++i)
+#else
+            for (int i = 0; i < LOCAL_SIZE; ++i)
 #endif
-                ; ++i)
+                MUL(a_local[mad24(lidy, LOCAL_SIZE, i)], b_local[mad24(i, LOCAL_SIZE, lidx)]);
-                MUL(i, a_local[mad24(lidy, LOCAL_SIZE, i)], b_local[mad24(i, LOCAL_SIZE, lidx)]);
         }
-
         barrier(CLK_LOCAL_MEM_FENCE);
     }