Merge pull request #111 from guacamoleo/develop

enables apiCallCount for zgemm within client and an zgemm bug fix

Author: Timmy

src/client/clfunc_xgemm.hpp

@@ -1013,12 +1013,15 @@ void
 xGemm<cl_double>::
 xGemm_Function(bool flush, cl_uint apiCallCount )
 {
-	clblasDgemm(order_, buffer_.trans_a_, buffer_.trans_b_,
+  for (int i = 0; i < apiCallCount; i++)
+	{
+	  clblasDgemm(order_, buffer_.trans_a_, buffer_.trans_b_,
                      buffer_.m_, buffer_.n_, buffer_.k_, buffer_.alpha_,
                      buffer_.buf_a_, buffer_.offA_, buffer_.lda_,
                      buffer_.buf_b_, buffer_.offB_, buffer_.ldb_,
                      buffer_.beta_, buffer_.buf_c_, buffer_.offC_,
                      buffer_.ldc_, 1, &queue_, 0, NULL, &event_);
+  }
 	//flush==true if only the kernel time (library call) is timed
 	//flush==false if memory time is also timed
 	if (flush==true)
@@ -1032,12 +1035,15 @@ void
 xGemm<cl_float2>::
 xGemm_Function(bool flush, cl_uint apiCallCount )
 {
-	clblasCgemm(order_, buffer_.trans_a_, buffer_.trans_b_,
+  for (int i = 0; i < apiCallCount; i++)
+	{
+	  clblasCgemm(order_, buffer_.trans_a_, buffer_.trans_b_,
                      buffer_.m_, buffer_.n_, buffer_.k_, buffer_.alpha_,
                      buffer_.buf_a_, buffer_.offA_, buffer_.lda_,
                      buffer_.buf_b_, buffer_.offB_, buffer_.ldb_,
                      buffer_.beta_, buffer_.buf_c_, buffer_.offC_,
                      buffer_.ldc_, 1, &queue_, 0, NULL, &event_);
+  }
 	//flush==true if only the kernel time (library call) is timed
 	//flush==false if memory time is also timed
 	if (flush==true)
@@ -1051,12 +1057,15 @@ void
 xGemm<cl_double2>::
 xGemm_Function(bool flush, cl_uint apiCallCount )
 {
-	clblasZgemm(order_, buffer_.trans_a_, buffer_.trans_b_,
+  for (int i = 0; i < apiCallCount; i++)
+	{
+	  clblasZgemm(order_, buffer_.trans_a_, buffer_.trans_b_,
                      buffer_.m_, buffer_.n_, buffer_.k_, buffer_.alpha_,
                      buffer_.buf_a_, buffer_.offA_, buffer_.lda_,
                      buffer_.buf_b_, buffer_.offB_, buffer_.ldb_,
                      buffer_.beta_, buffer_.buf_c_, buffer_.offC_,
                      buffer_.ldc_, 1, &queue_, 0, NULL, &event_);
+  }
 	//flush==true if only the kernel time (library call) is timed
 	//flush==false if memory time is also timed
 	if (flush==true)
@@ -1070,15 +1079,15 @@ double
 xGemm<cl_float2>::
 gflops()
 {
-    return (8.0*buffer_.m_*buffer_.n_*buffer_.k_)/time_in_ns();
+    return (8.0*buffer_.m_*buffer_.n_*buffer_.k_)/(time_in_ns() / buffer_.apiCallCount);
 }
 
 template<>
 double
 xGemm<cl_double2>::
 gflops()
 {
-    return (8.0*buffer_.m_*buffer_.n_*buffer_.k_)/time_in_ns();
+    return (8.0*buffer_.m_*buffer_.n_*buffer_.k_)/(time_in_ns() / buffer_.apiCallCount);
 }
 
 template<>

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