From 65b8a1cb3756ba908440efaa4b18eed4037e25c6 Mon Sep 17 00:00:00 2001 From: ElenaGvozdeva Date: Thu, 16 Oct 2014 10:24:44 +0300 Subject: [PATCH] Some small fixes --- modules/core/src/matmul.cpp | 7 ++----- modules/core/src/opencl/gemm.cl | 18 +++++++++--------- 2 files changed, 11 insertions(+), 14 deletions(-) diff --git a/modules/core/src/matmul.cpp b/modules/core/src/matmul.cpp index 6d2adc8690..fec2ecb8b8 100644 --- a/modules/core/src/matmul.cpp +++ b/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", diff --git a/modules/core/src/opencl/gemm.cl b/modules/core/src/opencl/gemm.cl index ddd18adaf2..0961628a49 100644 --- a/modules/core/src/opencl/gemm.cl +++ b/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(i, a_local[mad24(lidy, LOCAL_SIZE, i)], b_local[mad24(i, LOCAL_SIZE, lidx)]); + MUL(a_local[mad24(lidy, LOCAL_SIZE, i)], b_local[mad24(i, LOCAL_SIZE, lidx)]); } - barrier(CLK_LOCAL_MEM_FENCE); } -- GitLab