[clblas] 49/61: fixed zgemm offset bug; removed profiling from client

[Ghislain Vaillant]

This is an automated email from the git hooks/post-receive script.

ghisvail-guest pushed a commit to branch master
in repository clblas.

commit 03ae1873f75af45d847afe70ecb1b39a5e5d2d88
Author: David Tanner <guacamoleo at gmail.com>
Date:   Tue Jun 23 10:00:14 2015 -0500

    fixed zgemm offset bug; removed profiling from client
---
 src/client/clfunc_common.hpp                   |  2 +-
 src/client/clfunc_xgemm.hpp                    | 13 -------------
 src/library/blas/functor/hawaii.cc             |  5 +----
 src/library/blas/gens/clTemplates/zgemm_gcn.cl | 11 ++++++++---
 4 files changed, 10 insertions(+), 21 deletions(-)

diff --git a/src/client/clfunc_common.hpp b/src/client/clfunc_common.hpp
index 86f93c0..01363a9 100644
--- a/src/client/clfunc_common.hpp
+++ b/src/client/clfunc_common.hpp
@@ -246,7 +246,7 @@ public:
         props_[2] = 0;
         ctx_ = clCreateContext(props_, 1, &device_, NULL, NULL, &err);
         OPENCL_V_THROW(err, "creating context");
-        queue_ = clCreateCommandQueue(ctx_, device_, CL_QUEUE_PROFILING_ENABLE, &err);
+        queue_ = clCreateCommandQueue(ctx_, device_, 0, &err);
 
 
         timer_id = timer.getUniqueID( "clfunc", 0 );
diff --git a/src/client/clfunc_xgemm.hpp b/src/client/clfunc_xgemm.hpp
index 27fa1f2..b676f13 100644
--- a/src/client/clfunc_xgemm.hpp
+++ b/src/client/clfunc_xgemm.hpp
@@ -1065,19 +1065,6 @@ xGemm_Function(bool flush, cl_uint apiCallCount )
                      buffer_.buf_b_, buffer_.offB_, buffer_.ldb_,
                      buffer_.beta_, buffer_.buf_c_, buffer_.offC_,
                      buffer_.ldc_, 1, &queue_, 0, NULL, &event_);
-#if 0
-    // print kernel time
-    clFinish(queue_);
-    cl_ulong start, stop;
-    double time;
-    cl_int err;
-    err = clGetEventProfilingInfo( event_, CL_PROFILING_COMMAND_START, sizeof(start), &start, NULL );
-    if (err) printf("err = %i\n", err);
-    err = clGetEventProfilingInfo( event_, CL_PROFILING_COMMAND_END,   sizeof(stop),  &stop,  NULL );
-    if (err) printf("err = %i\n", err);
-    time = (stop - start) / 1000000.0; // milliseconds
-    printf("kernel %lu -> %lu = %.f ms\n", start, stop, time );
-#endif
   }
 	//flush==true if only the kernel time (library call) is timed
 	//flush==false if memory time is also timed
diff --git a/src/library/blas/functor/hawaii.cc b/src/library/blas/functor/hawaii.cc
index 0d2b2c0..94a23a5 100644
--- a/src/library/blas/functor/hawaii.cc
+++ b/src/library/blas/functor/hawaii.cc
@@ -168,16 +168,13 @@ clblasSgemmFunctor * FunctorSelectorHawaii::select_sgemm_specific(clblasSgemmFun
 clblasZgemmFunctor * FunctorSelectorHawaii::select_zgemm_specific(clblasZgemmFunctor::Args & args)
 {
 
-  //TODO: the logic below is complicated; Needs cleanup;
-  clblasZgemmFunctor * functor;
-
   if ( args.M%32==0
     && args.N%64==0
     && args.K%8==0
     && args.transA==clblasNoTrans
     && args.transB==clblasTrans
     && args.order==clblasColumnMajor) {
-    functor = clblasZgemmFunctorGCN::provide(args, "Hawaii");
+    return clblasZgemmFunctorGCN::provide(args, "Hawaii");
   } else {
     return this->clblasFunctorSelector::select_zgemm_specific(args);
   }
diff --git a/src/library/blas/gens/clTemplates/zgemm_gcn.cl b/src/library/blas/gens/clTemplates/zgemm_gcn.cl
index fc28caf..1769639 100644
--- a/src/library/blas/gens/clTemplates/zgemm_gcn.cl
+++ b/src/library/blas/gens/clTemplates/zgemm_gcn.cl
@@ -200,10 +200,15 @@ __kernel void KERNEL_NAME(DATA_TYPE_CHAR,TRANSPOSE_A,TRANSPOSE_B,MACRO_TILE_NUM_
   uint const lda,
   uint const ldb,
   uint const ldc,
-  uint const offA,
-  uint const offB,
-  uint const offC )
+  uint const offsetA,
+  uint const offsetB,
+  uint const offsetC )
 {
+  // apply offsets
+  A += offsetA;
+  B += offsetB;
+  C += offsetC;
+
   // registers
   DATA_TYPE_STR rC[MICRO_TILE_NUM_ROWS][MICRO_TILE_NUM_COLS]  = {0};
   DATA_TYPE_STR rA[MICRO_TILE_NUM_ROWS];

-- 
Alioth's /usr/local/bin/git-commit-notice on /srv/git.debian.org/git/debian-science/packages/clblas.git