use vectors

modules/core/src/matmul.cpp

@@ -782,7 +782,6 @@ static bool ocl_gemm( InputArray matA, InputArray matB, double alpha,
 {
     int depth = matA.depth(), cn = matA.channels();
     int type = CV_MAKETYPE(depth, cn);
-    const int block_size = 16;
 
     CV_Assert( type == matB.type() && (type == CV_32FC1 || type == CV_64FC1 || type == CV_32FC2 || type == CV_64FC2) );
 
@@ -808,14 +807,8 @@ static bool ocl_gemm( InputArray matA, InputArray matB, double alpha,
     CV_Assert( matB.type() == type && (!haveC || matC.type() == type) );
     CV_Assert( sizeA.width == sizeB.height && (!haveC || sizeC == sizeD) );
 
-    String opts = format("-D T=%s -D T1=%s -D cn=%d -D LOCAL_SIZE=%d %s %s",
-                          ocl::typeToStr(type), ocl::typeToStr(depth), cn, block_size,
-                          haveC ? "-D HAVE_C" : "",
-                          doubleSupport ? " -D DOUBLE_SUPPORT" : "");
-
-    ocl::Kernel k("gemm", cv::ocl::core::gemm_oclsrc, opts);
-    if (k.empty())
-        return false;
+    int max_wg_size = (int)dev.maxWorkGroupSize();
+    int block_size = (max_wg_size / (32*cn) < 32) ? (max_wg_size / (16*cn) < 16) ? (max_wg_size / (8*cn) < 8) ? 1 : 8 : 16 : 32;
 
     matD.create(sizeD, type);
 
@@ -832,24 +825,37 @@ static bool ocl_gemm( InputArray matA, InputArray matB, double alpha,
     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",
+                          ocl::typeToStr(type), ocl::typeToStr(depth), ocl::typeToStr(CV_MAKETYPE(depth, kercn)),
+                          cn, kercn, block_size,
+                          (sizeA.width % block_size !=0) ? "-D NO_MULT" : "",
+                          haveC ? "-D HAVE_C" : "",
+                          doubleSupport ? " -D DOUBLE_SUPPORT" : "");
+
+    ocl::Kernel k("gemm", cv::ocl::core::gemm_oclsrc, opts);
+    if (k.empty())
+        return false;
+
     if (depth == CV_64F)
         k.args(ocl::KernelArg::ReadOnlyNoSize(A),
-               ocl::KernelArg::ReadOnlyNoSize(B),
-               ocl::KernelArg::ReadWrite(D),
+               ocl::KernelArg::ReadOnlyNoSize(B, cn, kercn),
+               ocl::KernelArg::ReadWrite(D, cn, kercn),
                sizeA.width, alpha, beta);
     else
         k.args(ocl::KernelArg::ReadOnlyNoSize(A),
-               ocl::KernelArg::ReadOnlyNoSize(B),
-               ocl::KernelArg::ReadWrite(D),
+               ocl::KernelArg::ReadOnlyNoSize(B, cn, kercn),
+               ocl::KernelArg::ReadWrite(D, cn, kercn),
                sizeA.width, (float)alpha, (float)beta);
 
-    size_t globalsize[2] = { sizeD.width, sizeD.height};
+    size_t globalsize[2] = { sizeD.width * cn / kercn, sizeD.height};
     size_t localsize[2] = { block_size, block_size};
-    return k.run(2, globalsize, localsize, false);
+    return k.run(2, globalsize, block_size!=1 ? localsize : NULL, false);
 }
-
 #endif
-
 }
 
 void cv::gemm( InputArray matA, InputArray matB, double alpha,

modules/core/src/opencl/gemm.cl

@@ -13,20 +13,29 @@
 #endif
 #endif
 
-#define TSIZE (int)sizeof(T)
+#define TSIZE  (int)sizeof(T)
+#define WTSIZE (int)sizeof(WT)
 
 #define IND_A mad24(y, A_step, A_offset)
-#define STEP_A 1
-
-#define IND_B mad24(x, TSIZE, B_offset)
-#define STEP_B B_step / TSIZE
+#define IND_B mad24(x, WTSIZE, B_offset)
+#define STEP_B B_step / WTSIZE
 
 #if cn==2
+#if kercn==2
+#define MUL(i, 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)\
     {\
     sum.x += fma(a.x, b.x, - a.y * b.y);\
     sum.y += fma(a.x, b.y, a.y * b.x);\
+    sum.z += fma(a.x, b.z, - a.y * b.w);\
+    sum.w += fma(a.x, b.w, a.y * b.z);\
     }
+#endif
 #else
 #define MUL(i, a, b) sum = fma(a, b, sum);
 #endif
@@ -44,22 +53,44 @@ __kernel void gemm(__global const uchar * A_ptr, int A_step, int A_offset,
     int lidy = get_local_id(1);
 
     __global const T* A = (__global const T*)(A_ptr + IND_A);
-    __global const T* B = (__global const T*)(B_ptr + IND_B);
+    __global const WT* B = (__global const WT*)(B_ptr + IND_B);
 
-    T sum = (T)(0);
-    __local T a_local[LOCAL_SIZE*LOCAL_SIZE];
-    __local T b_local[LOCAL_SIZE*LOCAL_SIZE];
+    WT sum = (WT)(0);
 
-    for (int p = 0; p < (n+LOCAL_SIZE-1)/LOCAL_SIZE; ++p)
+#if LOCAL_SIZE == 1
+
+    if (x < D_cols && y < D_rows)
     {
-        a_local[mad24(lidy, LOCAL_SIZE, lidx)] = A[mad24(p, LOCAL_SIZE, lidx)];
-        b_local[mad24(lidy, LOCAL_SIZE, lidx)] = B[mad24(p, LOCAL_SIZE, lidy)*STEP_B];
+        for (int i = 0; i < n; ++i)
+            MUL(i, A[i], B[i*STEP_B]);
+#else
+
+    __local T  a_local[LOCAL_SIZE*LOCAL_SIZE];
+    __local WT b_local[LOCAL_SIZE*LOCAL_SIZE];
+
+    int reps;
+#if NO_MULT
+    reps = (n + LOCAL_SIZE-1)/LOCAL_SIZE;
+#else
+    reps = n/LOCAL_SIZE;
+#endif
+
+    for (int p = 0; p < reps; ++p)
+    {
+        if (p * LOCAL_SIZE + lidx < n && y < D_rows)
+            a_local[mad24(lidy, LOCAL_SIZE, lidx)] = A[mad24(p, LOCAL_SIZE, lidx)];
+        if (p * LOCAL_SIZE + lidy < n && x < D_cols)
+            b_local[mad24(lidy, LOCAL_SIZE, lidx)] = B[mad24(p, LOCAL_SIZE, lidy)*STEP_B];
 
         barrier(CLK_LOCAL_MEM_FENCE);
 
         if (x < D_cols && y < D_rows)
         {
-            for (int i = 0; i < LOCAL_SIZE && p * LOCAL_SIZE + i < n; ++i)
+            for (int i = 0; i < LOCAL_SIZE
+#if NO_MULT
+                && p * LOCAL_SIZE + i < n
+#endif
+                ; ++i)
                 MUL(i, a_local[mad24(lidy, LOCAL_SIZE, i)], b_local[mad24(i, LOCAL_SIZE, lidx)]);
         }
 
@@ -68,11 +99,12 @@ __kernel void gemm(__global const uchar * A_ptr, int A_step, int A_offset,
 
     if (x < D_cols && y < D_rows)
     {
-        __global T* D = (__global T*)(D_ptr + mad24(y, D_step, mad24(x, TSIZE, D_offset)));
+#endif
+        __global WT* D = (__global WT*)(D_ptr + mad24(y, D_step, mad24(x, WTSIZE, D_offset)));
 #if HAVE_C
         D[0] = mad(alpha, sum, D[0]*beta);
 #else
         D[0] = alpha * sum;
 #endif
     }
-}
+}
\ No newline at end of file

modules/ts/src/ocl_test.cpp

@@ -48,7 +48,7 @@ namespace ocl {
 
 using namespace cv;
 
-int test_loop_times = 10; // TODO Read from command line / environment
+int test_loop_times = 1; // TODO Read from command line / environment
 
 #ifdef HAVE_OPENCL