[bugfix] Used global memmory for result in dot_cl kernel

Fixed kernel argument bug for dot_cl kernel

Signed-off-by: Debadri Samaddar <s.debadri@samsung.com>

diff --git a/nntrainer/layers/cl_layers/blas_kernels.cpp b/nntrainer/layers/cl_layers/blas_kernels.cpp
index f901179d5954a53f644e935c0eb43d14c708cae4..06b8754718f1dc1a7d9fc4b6522fbb4345d85109 100644 (file)
--- a/nntrainer/layers/cl_layers/blas_kernels.cpp
+++ b/nntrainer/layers/cl_layers/blas_kernels.cpp
@@ -18,9 +18,9 @@ namespace nntrainer {
 std::string sgemv_cl_kernel_ =
   R"(__kernel void sgemv_cl(const __global float* A, const __global float* X,
                       __global float* Y, unsigned int M, unsigned int N) {                                            
-        unsigned int i, j;
+        unsigned int i;
         i = get_global_id(0);                         
-        float y0 = Y[i] * 0.0f;
+        float y0 = 0.0f;
         for (unsigned int j = 0; j < M; j++)                         
             y0 += A[i + j * N] * X[j]; 
         Y[i] = y0;                            
@@ -28,10 +28,10 @@ std::string sgemv_cl_kernel_ =
     })";
 
 std::string dot_cl_kernel_ =
-  R"(__kernel void dot_cl(const __global float* A, const __global float* X, unsigned int K, float res) {
-        res = 0;
+  R"(__kernel void dot_cl(const __global float* A, const __global float* X, unsigned int K, __global float* res) {
+        *res = 0;
         for (unsigned int i = 0; i < K; i++){
-            res += A[i] * X[i];
+            *res += A[i] * X[i];
         }
     })";
 
@@ -41,8 +41,7 @@ std::string sgemm_cl_kernel_ =
         
         unsigned int m = get_global_id(0);
         for (unsigned int n = 0; n < N; ++n) {
-          float c = 0.0;
-          float c_old = C[m * ldc + n];
+          float c = 0.0f;
           for (unsigned int k = 0; k < K; ++k) {
             float a, b;
             a = A[m * lda + k];
@@ -159,6 +158,9 @@ float dot_cl(const float *matAdata, const float *vecXdata, unsigned int dim1,
 
     opencl::Buffer inputX(context.context_inst_, dim1_size, true, nullptr);
 
+    opencl::Buffer dotResult(context.context_inst_, sizeof(float), true,
+                             &cl_ret);
+
     result = inputA.WriteData(context.command_queue_inst_, matAdata);
     if (!result) {
       break;
@@ -184,7 +186,7 @@ float dot_cl(const float *matAdata, const float *vecXdata, unsigned int dim1,
       break;
     }
 
-    result = kernel_dot.SetKernelArguments(3, &cl_ret, sizeof(float));
+    result = kernel_dot.SetKernelArguments(3, &dotResult, sizeof(cl_mem));
     if (!result) {
       break;
     }
@@ -198,6 +200,11 @@ float dot_cl(const float *matAdata, const float *vecXdata, unsigned int dim1,
       break;
     }
 
+    result = dotResult.ReadData(context.command_queue_inst_, &cl_ret);
+    if (!result) {
+      break;
+    }
+
   } while (false);
 
   return cl_ret;

diff --git a/nntrainer/layers/cl_layers/blas_kernels.h b/nntrainer/layers/cl_layers/blas_kernels.h
index d44b369a158eb7ac05b1ce0d93afaf5854fccb02..e3f1f870d382de701ecef34ec95a9f26cd99144f 100644 (file)
--- a/nntrainer/layers/cl_layers/blas_kernels.h
+++ b/nntrainer/layers/cl_layers/blas_kernels.h
@@ -61,6 +61,9 @@ float dot_cl(const float *matAdata, const float *vecXdata, unsigned int dim1,
  * @param[in] M number of op(A)'s and C's row
  * @param[in] N number of op(B)'s and C's columns
  * @param[in] K number of op(A)'s and columns and op(B)'s rows
+ * @param[in] lda number of A's columns
+ * @param[in] ldb number of B's columns
+ * @param[in] ldc number of C's columns
  * @param[in] context RunLayerContext reference
  */
 void sgemm_cl(const float *A, const float *B, float *C, unsigned int M,