Merge pull request #175 from hughperkins/fix-teardown-userkernels-rebase-develop

Fix teardown userkernels rebase develop

Author: TimmyLiu

src/library/CMakeLists.txt

@@ -43,6 +43,7 @@ set(AUTOGEMM_HEADERS
 )
 
 set(AUTOGEMM_SRC
+  ${CMAKE_SOURCE_DIR}/library/blas/AutoGemm/UserGemmKernelSources/UserGemmClKernels.cc
   ${CMAKE_BINARY_DIR}/include/AutoGemmIncludes/AutoGemmClKernels.cpp
   ${CMAKE_BINARY_DIR}/include/AutoGemmIncludes/AutoGemmKernelBuildOptionsBinary.cpp
   ${CMAKE_BINARY_DIR}/include/AutoGemmIncludes/AutoGemmKernelBinaries.cpp
@@ -851,6 +852,7 @@ set(CLBLAS_ALL_SOURCES
 	#${USERGEMM_SRC}
 	#${USERGEMM_HEADERS}
 )
+add_definitions(-DOPENCL_VERSION="${OPENCL_VERSION}")
 add_library(clBLAS ${CLBLAS_ALL_SOURCES})
 add_dependencies(clBLAS GENERATE_CLT)
 

src/library/blas/AutoGemm/UserGemmKernelSources/UserGemmClKernels.cc

@@ -0,0 +1,54 @@
+// GENERATED using create_user_gemm_cl_kernels.py
+
+#if defined( __APPLE__ ) || defined( __MACOSX )
+#include <OpenCL/cl.h>
+#else
+#include <CL/cl.h>
+#endif
+
+cl_kernel sgemm_Col_NT_B1_MX128_NX128_KX16_clKernel = NULL;
+cl_kernel sgemm_Col_NT_B1_MX032_NX064_KX16_ROW_clKernel = NULL;
+cl_kernel sgemm_Col_NT_B1_MX064_NX032_KX16_COLUMN_clKernel = NULL;
+cl_kernel sgemm_Col_NT_B1_MX032_NX032_KX16_SINGLE_clKernel = NULL;
+cl_kernel sgemm_Col_NN_B1_MX032_NX032_KX16_BRANCH_clKernel = NULL;
+cl_kernel sgemm_Col_NT_B1_MX032_NX032_KX16_BRANCH_clKernel = NULL;
+cl_kernel sgemm_Col_TN_B1_MX032_NX032_KX16_BRANCH_clKernel = NULL;
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+    void initUserGemmClKernels(void);
+#ifdef __cplusplus
+}
+#endif
+
+void initUserGemmClKernels(void) {
+    if(sgemm_Col_NT_B1_MX128_NX128_KX16_clKernel != NULL) {
+        clReleaseKernel(sgemm_Col_NT_B1_MX128_NX128_KX16_clKernel);
+        sgemm_Col_NT_B1_MX128_NX128_KX16_clKernel = NULL;
+    }
+    if(sgemm_Col_NT_B1_MX032_NX064_KX16_ROW_clKernel != NULL) {
+        clReleaseKernel(sgemm_Col_NT_B1_MX032_NX064_KX16_ROW_clKernel);
+        sgemm_Col_NT_B1_MX032_NX064_KX16_ROW_clKernel = NULL;
+    }
+    if(sgemm_Col_NT_B1_MX064_NX032_KX16_COLUMN_clKernel != NULL) {
+        clReleaseKernel(sgemm_Col_NT_B1_MX064_NX032_KX16_COLUMN_clKernel);
+        sgemm_Col_NT_B1_MX064_NX032_KX16_COLUMN_clKernel = NULL;
+    }
+    if(sgemm_Col_NT_B1_MX032_NX032_KX16_SINGLE_clKernel != NULL) {
+        clReleaseKernel(sgemm_Col_NT_B1_MX032_NX032_KX16_SINGLE_clKernel);
+        sgemm_Col_NT_B1_MX032_NX032_KX16_SINGLE_clKernel = NULL;
+    }
+    if(sgemm_Col_NN_B1_MX032_NX032_KX16_BRANCH_clKernel != NULL) {
+        clReleaseKernel(sgemm_Col_NN_B1_MX032_NX032_KX16_BRANCH_clKernel);
+        sgemm_Col_NN_B1_MX032_NX032_KX16_BRANCH_clKernel = NULL;
+    }
+    if(sgemm_Col_NT_B1_MX032_NX032_KX16_BRANCH_clKernel != NULL) {
+        clReleaseKernel(sgemm_Col_NT_B1_MX032_NX032_KX16_BRANCH_clKernel);
+        sgemm_Col_NT_B1_MX032_NX032_KX16_BRANCH_clKernel = NULL;
+    }
+    if(sgemm_Col_TN_B1_MX032_NX032_KX16_BRANCH_clKernel != NULL) {
+        clReleaseKernel(sgemm_Col_TN_B1_MX032_NX032_KX16_BRANCH_clKernel);
+        sgemm_Col_TN_B1_MX032_NX032_KX16_BRANCH_clKernel = NULL;
+    }
+}

src/library/blas/AutoGemm/UserGemmKernelSources/UserGemmClKernels.h

@@ -8,16 +8,24 @@
 #include <CL/cl.h>
 #endif
 
-static cl_kernel sgemm_Col_NT_B1_MX128_NX128_KX16_clKernel = NULL;
+extern cl_kernel sgemm_Col_NT_B1_MX128_NX128_KX16_clKernel;
 
-static cl_kernel sgemm_Col_NT_B1_MX032_NX064_KX16_ROW_clKernel = NULL;
-static cl_kernel sgemm_Col_NT_B1_MX064_NX032_KX16_COLUMN_clKernel = NULL;
-static cl_kernel sgemm_Col_NT_B1_MX032_NX032_KX16_SINGLE_clKernel = NULL;
+extern cl_kernel sgemm_Col_NT_B1_MX032_NX064_KX16_ROW_clKernel;
+extern cl_kernel sgemm_Col_NT_B1_MX064_NX032_KX16_COLUMN_clKernel;
+extern cl_kernel sgemm_Col_NT_B1_MX032_NX032_KX16_SINGLE_clKernel;
 
-static cl_kernel sgemm_Col_NN_B1_MX032_NX032_KX16_BRANCH_clKernel = NULL;
-static cl_kernel sgemm_Col_NT_B1_MX032_NX032_KX16_BRANCH_clKernel = NULL;
-static cl_kernel sgemm_Col_TN_B1_MX032_NX032_KX16_BRANCH_clKernel = NULL;
+extern cl_kernel sgemm_Col_NN_B1_MX032_NX032_KX16_BRANCH_clKernel;
+extern cl_kernel sgemm_Col_NT_B1_MX032_NX032_KX16_BRANCH_clKernel;
+extern cl_kernel sgemm_Col_TN_B1_MX032_NX032_KX16_BRANCH_clKernel;
 
 static const int user_kernel_count = 7;
 
+#ifdef __cplusplus
+extern "C" {
+#endif
+    void initUserGemmClKernels(void);
+#ifdef __cplusplus
+}
+#endif
+
 #endif

src/library/blas/AutoGemm/UserGemmKernelSources/UserGemmKernelSourceIncludes.h

@@ -9,8 +9,8 @@
 
 //**** compiler flags
 //**** online compilation flags
-const char * const User_srcBuildOptions = "-cl-std=CL2.0";
-const char * const User_binBuildOptions = "-cl-std=CL2.0";
+const char * const User_srcBuildOptions = "-cl-std=CL" OPENCL_VERSION;
+const char * const User_binBuildOptions = "-cl-std=CL" OPENCL_VERSION;
 
 
 extern const unsigned int sgemm_Col_NT_B1_MX032_NX064_KX16_ROW_workGroupNumRows;

src/library/blas/AutoGemm/UserGemmKernelSources/create_user_gemm_cl_kernels.py

@@ -0,0 +1,44 @@
+"""
+Run from same directory as this file is in
+Generates UserGemmClKernels.cc
+
+Would be nice to use Jinja2 for this, but using print for now, for consistency
+"""
+
+kernelNames = []
+# lets just read the kernel names from UserGemmClKernels.h:
+ifile = open('UserGemmClKernels.h', 'r')
+contents = ifile.read()
+for line in contents.split('\n'):
+    if line.find('cl_kernel') < 0:
+        continue
+    kernelName = line.split()[2].split(';')[0]  # probably not terribly un-fragile, but works for now
+    kernelNames.append(kernelName)
+ifile.close()
+
+ofile = open('UserGemmClKernels.cc', 'w')
+
+ofile.write('// GENERATED using create_user_gemm_cl_kernels.py\n')
+ofile.write('\n')
+
+ofile.write('#if defined( __APPLE__ ) || defined( __MACOSX )\n')
+ofile.write('#include <OpenCL/cl.h>\n')
+ofile.write('#else\n')
+ofile.write('#include <CL/cl.h>\n')
+ofile.write('#endif\n')
+ofile.write('\n')
+
+for kernelName in kernelNames:
+    ofile.write('cl_kernel %s = NULL;\n' % kernelName)
+ofile.write('\n')
+
+ofile.write('void initUserGemmClKernels(void) {\n')
+
+for kernelName in kernelNames:
+    ofile.write('    if(%s != NULL) {\n' % kernelName)
+    ofile.write('        clReleaseKernel(%s);\n' % kernelName)
+    ofile.write('        %s = NULL;\n' % kernelName)
+    ofile.write('    }\n')
+
+ofile.write('}\n')
+ofile.close()

src/library/blas/AutoGemm/UserGemmKernelSources/dgemm_Col_NN_B0_MX048_NX048_KX08_src.cpp

@@ -96,7 +96,7 @@ __kernel void dgemm_Col_NN_B0_MX048_NX048_KX08 (
     C += offsetC;
 
 
-    double rC[6][6] = {(double)0};
+    double rC[6][6] = { {(double)0} };
     double rA[6];
     double rB[6];
   __local double lA[392];

src/library/blas/AutoGemm/UserGemmKernelSources/dgemm_Col_NN_B1_MX048_NX048_KX08_src.cpp

@@ -94,7 +94,7 @@ __kernel void dgemm_Col_NN_B1_MX048_NX048_KX08 (
     C += offsetC;
 
 
-    double rC[6][6] = {(double)0};
+    double rC[6][6] = { {(double)0} };
     double rA[6];
     double rB[6];
 

src/library/blas/AutoGemm/UserGemmKernelSources/dgemm_Col_NT_B0_MX048_NX048_KX08_src.cpp

@@ -108,7 +108,7 @@ const char * const dgemm_Col_NT_B0_MX048_NX048_KX08_src = STRINGIFY(
 \n    C    += offsetC;
 \n
 \n
-\n    double rC[6][6] = {(double)0};
+\n    double rC[6][6] = { {(double)0} };
 \n    double rA[6];
 \n    double rB[6];
 \n

src/library/blas/AutoGemm/UserGemmKernelSources/dgemm_Col_NT_B1_MX048_NX048_KX08_src.cpp

@@ -107,7 +107,7 @@ const char * const dgemm_Col_NT_B1_MX048_NX048_KX08_src = STRINGIFY(
 \n    C    += offsetC;
 \n
 \n
-\n    double rC[6][6] = {(double)0};
+\n    double rC[6][6] = { {(double)0} };
 \n    double rA[6];
 \n    double rB[6];
 \n

src/library/blas/AutoGemm/UserGemmKernelSources/dgemm_Col_TN_B0_MX048_NX048_KX08_src.cpp

@@ -36,7 +36,7 @@ __kernel void dgemm_Col_TN_B0_MX048_NX048_KX08_src (
   uint const offsetB,
   uint const offsetC )
 {
-    double rC[6][6]  = {(double)0};
+    double rC[6][6]  = { {(double)0} };
     double rA[1][6];
     double rB[1][6];