fixed zgemm offset bug; removed profiling from client

Author: guacamoleo

src/client/clfunc_common.hpp

@@ -246,7 +246,7 @@ class clblasFunc
         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 );

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

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);
   }

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];