merge master branch to develop branch. please only make pull requests to develop branch

Author: Timmy

.gitignore

@@ -23,3 +23,6 @@
 
 # vim temp files
 .*.swp
+
+src/build/
+

.travis.yml

@@ -120,15 +120,15 @@ before_script:
 script:
   - make package
 
-deploy:
-  provider: releases
-  prerelease: true
-  draft: true
-  skip_cleanup: true
-  api_key:
-    secure: MBkxtcfSk+4UvGRO+WRhmS86vIVzAs0LIF2sAtr/S+Ed+OdUAuhZypUsDXGWtK3mL55v9c8BZXefFfHfJqElcNmyHKwCptbCR/JiM8YBtjoy2/RW1NcJUZp+QuRlk23xPADj7QkPjv7dfrQUMitkLUXAD+uTmMe2l8gmlbhMrQqPBKhb+31FNv6Lmo6oa6GjbiGi7qjsrJc7uQjhppLam+M7BZbBALGbIqMIrb2BMDMMhBoDbb4zSKrSg3+krd3kKiCClJlK7xjIlyFXZ527ETQ+PMtIeQb0eJ3aQwa4caBRCm5BDzt8GnJ48S88EkynbQioCEE87ebcyOM7M+wfslW/Fm1Y86X5odIljkOmTNKoDvgLxc9vUCBtMyVHNIgZcToPdsrMsGxcHV+JtU3yVQVm6dnA5P/zG5bA+aBjsd7p7BdOE4fdhvZV5XRAk/wmiyWalF7hKJxHIiWAKknL+tpPDDUF+fHmDDsdf7yRDJBegNcKfw4+m19MIvLn9fbiNVCtwCAL1T4yWkIEpi4MRMDPtftmkZPbi6UwluOJUTeCeHe4en99Yu2haemNPqXs6rR0LlXGk31GQwzlrNfb+94F5tT2a4Ka4PsruA2NMW/IYCYEE5Gu7PihVDR031Fn9cdCU9kefUgyB07rJD6q/W+ljsU0osyg7VxyfMg8rkw=
-  file: ${CLBLAS_ROOT}/clBLAS-build/*.tar.gz
-  file_glob: true
-  on:
-    all_branches: true
-    tags: true
+#deploy:
+#  provider: releases
+#  prerelease: true
+#  draft: true
+#  skip_cleanup: true
+#  api_key:
+#    secure: MBkxtcfSk+4UvGRO+WRhmS86vIVzAs0LIF2sAtr/S+Ed+OdUAuhZypUsDXGWtK3mL55v9c8BZXefFfHfJqElcNmyHKwCptbCR/JiM8YBtjoy2/RW1NcJUZp+QuRlk23xPADj7QkPjv7dfrQUMitkLUXAD+uTmMe2l8gmlbhMrQqPBKhb+31FNv6Lmo6oa6GjbiGi7qjsrJc7uQjhppLam+M7BZbBALGbIqMIrb2BMDMMhBoDbb4zSKrSg3+krd3kKiCClJlK7xjIlyFXZ527ETQ+PMtIeQb0eJ3aQwa4caBRCm5BDzt8GnJ48S88EkynbQioCEE87ebcyOM7M+wfslW/Fm1Y86X5odIljkOmTNKoDvgLxc9vUCBtMyVHNIgZcToPdsrMsGxcHV+JtU3yVQVm6dnA5P/zG5bA+aBjsd7p7BdOE4fdhvZV5XRAk/wmiyWalF7hKJxHIiWAKknL+tpPDDUF+fHmDDsdf7yRDJBegNcKfw4+m19MIvLn9fbiNVCtwCAL1T4yWkIEpi4MRMDPtftmkZPbi6UwluOJUTeCeHe4en99Yu2haemNPqXs6rR0LlXGk31GQwzlrNfb+94F5tT2a4Ka4PsruA2NMW/IYCYEE5Gu7PihVDR031Fn9cdCU9kefUgyB07rJD6q/W+ljsU0osyg7VxyfMg8rkw=
+#  file: ${CLBLAS_ROOT}/clBLAS-build/*.tar.gz
+#  file_glob: true
+#  on:
+#    all_branches: true
+#    tags: true

README.md

@@ -68,13 +68,12 @@ The simple example below shows how to use clBLAS to compute an OpenCL accelerate
     #include <sys/types.h>
     #include <stdio.h>
 
-    /* Include the clBLAS header. It includes the appropriate OpenCL headers
-     */
+    /* Include the clBLAS header. It includes the appropriate OpenCL headers */
     #include <clBLAS.h>
 
     /* This example uses predefined matrices and their characteristics for
      * simplicity purpose.
-     */
+    */
 
     #define M  4
     #define N  3
@@ -83,19 +82,19 @@ The simple example below shows how to use clBLAS to compute an OpenCL accelerate
     static const cl_float alpha = 10;
 
     static const cl_float A[M*K] = {
-        11, 12, 13, 14, 15,
-        21, 22, 23, 24, 25,
-        31, 32, 33, 34, 35,
-        41, 42, 43, 44, 45,
+    11, 12, 13, 14, 15,
+    21, 22, 23, 24, 25,
+    31, 32, 33, 34, 35,
+    41, 42, 43, 44, 45,
     };
     static const size_t lda = K;        /* i.e. lda = K */
 
     static const cl_float B[K*N] = {
-        11, 12, 13,
-        21, 22, 23,
-        31, 32, 33,
-        41, 42, 43,
-        51, 52, 53,
+    11, 12, 13,
+    21, 22, 23,
+    31, 32, 33,
+    41, 42, 43,
+    51, 52, 53,
     };
     static const size_t ldb = N;        /* i.e. ldb = N */
 
@@ -113,41 +112,41 @@ The simple example below shows how to use clBLAS to compute an OpenCL accelerate
 
     int main( void )
     {
-        cl_int err;
-        cl_platform_id platform = 0;
-        cl_device_id device = 0;
-        cl_context_properties props[3] = { CL_CONTEXT_PLATFORM, 0, 0 };
-        cl_context ctx = 0;
-        cl_command_queue queue = 0;
-        cl_mem bufA, bufB, bufC;
-        cl_event event = NULL;
-        int ret = 0;
-
-        /* Setup OpenCL environment. */
-        err = clGetPlatformIDs( 1, &platform, NULL );
-        err = clGetDeviceIDs( platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL );
-
-        props[1] = (cl_context_properties)platform;
-        ctx = clCreateContext( props, 1, &device, NULL, NULL, &err );
-        queue = clCreateCommandQueue( ctx, device, 0, &err );
-
-        /* Setup clBLAS */
-        err = clblasSetup( );
-
-        /* Prepare OpenCL memory objects and place matrices inside them. */
-        bufA = clCreateBuffer( ctx, CL_MEM_READ_ONLY, M * K * sizeof(*A),
-                              NULL, &err );
-        bufB = clCreateBuffer( ctx, CL_MEM_READ_ONLY, K * N * sizeof(*B),
-                              NULL, &err );
-        bufC = clCreateBuffer( ctx, CL_MEM_READ_WRITE, M * N * sizeof(*C),
-                              NULL, &err );
-
-        err = clEnqueueWriteBuffer( queue, bufA, CL_TRUE, 0,
-            M * K * sizeof( *A ), A, 0, NULL, NULL );
-        err = clEnqueueWriteBuffer( queue, bufB, CL_TRUE, 0,
-            K * N * sizeof( *B ), B, 0, NULL, NULL );
-        err = clEnqueueWriteBuffer( queue, bufC, CL_TRUE, 0,
-            M * N * sizeof( *C ), C, 0, NULL, NULL );
+    cl_int err;
+    cl_platform_id platform = 0;
+    cl_device_id device = 0;
+    cl_context_properties props[3] = { CL_CONTEXT_PLATFORM, 0, 0 };
+    cl_context ctx = 0;
+    cl_command_queue queue = 0;
+    cl_mem bufA, bufB, bufC;
+    cl_event event = NULL;
+    int ret = 0;
+
+    /* Setup OpenCL environment. */
+    err = clGetPlatformIDs( 1, &platform, NULL );
+    err = clGetDeviceIDs( platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL );
+
+    props[1] = (cl_context_properties)platform;
+    ctx = clCreateContext( props, 1, &device, NULL, NULL, &err );
+    queue = clCreateCommandQueue( ctx, device, 0, &err );
+
+    /* Setup clBLAS */
+    err = clblasSetup( );
+
+    /* Prepare OpenCL memory objects and place matrices inside them. */
+    bufA = clCreateBuffer( ctx, CL_MEM_READ_ONLY, M * K * sizeof(*A),
+                          NULL, &err );
+    bufB = clCreateBuffer( ctx, CL_MEM_READ_ONLY, K * N * sizeof(*B),
+                          NULL, &err );
+    bufC = clCreateBuffer( ctx, CL_MEM_READ_WRITE, M * N * sizeof(*C),
+                          NULL, &err );
+
+    err = clEnqueueWriteBuffer( queue, bufA, CL_TRUE, 0,
+        M * K * sizeof( *A ), A, 0, NULL, NULL );
+    err = clEnqueueWriteBuffer( queue, bufB, CL_TRUE, 0,
+        K * N * sizeof( *B ), B, 0, NULL, NULL );
+    err = clEnqueueWriteBuffer( queue, bufC, CL_TRUE, 0,
+        M * N * sizeof( *C ), C, 0, NULL, NULL );
 
         /* Call clBLAS extended function. Perform gemm for the lower right sub-matrices */
         err = clblasSgemm( clblasRowMajor, clblasNoTrans, clblasNoTrans,
@@ -157,27 +156,27 @@ The simple example below shows how to use clBLAS to compute an OpenCL accelerate
                                 bufC, 0, ldc,
                                 1, &queue, 0, NULL, &event );
 
-        /* Wait for calculations to be finished. */
-        err = clWaitForEvents( 1, &event );
+    /* Wait for calculations to be finished. */
+    err = clWaitForEvents( 1, &event );
 
-        /* Fetch results of calculations from GPU memory. */
-        err = clEnqueueReadBuffer( queue, bufC, CL_TRUE, 0,
-                                    M * N * sizeof(*result),
-                                    result, 0, NULL, NULL );
+    /* Fetch results of calculations from GPU memory. */
+    err = clEnqueueReadBuffer( queue, bufC, CL_TRUE, 0,
+                                M * N * sizeof(*result),
+                                result, 0, NULL, NULL );
 
-        /* Release OpenCL memory objects. */
-        clReleaseMemObject( bufC );
-        clReleaseMemObject( bufB );
-        clReleaseMemObject( bufA );
+    /* Release OpenCL memory objects. */
+    clReleaseMemObject( bufC );
+    clReleaseMemObject( bufB );
+    clReleaseMemObject( bufA );
 
-        /* Finalize work with clBLAS */
-        clblasTeardown( );
+    /* Finalize work with clBLAS */
+    clblasTeardown( );
 
-        /* Release OpenCL working objects. */
-        clReleaseCommandQueue( queue );
-        clReleaseContext( ctx );
+    /* Release OpenCL working objects. */
+    clReleaseCommandQueue( queue );
+    clReleaseContext( ctx );
 
-        return ret;
+    return ret;
     }
 ```
 

appveyor.yml

@@ -91,15 +91,15 @@ on_finish:
   - echo on_finish step
 
 # Appveyor will push the artifacts it has saved to GitHub 'releases' tab
-deploy:
-  provider: GitHub
-  auth_token:
-    secure: dRXIWJKpU7h2RsHX7RqmyYCtCw+Q9O3X5MArloY6p34GZC1w7bp+jQYTZqbdO7bw
-  artifact: binary_zip
-  draft: true
-  prerelease: true
-  on:
-    appveyor_repo_tag: true
+# deploy:
+#  provider: GitHub
+#  auth_token:
+#    secure: dRXIWJKpU7h2RsHX7RqmyYCtCw+Q9O3X5MArloY6p34GZC1w7bp+jQYTZqbdO7bw
+#  artifact: binary_zip
+#  draft: true
+#  prerelease: true
+#  on:
+#    appveyor_repo_tag: true
 
   # Uncomment the following to pause the VM and wait for RDP connetion to debug
   # - ps: $blockRdp = $true; iex ((new-object net.webclient).DownloadString('https://raw.githubusercontent.com/appveyor/ci/master/scripts/enable-rdp.ps1'))

src/CMakeLists.txt

@@ -199,7 +199,7 @@ if( BUILD_TEST )
 	    else()
 			find_package( Netlib COMPONENTS BLAS REQUIRED )
         endif()
-	else( )
+		else( )
 		# Find ACML BLAS implementation
 		# platform dependent ACML subdirectory
 		if (WIN32)

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/AutoGemmTeardown.h

@@ -0,0 +1,10 @@
+#pragma once
+
+#ifdef __cplusplus
+#extern "C" {
+#endif
+  void initAutoGemmClKernels(void);
+#ifdef __cplusplus
+}
+#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];