Merge pull request #163 from hughperkins/fix-teardown

Fix teardown

Author: guacamoleo

src/library/CMakeLists.txt

@@ -839,8 +839,7 @@ else()
   MESSAGE( STATUS "clBLAS will NOT depend on ${AUTOGEMM_PRECOMPILED_KERNELS}" )
 endif()
 
-
-add_library(clBLAS
+set(CLBLAS_ALL_SOURCES
     ${CLBLAS_SOURCES}
     ${GLOBAL_HEADERS}
     ${SRC_BLAS_HEADERS}
@@ -851,9 +850,23 @@ add_library(clBLAS
     ${AUTOGEMM_PRECOMPILED_KERNELS_CONDITIONAL}
 	#${USERGEMM_SRC}
 	#${USERGEMM_HEADERS}
-  )
+)
+add_library(clBLAS ${CLBLAS_ALL_SOURCES})
 add_dependencies(clBLAS GENERATE_CLT)
 
+function (add_target_definitions target)
+  get_target_property(defs ${target} COMPILE_DEFINITIONS)
+  if (defs MATCHES "NOTFOUND")
+    set(defs "")
+  endif ()
+  foreach (def ${defs} ${ARGN})
+    list(APPEND deflist ${def})
+  endforeach ()
+  set_target_properties(${target} PROPERTIES COMPILE_DEFINITIONS "${deflist}")
+endfunction ()
+
+add_target_definitions(clBLAS BUILDING_CLBLAS)
+
 if (PRECOMPILE_TRSM_DTRSM OR PRECOMPILE_TRSM_STRSM)
 add_dependencies(clBLAS OCLBinaryGenerator_GEN)
 endif()

src/library/blas/AutoGemm/Includes.py

@@ -179,6 +179,14 @@ def __init__(self):
     self.incStr += "#include <CL/cl.h>\n"
     self.incStr += "#endif\n"
     self.incStr += "\n"
+    self.incStr += "#ifdef __cplusplus\n"
+    self.incStr += "extern \"C\" {\n"
+    self.incStr += "#endif\n"
+    self.incStr += "    void initAutoGemmClKernels(void);\n";
+    self.incStr += "#ifdef __cplusplus\n"
+    self.incStr += "}\n";
+    self.incStr += "#endif\n"
+    self.incStr += "\n";
 
     self.cppName = Common.getIncludePath() + "AutoGemmClKernels.cpp"
     self.cppFile = open(self.cppName, "w")
@@ -190,29 +198,50 @@ def __init__(self):
     self.cppStr += "#endif\n"
     self.cppStr += "\n"
 
+
+    self.initFunction = "";
+    self.initFunction += "extern \"C\" {\n";
+    self.initFunction += "    void initAutoGemmClKernels(void);\n";
+    self.initFunction += "}\n";
+    self.initFunction += "\n";
+    self.initFunction += "void initAutoGemmClKernels(void) {\n";
+    self.defines = "";
+
   def addKernel(self, kernel):
-    kernelName = kernel.getName()
-    self.incStr += "extern cl_kernel %s_clKernel;\n" % kernelName
-    self.cppStr += "cl_kernel %s_clKernel = NULL;\n" % kernelName
-    kernelName = kernel.getRowName()
-    self.incStr += "extern cl_kernel %s_clKernel;\n" % kernelName
-    self.cppStr += "cl_kernel %s_clKernel = NULL;\n" % kernelName
-    kernelName = kernel.getColName()
-    self.incStr += "extern cl_kernel %s_clKernel;\n" % kernelName
-    self.cppStr += "cl_kernel %s_clKernel = NULL;\n" % kernelName
-    kernelName = kernel.getCornerName()
-    self.incStr += "extern cl_kernel %s_clKernel;\n" % kernelName
-    self.cppStr += "cl_kernel %s_clKernel = NULL;\n" % kernelName
+    kernelNames = [
+      kernel.getName(),
+      kernel.getRowName(),
+      kernel.getColName(),
+      kernel.getCornerName()
+    ]
+    for kernelName in kernelNames:
+      self.incStr += "extern cl_kernel %s_clKernel;\n" % kernelName
+
+      self.defines += "cl_kernel %s_clKernel = NULL;\n" % kernelName
+
+      self.initFunction += "    if(%s_clKernel != NULL) {\n" % kernelName
+      self.initFunction += "        clReleaseKernel(%s_clKernel);\n" % kernelName
+      self.initFunction += "        %s_clKernel = NULL;\n" % kernelName
+      self.initFunction += "    }\n"
 
     self.incFile.write( self.incStr )
     self.incStr = ""
-    self.cppFile.write( self.cppStr )
-    self.cppStr = ""
+#   self.cppFile.write( self.cppStr )
+#   self.cppStr = ""
 
   def writeToFile(self):
     self.incFile.write( self.incStr )
     self.incFile.write( "\n#endif\n" )
     self.incFile.close()
+
+    self.initFunction += "}\n";
+    self.cppStr += self.defines + "\n";
+    self.defines = "";
+    self.cppStr += self.initFunction + "\n";
+    self.initFunction = "";
+
+#    self.cppStr += "\n";
+#    self.cppStr += "initAutoGemmClKernels();\n";
     self.cppFile.write( self.cppStr )
     self.cppFile.close()
 

src/library/blas/AutoGemm/KernelOpenCL.py

@@ -218,7 +218,7 @@ def makeOpenCLKernelString(kernel):
   kStr += endLine
   kStr += (
     "  /* allocate registers */" + endLine +
-    "  DATA_TYPE_STR rC[MICRO_TILE_NUM_ROWS][MICRO_TILE_NUM_COLS] = {0};" + endLine +
+    "  DATA_TYPE_STR rC[MICRO_TILE_NUM_ROWS][MICRO_TILE_NUM_COLS] = { {0} };" + endLine +
     "  DATA_TYPE_STR rA[MICRO_TILE_NUM_ROWS];" + endLine +
     "  DATA_TYPE_STR rB[MICRO_TILE_NUM_COLS];" + endLine )
 

src/library/blas/gens/clTemplates/zgemm_gcn.cl

@@ -210,7 +210,7 @@ __kernel void KERNEL_NAME(DATA_TYPE_CHAR,TRANSPOSE_A,TRANSPOSE_B,MACRO_TILE_NUM_
   C += offsetC;
 
   // registers
-  DATA_TYPE_STR rC[MICRO_TILE_NUM_ROWS][MICRO_TILE_NUM_COLS]  = {0};
+  DATA_TYPE_STR rC[MICRO_TILE_NUM_ROWS][MICRO_TILE_NUM_COLS]  = { {0} };
   DATA_TYPE_STR rA[MICRO_TILE_NUM_ROWS];
   DATA_TYPE_STR rB[MICRO_TILE_NUM_COLS];
 

src/library/blas/init.c

@@ -25,6 +25,9 @@
 #include <events.h>
 #include <stdlib.h>
 #include <stdio.h>
+#ifdef BUILDING_CLBLAS
+#include "AutoGemmIncludes/AutoGemmClKernels.h"
+#endif
 
 clblasStatus
 clblasGetVersion(cl_uint* major, cl_uint* minor, cl_uint* patch)
@@ -250,5 +253,9 @@ clblasTeardown(void)
     printMemLeaksInfo();
     releaseMallocTrace();
 
+#ifdef BUILDING_CLBLAS
+   initAutoGemmClKernels();
+#endif
+
     clblasInitialized = 0;
 }