[clblas] 37/61: better handle sgemm NT where M and N are mod32 and not mod64. M and N are within range from 1184 to 3872
Ghislain Vaillant
ghisvail-guest at moszumanska.debian.org
Fri Jul 24 22:49:46 UTC 2015
This is an automated email from the git hooks/post-receive script.
ghisvail-guest pushed a commit to branch master
in repository clblas.
commit f4af83823419e8d371db2cbca8cd01c15676a1d6
Author: Timmy <timmy.liu at amd.com>
Date: Fri May 29 13:09:59 2015 -0500
better handle sgemm NT where M and N are mod32 and not mod64. M and N are within range from 1184 to 3872
---
src/library/CMakeLists.txt | 3 +
src/library/bingen.cmake | 1 +
src/library/blas/functor/hawaii.cc | 20 +-
src/library/blas/functor/hawaii_sgemmSplit64_32.cc | 423 ++++++++++++++++
.../blas/functor/hawaii_sgemmSplitKernel.cc | 5 +-
.../blas/functor/include/hawaii_sgemmSplit64_32.h | 46 ++
.../gens/clTemplates/sgemm_hawaiiSplit64_32.cl | 530 +++++++++++++++++++++
7 files changed, 1026 insertions(+), 2 deletions(-)
diff --git a/src/library/CMakeLists.txt b/src/library/CMakeLists.txt
index 3378d19..b5f5d66 100644
--- a/src/library/CMakeLists.txt
+++ b/src/library/CMakeLists.txt
@@ -73,6 +73,7 @@ set(SRC_BLAS
blas/functor/hawaii_dgemmChannelConflict.cc
blas/functor/hawaii_dgemmSplitKernel.cc
blas/functor/hawaii_sgemmSplitKernel.cc
+ blas/functor/hawaii_sgemmSplit64_32.cc
blas/functor/gcn_dgemmCommon.cc
blas/functor/gcn_sgemm.cc
blas/functor/gcn_dgemmSmallMatrices.cc
@@ -103,6 +104,7 @@ set(SRC_BLAS_HEADERS
blas/functor/include/hawaii_dgemmChannelConflict.h
blas/functor/include/hawaii_dgemmSplitKernel.h
blas/functor/include/hawaii_sgemmSplitKernel.h
+ blas/functor/include/hawaii_sgemmSplit64_32.h
blas/functor/include/gcn_dgemmCommon.h
blas/functor/include/gcn_sgemm.h
blas/functor/include/gcn_dgemmSmallMatrices.h
@@ -248,6 +250,7 @@ set(SRC_CL_TEMPLATES_GEN
sgemm_gcn.clHawaii_64.bin.cl
sgemm_gcn.clBonaire_64.bin.cl
sgemm_gcn.clTahiti_64.bin.cl
+ sgemm_hawaiiSplit64_32.clHawaii_64.bin.cl
)
set(SRC_BLAS_GENERIC_HEADERS
diff --git a/src/library/bingen.cmake b/src/library/bingen.cmake
index 6a3e778..bb6515d 100644
--- a/src/library/bingen.cmake
+++ b/src/library/bingen.cmake
@@ -14,6 +14,7 @@ ${CLTEMPLATE_PATH}/dgemm_hawaiiSplitKernel.cl
${CLTEMPLATE_PATH}/sgemm_hawaiiSplitKernel.cl
${CLTEMPLATE_PATH}/sgemm_gcn.cl
${CLTEMPLATE_PATH}/sgemm_gcn_SmallMatrices.cl
+${CLTEMPLATE_PATH}/sgemm_hawaiiSplit64_32.cl
)
diff --git a/src/library/blas/functor/hawaii.cc b/src/library/blas/functor/hawaii.cc
index c91e64e..643ba9e 100644
--- a/src/library/blas/functor/hawaii.cc
+++ b/src/library/blas/functor/hawaii.cc
@@ -24,6 +24,7 @@
#include "gcn_dgemmSmallMatrices.h"
#include "gcn_sgemmSmallMatrices.h"
#include "hawaii_sgemmBranchKernel.h"
+#include "hawaii_sgemmSplit64_32.h"
FunctorSelectorHawaii FunctorSelectorHawaii::instance ;
@@ -106,8 +107,24 @@ clblasSgemmFunctor * FunctorSelectorHawaii::select_sgemm_specific(clblasSgemmFun
SmallMatricesMod32 = SmallMatricesMod32&&Not_TT&&args.K % 16 == 0;
//SmallMatrices= false;
- bool useSpliKernel=((args.M%96==0 && args.N%96==0) || !(args.M%64==0 && args.N%64==0&& args.M<4000 &&args.N<4000)) /*&&args.K%16==0*/;
+ bool useSpliKernel=((args.M%96==0 && args.N%96==0) || !(args.M%64==0 && args.N%64==0&& args.M<4000 &&args.N<4000)) ;
useSpliKernel=useSpliKernel&&Not_TT;
+
+ //functor = clBlashawaiiSgemmSplit64_32Functor::provide(args, "Hawaii");
+ //if (functor)
+ // return functor;
+
+ if ((args.M >= 1184 && args.N >= 1184) && (args.M <= 3872 && args.N <= 3872) && (args.M % 64 != 0 && args.N % 64 != 0) && (args.M % 96 != 0 && args.N % 96 != 0) && (args.K % 16 == 0))
+ {
+ //all the mod32 sizes that is not mod64 or mod96 ranging from 1184 to 3872
+ //non mod32 cases are not implemented in this approach and are of less interest
+ if ((args.M % 32 == 0 && args.N % 32 == 0) && (args.transA == clblasNoTrans && args.transB == clblasTrans))
+ {
+ functor = clBlashawaiiSgemmSplit64_32Functor::provide(args, "Hawaii");
+ if (functor)
+ return functor;
+ }
+ }
//the English translation of below is: if small matrix that is (not mod32) and (not_TT) and K has to be mod 16
if (SmallMatrices && (!SmallMatricesMod32) && (Not_TT) && (args.K%16 == 0))
@@ -141,6 +158,7 @@ clblasSgemmFunctor * FunctorSelectorHawaii::select_sgemm_specific(clblasSgemmFun
// else use the fallback implementation
return this->clblasFunctorSelector::select_sgemm_specific(args);
+
#endif
}
diff --git a/src/library/blas/functor/hawaii_sgemmSplit64_32.cc b/src/library/blas/functor/hawaii_sgemmSplit64_32.cc
new file mode 100644
index 0000000..a0766ce
--- /dev/null
+++ b/src/library/blas/functor/hawaii_sgemmSplit64_32.cc
@@ -0,0 +1,423 @@
+#if !defined CLBLAS_HAWAII_DYNAMIC_KERNEL || !defined CLBLAS_BONAIRE_DYNAMIC_KERNEL
+//this split kernel algorithm solves the main matrix with 64x64 micro tile size
+//solves the row boundry with 32x64 micro tile size
+//solves the column boundry with 64x32 micro tile size
+//solves the rest boundry with 32x32 micro tile size
+//assumption : after the main matrix being computed by kernels with 64x64 micro tile size, the boundary are of size 32.
+//in other words, M and N are mod32 and not mod64
+#include <stdio.h>
+#include <string.h>
+#include <clBLAS.h>
+
+#include <devinfo.h>
+#include "clblas-internal.h"
+#include "solution_seq.h"
+
+#include <functor.h>
+#include <binary_lookup.h>
+#include <iostream>
+
+#include <functor_xgemm.h>
+#include <tahiti.h>
+#include <hawaii.h>
+
+#include "BinaryBuild.h"
+#include "hawaii_sgemmSplit64_32.h"
+
+
+#if BUILD_KERNEL_FROM_STRING
+//#include "sgemm_hawaiiSplitKernel.clT"
+#else
+
+#ifndef CLBLAS_HAWAII_DYNAMIC_KERNEL
+#include "sgemm_hawaiiSplit64_32.clHawaii_64.bin.clT"
+#include "sgemm_gcn.clHawaii_64.bin.clT"
+#endif//CLBLAS_HAWAII_DYNAMIC_KERNEL
+
+#ifndef CLBLAS_BONAIRE_DYNAMIC_KERNEL
+//#include "sgemm_hawaiiSplitKernel.clBonaire_64.bin.clT"
+#endif //CLBLAS_BONAIRE_DYNAMIC_KERNEL
+
+#endif //BUILD_KERNEL_FROM_STRING
+
+// Just because the full name is too long
+typedef clBlashawaiiSgemmSplit64_32Functor::Variant Variant;
+
+//define the string name of the soure/binary code
+#define SGEMM_SRC_NAME(TA,TB, DIVK, MULT) sgemm_##TA##TB##_##DIVK##_SPLIT##MULT
+#define SGEMM_SRC_NAME_HAWAII(TA,TB, DIVK, MULT, BITS) sgemm_##TA##TB##_##DIVK##_SPLIT##MULT##_##BITS##_bin_Hawaii
+#define SGEMM_SRC_NAME_BONAIRE(TA,TB, DIVK, MULT, BITS) sgemm_##TA##TB##_##DIVK##_SPLIT##MULT##_##BITS##_bin_Bonaire
+#define SGEMM_SRC_NAME_BIN(TA,TB, DIVK, MULT, BITS, DEVICE) SGEMM_SRC_NAME##_##DEVICE(TA,TB, DIVK, MULT, BITS)
+
+
+//variant name used to differentiate the different ones
+#define SGEMM_VARIANT_NAME(TA,TB, DIVK, MULT) "sgemm_" #TA #TB "_" #DIVK "_SPLIT64_32" #MULT
+//SGEMM_VARIANT_NAME(TA, TB, DIVM , DIVN, DIVK, GREATER48M, GREATER48N, NBKERNEL),
+
+#define SGEMM_KERNEL_NAME(TA,TB,DIVM,DIVN,DIVK,BS0,BS1,NV0,NV1,MULT, BLOC) "sgemm_" #TA #TB "_" #DIVM "_" #DIVN "_" #DIVK "_" #BS0 "x" #BS1 "_" #NV0 "x" #NV1 #MULT "_SPLIT_" #BLOC
+
+
+#define trans_N clblasNoTrans
+#define trans_T clblasTrans
+
+// Fill a variant descriptor using OpenCL source
+#define SGEMM_VARIANT_OBJ(TA,TB,DIVK,BS0,BS1,NV0,NV1, BITS, MULT, \
+ KERNEL_NAME_MAIN, KERNEL_NAME_ROW, KERNEL_NAME_COLUMN, KERNEL_NAME_SINGLE, \
+ KERNELS_SRC, \
+ KERNEL_BUILD_OPTIONS, \
+ KERNELS_BIN, \
+ KERNEL_BIN_SIZE) { \
+ SGEMM_VARIANT_NAME(TA,TB, DIVK, MULT), \
+{ KERNEL_NAME_MAIN, KERNEL_NAME_ROW, KERNEL_NAME_COLUMN, KERNEL_NAME_SINGLE } , \
+ KERNELS_SRC, \
+ KERNEL_BUILD_OPTIONS, \
+ KERNELS_BIN, \
+ KERNEL_BIN_SIZE, \
+ trans_##TA, trans_##TB, \
+ DIVK , \
+{ BS0, BS1 } , \
+{ NV0, NV1 } , \
+#MULT \
+}
+
+typedef clblasFunctorCache<clBlashawaiiSgemmSplit64_32Functor, const Variant *> CacheSplit;
+static CacheSplit cachesplit ;
+
+// Make it 1 to enable additional debug 'print'
+#define VERB 0
+
+
+
+//static bool applicable( const Variant & var, clblasSgemmFunctor::Args & args, int RefMultiple )
+//{
+//#if 0
+// // Transpose values are tested in select_variant
+// if ( args.transA != var.transA ) return false ;
+// if ( args.transB != var.transB ) return false ;
+//#endif
+//
+// //if (args.N>=var.divN && args.N % var.divN != 0 )
+// if ( args.N % var.divN != 0 )
+// return false ;
+// if ( args.M % var.divM != 0 )
+// return false ;
+// if(var.Greater[0]?args.M<RefMultiple:args.M>=RefMultiple)
+// return false;
+// if(var.Greater[1]?args.N<RefMultiple:args.N>=RefMultiple)
+// return false;
+// if ( args.beta==0 && var.mult.compare("__ALPHA")!=0)
+// return false ;
+// return true ;
+//}
+
+static void to_upper(char* input)
+{
+ while(*input)
+ {
+ *input=toupper(*input);
+ input++;
+ }
+}
+
+
+static const Variant * select_variant_SplitKernel( clblasSgemmFunctor::Args & args, const char* DevName, cl_uint _64BitsUse )
+{
+ if(_64BitsUse!=64)
+ {
+ std::cout<<"we don't support clblas on 32 bits"<< std::endl;
+ assert(1);
+ return NULL;
+ }
+
+
+ if ( args.transA == clblasNoTrans )
+ {
+ if ( args.transB == clblasNoTrans )
+ {
+
+ // ===== sgemm NN ======
+ // NN not implemented yet
+ return NULL;
+
+
+ }
+ if (args.transB == clblasTrans)
+ {
+
+ const char* KName_NTMain = "sgemm_NT_64_64_16_16x16_4x4__ALPHABETA_SPLIT_MAIN" ;
+ const char* KName_NTRow = "sgemm_NT_32_64_16_16x16_2x4__ALPHABETA_SPLIT_ROW" ;
+ const char* KName_NTColumn = "sgemm_NT_64_32_16_16x16_4x2__ALPHABETA_SPLIT_COLUMN" ;
+ const char* KName_NTSingleWave = "sgemm_NT_32_32_16_16x16_2x2__ALPHABETA_SPLIT_SINGLE" ;
+
+
+
+ const char* KBin_NTMain64 ;
+ size_t KBin_NTMainSize64 = 0;
+
+
+
+
+ if (!strcmp(DevName, "Hawaii"))
+ {
+#ifndef CLBLAS_HAWAII_DYNAMIC_KERNEL
+ //KBin_NTMain64 = SGEMM_SRC_NAME_BIN(N, T, 16, __ALPHABETA, 64, HAWAII) ;
+ //KBin_NTMainSize64 = sizeof(SGEMM_SRC_NAME_BIN(N, T, 16, __ALPHABETA, 64, HAWAII)) ;
+ KBin_NTMain64 = sgemm_NT_64_32_SPLIT__ALPHABETA_64_bin_Hawaii;
+ KBin_NTMainSize64 = sizeof(sgemm_NT_64_32_SPLIT__ALPHABETA_64_bin_Hawaii);
+
+#endif //CLBLAS_HAWAII_DYNAMIC_KERNEL
+ }
+
+ else if (!strcmp(DevName, "Bonaire"))
+ {
+#ifndef CLBLAS_BONAIRE_DYNAMIC_KERNEL
+ //not implemented for Bonaire yet
+#endif //#ifndef CLBLAS_BONAIRE_DYNAMIC_KERNEL
+ }
+
+
+ // ===== SGEMM NT ======
+
+
+ static const Variant variant = SGEMM_VARIANT_OBJ(N,T,16,16,16,4,4,64,__ALPHABETA,
+ KName_NTMain,KName_NTRow, KName_NTColumn, KName_NTSingleWave ,
+ NULL,
+ NULL,
+ KBin_NTMain64,
+ KBin_NTMainSize64) ;
+
+ return &variant ;
+
+ }
+ }
+ else
+ {
+ // TN and TT are not implemented yet
+ return NULL;
+ }
+
+ return NULL;
+}
+
+clBlashawaiiSgemmSplit64_32Functor::clBlashawaiiSgemmSplit64_32Functor(Args & args, const Variant * variant, cl_int & err)
+{
+
+ cl_device_id device;
+ cl_context context;
+ m_program=NULL;
+ m_variantSplit = variant;
+
+
+ cl_command_queue queue = args.queue;
+ err = getDeviceAndContext(queue, device, context);
+ if( err != CL_SUCCESS )
+ {
+ return;
+ }
+
+ if (VERB) printf(" ===> GET KERNEL %s\n", this->m_variantSplit->variantName) ;
+
+ //Ben do I use the correct "kernel_name"?
+ BinaryLookup bl(context, device, "clBlashawaiiSgemmSplitKernelFunctor");
+
+ bl.variantRaw( this->m_variantSplit->variantName, strlen(this->m_variantSplit->variantName)+1 ) ;
+
+ if ( !bl.found() ) // may create empty file or may wait until file is ready
+ {
+ if ( this->m_variantSplit->bin != NULL )
+ {
+ // build from a pre-compiled version of the kernel (SPIR or cl binaries)
+ //only 1 binary containing all the kernel
+ err = bl.buildFromBinary(this->m_variantSplit->bin, this->m_variantSplit->bin_size, /*this->m_variantSplit->build_options[i]*/ "-cl-std=2.0");
+ }
+ else
+ {
+ //// directly build from a char*
+ //for (int i=0; i<4; i++)
+ // if(this->m_variantSplit->source[i] != 0)
+ // err = bl.buildFromSource(this->m_variantSplit->source[i]);
+ if (VERB) printf(" ===> BUILD PROBLEM WE DON'T SUPPORT SOURCE BUILD FOR SPLIT SGEMM\n") ;
+ return;
+ }
+
+ if ( err != CL_SUCCESS )
+ {
+ if (VERB) printf(" ===> BUILD PROBLEM\n") ;
+
+ return;
+ }
+ }
+
+ this->m_program = bl.getProgram();
+}
+
+
+
+clBlashawaiiSgemmSplit64_32Functor *
+clBlashawaiiSgemmSplit64_32Functor::provide(clblasSgemmFunctor::Args & args, char* DevName)
+{
+
+ if ( args.order == clblasRowMajor )
+ return NULL ; // The RowMajor case shall never occur.
+
+ cl_device_id dev;
+ cl_context ctxt;
+
+ cl_int err = getDeviceAndContext(args.queue, dev, ctxt);
+ if (err != CL_SUCCESS)
+ {
+ return NULL;
+ }
+ cl_uint bitness = getAddressBits(dev);
+
+ int major;
+ int minor;
+
+ getCLVersion(dev, major, minor);
+
+ //if (major<2)
+ // return NULL;
+
+ // to_upper( DevName);
+ const Variant * variant = select_variant_SplitKernel( args, DevName, bitness ) ;
+ if ( variant == NULL )
+ return NULL ;
+
+
+
+ CacheSplit::Lookup lookup(cachesplit, ctxt, dev, variant) ;
+
+
+ if ( lookup.ok() )
+ {
+ clBlashawaiiSgemmSplit64_32Functor * functor = lookup.get();
+ functor->retain(); // increment the reference counter to avoid deletion while it is still beeing used
+ return functor;
+ }
+
+ clBlashawaiiSgemmSplit64_32Functor * functor = new clBlashawaiiSgemmSplit64_32Functor(args, variant, err);
+ if (err != CL_SUCCESS)
+ {
+ return NULL;
+ }
+
+ lookup.set(functor) ;
+
+ return functor;
+
+}
+
+
+cl_int clBlashawaiiSgemmSplit64_32Functor::KernelsLaunch(cl_command_queue queue, cl_kernel Kernel[4], Args &args)
+{
+ //GlobalX = ((Mvalue - 1) / 64) * 16
+ //GlobalY = ((Nvalue - 1) / 64) * 16
+ size_t GlobalX = ((args.M - 1) / (m_variantSplit->bwi[0] * m_variantSplit->ls[0])) * 16;
+ size_t GlobalY = ((args.N - 1) / (m_variantSplit->bwi[1] * m_variantSplit->ls[1])) * 16;
+
+
+ std::size_t gs[2] = {GlobalX, GlobalY};
+ cl_int error = 0;
+ //M and N are not mod64 and are mod32
+ if (args.M % 64 != 0 && args.N % 64 != 0 && args.M % 32 == 0 && args.N % 32 == 0 && args.M >= 64 && args.N >= 64)
+ {
+ if (VERB) printf(" ===> EXECUTE KERNEL 0, 1, 2, 3 \n") ;
+ error = clEnqueueNDRangeKernel(queue, Kernel[0], 2, NULL, gs, m_variantSplit->ls, args.numEventsInWaitList, args.eventWaitList,NULL);
+
+ gs[0] = 16;
+ error |= clEnqueueNDRangeKernel(queue, Kernel[1], 2, NULL, gs, m_variantSplit->ls, 0, NULL,NULL);
+
+ gs[1] = 16;
+ gs[0] = GlobalX;
+ error |= clEnqueueNDRangeKernel(queue, Kernel[2], 2, NULL, gs, m_variantSplit->ls, 0, NULL,NULL);
+
+ gs[0] = 16; gs[1] = 16;
+ error |= clEnqueueNDRangeKernel(queue, Kernel[3], 2, NULL, gs, m_variantSplit->ls, 0, NULL,args.events);
+ return error;
+ }
+
+
+ return clblasNotImplemented;
+}
+
+
+
+clblasStatus clBlashawaiiSgemmSplit64_32Functor::execute(Args &args)
+{
+ cl_int err;
+ cl_command_queue queue = args.queue;
+
+ if (VERB) printf(" ===> EXECUTE KERNEL %s, alpha =%f ,beta = %f\n", this->m_variantSplit->kernel_name, args.alpha, args.beta) ;
+
+ cl_kernel kernel[4];
+ int NBKernel = 0;
+
+ for (int i=0; i<4; i++)
+ {
+ if (this->m_variantSplit->kernel_name[i])
+ {
+ kernel[i ]= clCreateKernel( this->m_program, this->m_variantSplit->kernel_name[i], &err);
+ if (err != CL_SUCCESS)
+ return clblasStatus(err) ;
+ NBKernel++;
+ }
+ else
+ break;
+ }
+
+ if (NBKernel != 4) return clblasStatus(clblasBuildProgramFailure) ;
+
+ if (VERB)
+ {
+ for (int i=0; i<NBKernel; i++)
+ printf(" ===> FOUND %s\n", this->m_variantSplit->kernel_name[i]) ;
+ }
+
+ int M = args.M, N = args.N, K = args.K;
+ int lda = args.lda, ldb = args.ldb, ldc = args.ldc;
+
+ int offsetA = args.offA;
+ int offsetB = args.offB;
+ int offsetC = args.offC;
+
+ int arg[4]={0, 0, 0, 0} ;
+
+ //// All sgemm kernels shall have the same arguments: (A,B,C,M,N,K,alpha,beta,lda,ldb,ldc,offa,offb,offc)
+
+ for (int i=0; i<NBKernel; i++)
+ {
+ setKernelArg<cl_mem>(kernel[i], arg[i]++, args.A);
+ setKernelArg<cl_mem>(kernel[i], arg[i]++, args.B);
+ setKernelArg<cl_mem>(kernel[i], arg[i]++, args.C);
+
+ setKernelArg<int>(kernel[i], arg[i]++, M);
+ setKernelArg<int>(kernel[i], arg[i]++, N);
+ setKernelArg<int>(kernel[i], arg[i]++, K);
+
+ setKernelArg<cl_float>(kernel[i], arg[i]++, args.alpha);
+ //if (args.beta!=0 && this->m_variantSplit->mult.compare("__ALPHA")!=0)
+ setKernelArg<cl_float>(kernel[i], arg[i]++, args.beta);
+
+ setKernelArg<int>(kernel[i], arg[i]++, lda);
+ setKernelArg<int>(kernel[i], arg[i]++, ldb);
+ setKernelArg<int>(kernel[i], arg[i]++, ldc);
+
+ setKernelArg<int>(kernel[i], arg[i]++, offsetA);
+ setKernelArg<int>(kernel[i], arg[i]++, offsetB);
+ setKernelArg<int>(kernel[i], arg[i]++, offsetC);
+ }
+
+ err = KernelsLaunch(queue, kernel, args);
+
+
+
+ for (int i = 0; i<NBKernel; i++)
+ clReleaseKernel(kernel[i]) ;
+
+ if (VERB) printf(" ===> ERR=%d \n",(int)err) ;
+
+ // err= clFinish(queue);
+ return clblasStatus(err) ;
+
+}
+#endif
diff --git a/src/library/blas/functor/hawaii_sgemmSplitKernel.cc b/src/library/blas/functor/hawaii_sgemmSplitKernel.cc
index 637ef0f..3930cf1 100644
--- a/src/library/blas/functor/hawaii_sgemmSplitKernel.cc
+++ b/src/library/blas/functor/hawaii_sgemmSplitKernel.cc
@@ -1,5 +1,8 @@
#if !defined CLBLAS_HAWAII_DYNAMIC_KERNEL || !defined CLBLAS_BONAIRE_DYNAMIC_KERNEL
-
+//this split kernel algorithm solves the main matrix with 96x96 micro tile size
+//solves the row boundry with 16x96 micro tile size
+//solves the column boundry with 96x16 micro tile size
+//solves the rest boundry with 16x16 micro tile size
#include <stdio.h>
#include <string.h>
#include <clBLAS.h>
diff --git a/src/library/blas/functor/include/hawaii_sgemmSplit64_32.h b/src/library/blas/functor/include/hawaii_sgemmSplit64_32.h
new file mode 100644
index 0000000..d1e3c47
--- /dev/null
+++ b/src/library/blas/functor/include/hawaii_sgemmSplit64_32.h
@@ -0,0 +1,46 @@
+#ifndef HAWAII_SGEMMMSPLIT64_32
+#define HAWAII_SGEMMMSPLIT64_32
+
+#include "gcn_sgemm.h"
+
+
+class clBlashawaiiSgemmSplit64_32Functor : public clblasSgemmFunctorGCN
+{
+public:
+ struct Variant
+ {
+ const char * variantName;
+ const char * kernel_name[4] ; //order is main, row, column, single
+ const char * source ; // the kernel source (shall be unique)
+ const char * build_options;
+ const char * bin ;
+ size_t bin_size ;
+ clblasTranspose transA ; //
+ clblasTranspose transB ; //
+ unsigned divK ; // Required divisor of N (use 1 when N can be of any value)
+ size_t ls[2] ; // Local size (the work-group size)
+ size_t bwi[2] ; // Block size work-item: Number of elements calculated by each work items
+ // So basically each kernel is computing a block of
+ // (ls[0]*bwi[0]) x (ls[1]*bwi[1])
+ // elements of C.
+ std::string mult;
+
+ } ;
+
+
+ private: // Constructor & Destructor
+
+ clBlashawaiiSgemmSplit64_32Functor(Args & args, const Variant * variant, cl_int & err);
+ cl_int KernelsLaunch(cl_command_queue queue, cl_kernel Kernel[4], Args &args);
+ const Variant * m_variantSplit ; // Pointer to a 'const static' object describing the kernel variant.
+
+public:
+
+ // Provide a suitable hawaii_sgemmChannelConflict for the specified args
+ // or NULL if none
+ static clBlashawaiiSgemmSplit64_32Functor * provide(clblasSgemmFunctor::Args & args, char* DevName);
+ virtual clblasStatus execute(Args &args) ;
+
+};
+
+#endif
\ No newline at end of file
diff --git a/src/library/blas/gens/clTemplates/sgemm_hawaiiSplit64_32.cl b/src/library/blas/gens/clTemplates/sgemm_hawaiiSplit64_32.cl
new file mode 100644
index 0000000..fb40834
--- /dev/null
+++ b/src/library/blas/gens/clTemplates/sgemm_hawaiiSplit64_32.cl
@@ -0,0 +1,530 @@
+//below kernels work with an assumption: after the main matrix being computed by kernels with 64x64 micro tile size, the boundary are of size 32.
+//Thus, M and N are of mod32 and not necessarily of mod64.
+//use sgemm_NT_64_64_16_16x16_4x4__ALPHABETA() from sgemm_gcn.cl as the main solver
+
+static const char * sgemm_NT_64_32_SPLIT__ALPHABETA = "
+
+#define M4x4 \
+ rA[0][0] = lA[offA + 0]; \
+ rA[0][1] = lA[offA + 16]; \
+ rA[0][2] = lA[offA + 32]; \
+ rA[0][3] = lA[offA + 48]; \
+ rB[0][0] = lB[offB + 0]; \
+ rB[0][1] = lB[offB + 16]; \
+ rB[0][2] = lB[offB + 32]; \
+ rB[0][3] = lB[offB + 48]; \
+ offA += 65; \
+ offB += 65; \
+ rC[0][0]=mad(rA[0][0],rB[0][0],rC[0][0]); \
+ rC[1][0]=mad(rA[0][1],rB[0][0],rC[1][0]); \
+ rC[2][0]=mad(rA[0][2],rB[0][0],rC[2][0]); \
+ rC[3][0]=mad(rA[0][3],rB[0][0],rC[3][0]); \
+ rC[0][1]=mad(rA[0][0],rB[0][1],rC[0][1]); \
+ rC[1][1]=mad(rA[0][1],rB[0][1],rC[1][1]); \
+ rC[2][1]=mad(rA[0][2],rB[0][1],rC[2][1]); \
+ rC[3][1]=mad(rA[0][3],rB[0][1],rC[3][1]); \
+ rC[0][2]=mad(rA[0][0],rB[0][2],rC[0][2]); \
+ rC[1][2]=mad(rA[0][1],rB[0][2],rC[1][2]); \
+ rC[2][2]=mad(rA[0][2],rB[0][2],rC[2][2]); \
+ rC[3][2]=mad(rA[0][3],rB[0][2],rC[3][2]); \
+ rC[0][3]=mad(rA[0][0],rB[0][3],rC[0][3]); \
+ rC[1][3]=mad(rA[0][1],rB[0][3],rC[1][3]); \
+ rC[2][3]=mad(rA[0][2],rB[0][3],rC[2][3]); \
+ rC[3][3]=mad(rA[0][3],rB[0][3],rC[3][3]); \
+ mem_fence(CLK_LOCAL_MEM_FENCE);
+
+#define M2x4 \
+ rA[0][0] = lA[offA + 0]; \
+ rA[0][1] = lA[offA + 16]; \
+ rB[0][0] = lB[offB + 0]; \
+ rB[0][1] = lB[offB + 16]; \
+ rB[0][2] = lB[offB + 32]; \
+ rB[0][3] = lB[offB + 48]; \
+ offA += 33; \
+ offB += 65; \
+ rC[0][0]=mad(rA[0][0],rB[0][0],rC[0][0]); \
+ rC[1][0]=mad(rA[0][1],rB[0][0],rC[1][0]); \
+ rC[0][1]=mad(rA[0][0],rB[0][1],rC[0][1]); \
+ rC[1][1]=mad(rA[0][1],rB[0][1],rC[1][1]); \
+ rC[0][2]=mad(rA[0][0],rB[0][2],rC[0][2]); \
+ rC[1][2]=mad(rA[0][1],rB[0][2],rC[1][2]); \
+ rC[0][3]=mad(rA[0][0],rB[0][3],rC[0][3]); \
+ rC[1][3]=mad(rA[0][1],rB[0][3],rC[1][3]); \
+ mem_fence(CLK_LOCAL_MEM_FENCE);
+
+#define M4x2 \
+ rA[0][0] = lA[offA + 0]; \
+ rA[0][1] = lA[offA + 16]; \
+ rA[0][2] = lA[offA + 32]; \
+ rA[0][3] = lA[offA + 48]; \
+ rB[0][0] = lB[offB + 0]; \
+ rB[0][1] = lB[offB + 16]; \
+ offA += 65; \
+ offB += 33; \
+ rC[0][0]=mad(rA[0][0],rB[0][0],rC[0][0]); \
+ rC[1][0]=mad(rA[0][1],rB[0][0],rC[1][0]); \
+ rC[2][0]=mad(rA[0][2],rB[0][0],rC[2][0]); \
+ rC[3][0]=mad(rA[0][3],rB[0][0],rC[3][0]); \
+ rC[0][1]=mad(rA[0][0],rB[0][1],rC[0][1]); \
+ rC[1][1]=mad(rA[0][1],rB[0][1],rC[1][1]); \
+ rC[2][1]=mad(rA[0][2],rB[0][1],rC[2][1]); \
+ rC[3][1]=mad(rA[0][3],rB[0][1],rC[3][1]); \
+ mem_fence(CLK_LOCAL_MEM_FENCE);
+
+#define M2x2 \
+ rA[0][0] = lA[offA + 0]; \
+ rA[0][1] = lA[offA + 16]; \
+ rB[0][0] = lB[offB + 0]; \
+ rB[0][1] = lB[offB + 16]; \
+ offA += 33; \
+ offB += 33; \
+ rC[0][0]=mad(rA[0][0],rB[0][0],rC[0][0]); \
+ rC[1][0]=mad(rA[0][1],rB[0][0],rC[1][0]); \
+ rC[0][1]=mad(rA[0][0],rB[0][1],rC[0][1]); \
+ rC[1][1]=mad(rA[0][1],rB[0][1],rC[1][1]); \
+ rC[2][1]=mad(rA[0][2],rB[0][1],rC[2][1]); \
+ mem_fence(CLK_LOCAL_MEM_FENCE);
+
+__attribute__((reqd_work_group_size(16,16,1)))
+__kernel void sgemm_NT_64_64_16_16x16_4x4__ALPHABETA_SPLIT_MAIN( __global float const * restrict A,
+ __global float const * restrict B,
+ __global float * C,
+ uint const M,
+ uint const N,
+ uint const K,
+ float const alpha,
+ float const beta,
+ uint lda,
+ uint ldb,
+ uint ldc,
+ uint offsetA,
+ uint offsetB,
+ uint offsetC)
+{
+ float rC[4][4] = {(float)0};
+ float rA[1][4];
+ float rB[1][4];
+
+
+
+ A += offsetA;
+ B += offsetB;
+ C+=offsetC;
+
+ __local float lA[1040];
+ __local float lB[1040];
+
+ uint gidx = get_group_id(0);
+ uint gidy = get_group_id(1);
+ uint idx = get_local_id(0);
+ uint idy = get_local_id(1);
+
+ A += gidx*64+ idx + idy*lda;
+ B += gidy*64+ idx + idy*ldb;
+
+
+ uint block_k = K >> 4;
+ do
+ {
+ // for(unsigned int block_k=0 ; block_k< K ; block_k+=16)
+ //{
+ __local float* plA = lA + idy*65+idx;
+ __local float* plB = lB + idy*65+idx;
+ barrier(CLK_LOCAL_MEM_FENCE);
+ plB[0] = B[0+0*ldb];
+ plB[16] = B[16+0*ldb];
+ plB[32] = B[32+0*ldb];
+ plB[48] = B[48+0*ldb];
+
+ plA[0] = A[0+0*lda];
+ plA[16] = A[16+0*lda];
+ plA[32] = A[32+0*lda];
+ plA[48] = A[48+0*lda];
+
+
+ barrier(CLK_LOCAL_MEM_FENCE);
+ uint offA = idx;
+ uint offB = idy;
+
+// #pragma unroll 1
+// for(unsigned int k = 0 ; k < 16; k+=1){
+// }
+
+ M4x4
+ M4x4
+ M4x4
+ M4x4
+ M4x4
+ M4x4
+ M4x4
+ M4x4
+ M4x4
+ M4x4
+ M4x4
+ M4x4
+ M4x4
+ M4x4
+ M4x4
+ M4x4
+
+ A += lda<<4;
+ B += ldb<<4;
+ //}
+ } while (--block_k > 0);
+
+ C+= gidx*64+idx;
+ C+= gidy*64*ldc;
+ C+= idy*ldc;
+
+ C[0*ldc] = alpha*rC[0][0] + beta*C[0*ldc];
+ C[16*ldc] = alpha*rC[0][1] + beta*C[16*ldc];
+ C[32*ldc] = alpha*rC[0][2] + beta*C[32*ldc];
+ C[48*ldc] = alpha*rC[0][3] + beta*C[48*ldc];
+ C+=16;
+ C[0*ldc] = alpha*rC[1][0] + beta*C[0*ldc];
+ C[16*ldc] = alpha*rC[1][1] + beta*C[16*ldc];
+ C[32*ldc] = alpha*rC[1][2] + beta*C[32*ldc];
+ C[48*ldc] = alpha*rC[1][3] + beta*C[48*ldc];
+ C+=16;
+ C[0*ldc] = alpha*rC[2][0] + beta*C[0*ldc];
+ C[16*ldc] = alpha*rC[2][1] + beta*C[16*ldc];
+ C[32*ldc] = alpha*rC[2][2] + beta*C[32*ldc];
+ C[48*ldc] = alpha*rC[2][3] + beta*C[48*ldc];
+ C+=16;
+ C[0*ldc] = alpha*rC[3][0] + beta*C[0*ldc];
+ C[16*ldc] = alpha*rC[3][1] + beta*C[16*ldc];
+ C[32*ldc] = alpha*rC[3][2] + beta*C[32*ldc];
+ C[48*ldc] = alpha*rC[3][3] + beta*C[48*ldc];
+
+}
+
+__attribute__((reqd_work_group_size(16,16,1)))
+__kernel void sgemm_NT_32_64_16_16x16_2x4__ALPHABETA_SPLIT_ROW( __global float const * restrict A,
+ __global float const * restrict B,
+ __global float * C,
+ uint const M,
+ uint const N,
+ uint const K,
+ float const alpha,
+ float const beta,
+ uint lda,
+ uint ldb,
+ uint ldc,
+ uint offsetA,
+ uint offsetB,
+ uint offsetC)
+{
+ float rC[2][4] = {(float)0};
+ float rA[1][2];
+ float rB[1][4];
+
+
+ A += offsetA;
+ B += offsetB;
+ C+=offsetC;
+
+ __local float lA[528];//16*32+16
+ __local float lB[1040];//16*64+16
+
+ uint gidx = M/64;//get_group_id(0);
+ uint gidy = get_group_id(1);
+ uint idx = get_local_id(0);
+ uint idy = get_local_id(1);
+
+
+ int CurrentOffSetA = gidx*64+ idx;
+
+ A += gidx*64+ idx + idy*lda;
+ B += gidy*64+ idx + idy*ldb;
+
+
+ uint block_k = K >> 4;
+ do
+ {
+ __local float* plA = lA + idy*33+idx;
+ __local float* plB = lB + idy*65+idx;
+ barrier(CLK_LOCAL_MEM_FENCE);
+
+ plB[0] = B[0+0*ldb];
+ plB[16] = B[16+0*ldb];
+ plB[32] = B[32+0*ldb];
+ plB[48] = B[48+0*ldb];
+
+ //plA[0] = CurrentOffSetA>=M?0.0:A[0];
+ //plA[16] = CurrentOffSetA+16>=M?0.0:A[16];
+ //plA[32] = CurrentOffSetA+32>=M?0.0:A[32];
+ //plA[48] = CurrentOffSetA+48>=M?0.0:A[48];
+ plA[0] = A[0];
+ plA[16] = A[16];
+
+
+ barrier(CLK_LOCAL_MEM_FENCE);
+ uint offA = idx;
+ uint offB = idy;
+
+
+ M2x4
+ M2x4
+ M2x4
+ M2x4
+ M2x4
+ M2x4
+ M2x4
+ M2x4
+ M2x4
+ M2x4
+ M2x4
+ M2x4
+ M2x4
+ M2x4
+ M2x4
+ M2x4
+
+ A += lda<<4;
+ B += ldb<<4;
+ } while (--block_k > 0);
+
+
+ int offset_x = gidx*64+idx;
+ int offset_y = gidy*64+ idy;
+
+ //if(offset_x>=M )
+ // return;
+
+ C+=offset_x+offset_y*ldc;
+
+ int i = 0;
+ do
+ {
+ C[0 ] = mad(alpha, rC[i][0], beta*C[0]);
+ C[16*ldc] = mad(alpha, rC[i][1], beta*C[16*ldc]);
+ C[32*ldc] = mad(alpha, rC[i][2], beta*C[32*ldc]);
+ C[48*ldc] = mad(alpha, rC[i][3], beta*C[48*ldc]);
+ C+=16;
+ offset_x+=16;
+ //if(offset_x>=M )
+ // return;
+ }
+ while (++i < 2);
+}
+
+__attribute__((reqd_work_group_size(16,16,1)))
+__kernel void sgemm_NT_64_32_16_16x16_4x2__ALPHABETA_SPLIT_COLUMN( __global float const * restrict A,
+ __global float const * restrict B,
+ __global float * C,
+ uint const M,
+ uint const N,
+ uint const K,
+ float const alpha,
+ float const beta,
+ uint lda,
+ uint ldb,
+ uint ldc,
+ uint offsetA,
+ uint offsetB,
+ uint offsetC)
+{
+ float rC[4][2] = {(float)0};
+ float rA[1][4];
+ float rB[1][2];
+
+
+ A += offsetA;
+ B += offsetB;
+ C+=offsetC;
+
+ __local float lA[1040];//16*64+16
+ __local float lB[528];//16*32+16
+
+ uint gidx = get_group_id(0);
+ uint gidy = N/64;//get_group_id(1);
+ uint idx = get_local_id(0);
+ uint idy = get_local_id(1);
+
+ int CurrentOffSetB = gidy*64+ idx;
+
+ A += gidx*64+ idx + idy*lda;
+ B += gidy*64+ idx + idy*ldb;
+
+
+ uint block_k = K >> 4;
+ do
+ {
+ __local float* plA = lA + idy*65+idx;
+ __local float* plB = lB + idy*33+idx;
+ barrier(CLK_LOCAL_MEM_FENCE);
+
+ //plB[0] = CurrentOffSetB>=N?0.0:B[0];
+ //plB[16] = CurrentOffSetB+16>=N?0.0:B[16];
+ //plB[32] = CurrentOffSetB+32>=N?0.0:B[32];
+ //plB[48] = CurrentOffSetB+48>=N?0.0:B[48];
+ plB[0] = B[0];
+ plB[16] = B[16];
+
+ plA[0] = A[0];
+ plA[16] = A[16];
+ plA[32] = A[32];
+ plA[48] = A[48];
+
+
+ barrier(CLK_LOCAL_MEM_FENCE);
+ uint offA = idx;
+ uint offB = idy;
+
+
+ M4x2
+ M4x2
+ M4x2
+ M4x2
+ M4x2
+ M4x2
+ M4x2
+ M4x2
+ M4x2
+ M4x2
+ M4x2
+ M4x2
+ M4x2
+ M4x2
+ M4x2
+ M4x2
+
+ A += lda<<4;
+ B += ldb<<4;
+ } while (--block_k > 0);
+
+
+ int offset_x = gidx*64+idx;
+ int offset_y = gidy*64+ idy;
+
+ //if(offset_y>=N )
+ // return;
+
+ C+=offset_x+offset_y*ldc;
+
+ int i = 0;
+ do
+ {
+ C[0 ] = mad(alpha, rC[i][0], beta*C[0]);
+ C[16*ldc] = mad(alpha, rC[i][1], beta*C[16*ldc]);
+
+ C+=16;
+
+ }
+ while (++i < 4);
+}
+
+__attribute__((reqd_work_group_size(16,16,1)))
+__kernel void sgemm_NT_32_32_16_16x16_2x2__ALPHABETA_SPLIT_SINGLE( __global float const * restrict A,
+ __global float const * restrict B,
+ __global float * C,
+ uint const M,
+ uint const N,
+ uint const K,
+ float const alpha,
+ float const beta,
+ uint lda,
+ uint ldb,
+ uint ldc,
+ uint offsetA,
+ uint offsetB,
+ uint offsetC)
+{
+ float rC[2][2] = {(float)0};
+ float rA[1][2];
+ float rB[1][2];
+
+
+ A += offsetA;
+ B += offsetB;
+ C+=offsetC;
+
+ __local float lA[528];
+ __local float lB[528];
+
+ uint gidx = M/64;//get_group_id(0);
+ uint gidy = N/64;//get_group_id(1);
+ uint idx = get_local_id(0);
+ uint idy = get_local_id(1);
+
+ int CurrentOffSetA = gidx*64+ idx;
+ int CurrentOffSetB = gidy*64+ idx;
+
+ A += gidx*64+ idx + idy*lda;
+ B += gidy*64+ idx + idy*ldb;
+
+
+ uint block_k = K >> 4;
+ do
+ {
+ __local float* plA = lA + idy*33+idx;
+ __local float* plB = lB + idy*33+idx;
+ barrier(CLK_LOCAL_MEM_FENCE);
+
+ //plB[0] = CurrentOffSetB>=N?0.0:B[0];
+ //plB[16] = CurrentOffSetB+16>=N?0.0:B[16];
+ //plB[32] = CurrentOffSetB+32>=N?0.0:B[32];
+ //plB[48] = CurrentOffSetB+48>=N?0.0:B[48];
+ plB[0] = B[0];
+ plB[16] = B[16];
+
+ //plA[0] = CurrentOffSetA>=M?0.0:A[0];
+ //plA[16] = CurrentOffSetA+16>=M?0.0:A[16];
+ //plA[32] = CurrentOffSetA+32>=M?0.0:A[32];
+ //plA[48] = CurrentOffSetA+48>=M?0.0:A[48];
+ plA[0] = A[0];
+ plA[16] = A[16];
+
+ barrier(CLK_LOCAL_MEM_FENCE);
+ uint offA = idx;
+ uint offB = idy;
+
+
+ M2x2
+ M2x2
+ M2x2
+ M2x2
+ M2x2
+ M2x2
+ M2x2
+ M2x2
+ M2x2
+ M2x2
+ M2x2
+ M2x2
+ M2x2
+ M2x2
+ M2x2
+ M2x2
+
+ A += lda<<4;
+ B += ldb<<4;
+ } while (--block_k > 0);
+
+
+ int offset_x = gidx*64+idx;
+ int offset_y = gidy*64+ idy;
+
+ //if(offset_x>=M || offset_y>=N )
+ // return;
+
+ C+=offset_x+offset_y*ldc;
+
+ int i = 0;
+ do
+ {
+ C[0 ] = mad(alpha, rC[i][0], beta*C[0]);
+ C[16*ldc] = mad(alpha, rC[i][1], beta*C[16*ldc]);
+
+
+ C+=16;
+ offset_x+=16;
+ //if(offset_x>=M )
+ // return;
+
+
+ }
+ while (++i < 2);
+}
+
+";
\ No newline at end of file
--
Alioth's /usr/local/bin/git-commit-notice on /srv/git.debian.org/git/debian-science/packages/clblas.git
More information about the debian-science-commits
mailing list