Merge pull request clMathLibraries#80 from TimmyLiu/master

merge develop branch into master branch

Author: Timmy

.travis.yml

@@ -5,32 +5,24 @@ compiler:
 
 before_install:
   - sudo apt-get update -qq
-  - sudo apt-get install -qq fglrx opencl-headers libboost-program-options-dev libgtest-dev
+  - sudo apt-get install -qq fglrx opencl-headers libboost-program-options-dev
 # Uncomment below to help verify the installs above work
 #  - ls -la /usr/lib/libboost*
 #  - ls -la /usr/include/boost
-#  - ls -la /usr/src/gtest
-
-install:
-  - mkdir -p bin/gTest
-  - cd bin/gTest
-  - cmake -DCMAKE_BUILD_TYPE=Release /usr/src/gtest
-  - make
-  - sudo mv libg* /usr/lib
 
 before_script:
   - cd ${TRAVIS_BUILD_DIR}
   - mkdir -p bin/clBLAS
   - cd bin/clBLAS
-  - cmake -DCMAKE_BUILD_TYPE=Release -DBUILD_TEST=OFF -DBUILD_CLIENT=ON ../../src
+  - cmake -DCMAKE_BUILD_TYPE=Release -DBUILD_TEST=OFF -DBUILD_CLIENT=ON -DCMAKE_INSTALL_PREFIX:PATH=$PWD/package ../../src
 
 script: 
   - make install
 #  - ls -Rla package
 # Run a simple test to validate that the build works; CPU device in a VM
   - cd package/bin
   - export LD_LIBRARY_PATH=${TRAVIS_BUILD_DIR}/bin/clBLAS/package/lib64:${LD_LIBRARY_PATH}
-  - ./client --cpu
+  - ./clBLAS-client --cpu
 
 after_success:
   - cd ${TRAVIS_BUILD_DIR}/bin/clBLAS

src/CMakeLists.txt

@@ -23,6 +23,7 @@ option( BUILD_PERFORMANCE "Copy the performance scripts that can measure and gra
 option( BUILD_SAMPLE "Build the sample programs" OFF )
 option( BUILD_CLIENT "Build a command line clBLAS client program with a variety of configurable parameters (dependency on Boost)" OFF )
 option( BUILD_KTEST "A command line tool for testing single clBLAS kernel" ON )
+option( BUILD_SHARED_LIBS "Build shared libraries" ON )
 
 # By default test-correctness is linked and tested against ACML library.
 # However, test-correctness can instead use NETLIB as a reference library
@@ -50,7 +51,7 @@ if( NOT DEFINED clBLAS_VERSION_MAJOR )
 endif( )
 
 if( NOT DEFINED clBLAS_VERSION_MINOR )
-    set( clBLAS_VERSION_MINOR 2 )
+    set( clBLAS_VERSION_MINOR 4 )
 endif( )
 
 if( NOT DEFINED clBLAS_VERSION_PATCH )
@@ -67,7 +68,7 @@ set( CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} ${PROJECT_SOURCE_DIR} )
 
 # On windows, it's convenient to change the default install prefix such that it does NOT point to 'program files' (permissions problems)
 # Need to check out CMAKE_RUNTIME_OUTPUT_DIRECTORY variable, and see if that eliminates the need to modify install path
-if( CMAKE_INSTALL_PREFIX_INITIALIZED_TO_DEFAULT )
+if( WIN32 AND CMAKE_INSTALL_PREFIX_INITIALIZED_TO_DEFAULT )
 	set( CMAKE_INSTALL_PREFIX "${PROJECT_BINARY_DIR}/package" CACHE PATH "Install path prefix, prepended onto install directories" FORCE )
 endif( )
 
@@ -84,34 +85,33 @@ set( SUFFIX_BIN_DEFAULT "" )
 
 if(TARGET_PLATFORM EQUAL 32 OR TARGET_PLATFORM EQUAL 64)
     set(TARGET_PLATFORM ${TARGET_PLATFORM} CACHE STRING "Target platform type (32-bit or 64-bit)" FORCE)
-    if( TARGET_PLATFORM EQUAL 64 )
-        set( SUFFIX_LIB_DEFAULT "64" )
-    endif( )
 else()
     if(CMAKE_SIZEOF_VOID_P MATCHES 8)
         set(TARGET_PLATFORM "64" CACHE STRING "Target platform type (32-bit or 64-bit)" FORCE)
-        set( SUFFIX_LIB_DEFAULT "64" )
     else()
         set(TARGET_PLATFORM "32" CACHE STRING "Target platform type (32-bit or 64-bit)" FORCE)
     endif()
 endif()
 
-set( SUFFIX_LIB ${SUFFIX_LIB_DEFAULT} CACHE STRING "String to append to 'lib' install path" )
-set( SUFFIX_BIN ${SUFFIX_BIN_DEFAULT} CACHE STRING "String to append to 'bin' install path" )
-
-if( MSVC_IDE )
-    set_property( GLOBAL PROPERTY USE_FOLDERS TRUE )
-endif( )
-
 message(STATUS "Target platform: ${TARGET_PLATFORM}-bit")
 if(TARGET_PLATFORM EQUAL 32)
     set(_arch "x86" INTERNAL)
     set_property(GLOBAL PROPERTY FIND_LIBRARY_USE_LIB64_PATHS FALSE)
 else()
     set(_arch "x86_64" INTERNAL)
     set_property(GLOBAL PROPERTY FIND_LIBRARY_USE_LIB64_PATHS TRUE)
+    if( NOT APPLE )
+        set( SUFFIX_LIB_DEFAULT "64" )
+    endif( )
 endif()
 
+set( SUFFIX_LIB ${SUFFIX_LIB_DEFAULT} CACHE STRING "String to append to 'lib' install path" )
+set( SUFFIX_BIN ${SUFFIX_BIN_DEFAULT} CACHE STRING "String to append to 'bin' install path" )
+
+if( MSVC_IDE )
+    set_property( GLOBAL PROPERTY USE_FOLDERS TRUE )
+endif( )
+
 # add the math library for Linux
 if( UNIX ) 
     set(MATH_LIBRARY "m")
@@ -197,9 +197,6 @@ endif( )
 # This will define OPENCL_FOUND
 find_package( OpenCL )
 
-# Find Google Test package
-find_package( GTest )
-
 # Find Boost on the system, and configure the type of boost build we want
 set( Boost_USE_MULTITHREADED ON )
 set( Boost_USE_STATIC_LIBS   ON )
@@ -297,6 +294,26 @@ if( BUILD_TEST )
 	endif( )
 endif( )
 
+if(WIN32)
+  set(destdir CMake)
+else()
+  set(destdir share/clBLAS)
+endif()
+string(REGEX REPLACE "[^/]+" ".." reldir "${destdir}")
+configure_file(
+  ${CMAKE_CURRENT_SOURCE_DIR}/clBLASConfigVersion.cmake.in
+  ${CMAKE_CURRENT_BINARY_DIR}/clBLASConfigVersion.cmake
+  @ONLY)
+configure_file(
+  ${CMAKE_CURRENT_SOURCE_DIR}/clBLASConfig.cmake.in
+  ${CMAKE_CURRENT_BINARY_DIR}/clBLASConfig.cmake
+  @ONLY)
+install(EXPORT Library DESTINATION ${destdir} FILE clBLASTargets.cmake)
+install(FILES ${CMAKE_CURRENT_BINARY_DIR}/clBLASConfigVersion.cmake
+  ${CMAKE_CURRENT_BINARY_DIR}/clBLASConfig.cmake
+  DESTINATION ${destdir})
+
+
 # The following code is setting variables to control the behavior of CPack to generate our 
 if( WIN32 )
 	set( CPACK_SOURCE_GENERATOR "ZIP" )

src/FindOpenCL.cmake

@@ -56,6 +56,8 @@ find_path(OPENCL_INCLUDE_DIRS
     PATHS
         /usr/include
         /usr/local/include
+        /usr/local/cuda/include
+        /opt/cuda/include
     DOC "OpenCL header file path"
 )
 mark_as_advanced( OPENCL_INCLUDE_DIRS )
@@ -74,6 +76,8 @@ if( LIB64 )
         PATH_SUFFIXES x86_64 x64
         PATHS
             /usr/lib
+            /usr/local/cuda/lib
+            /opt/cuda/lib
     )
 else( )
     find_library( OPENCL_LIBRARIES
@@ -86,6 +90,8 @@ else( )
         PATH_SUFFIXES x86 Win32
         PATHS
             /usr/lib
+            /usr/local/cuda/lib
+            /opt/cuda/lib
     )
 endif( )
 mark_as_advanced( OPENCL_LIBRARIES )

src/clBLASConfig.cmake.in

@@ -0,0 +1,3 @@
+include(${CMAKE_CURRENT_LIST_DIR}/clBLASTargets.cmake)
+get_filename_component(CLBLAS_INCLUDE_DIRS ${CMAKE_CURRENT_LIST_DIR}/@reldir@/include ABSOLUTE)
+set(CLBLAS_LIBRARIES clBLAS)

src/clBLASConfigVersion.cmake.in

@@ -0,0 +1,46 @@
+# This is a basic version file for the Config-mode of find_package().
+# It is used by write_basic_package_version_file() as input file for configure_file()
+# to create a version-file which can be installed along a config.cmake file.
+#
+# The created file sets PACKAGE_VERSION_EXACT if the current version string and
+# the requested version string are exactly the same and it sets
+# PACKAGE_VERSION_COMPATIBLE if the current version is >= requested version,
+# but only if the requested major version is the same as the current one.
+# The variable CLBLAS_VERSION must be set before calling configure_file().
+
+
+set(PACKAGE_VERSION "@CLBLAS_VERSION@")
+
+if("${PACKAGE_VERSION}" VERSION_LESS "${PACKAGE_FIND_VERSION}" )
+  set(PACKAGE_VERSION_COMPATIBLE FALSE)
+else()
+
+  if("@CLBLAS_VERSION@" MATCHES "^([0-9]+)\\.")
+    set(CLBLAS_VERSION_MAJOR "${CMAKE_MATCH_1}")
+  else()
+    set(CLBLAS_VERSION_MAJOR "@CLBLAS_VERSION@")
+  endif()
+
+  if("${PACKAGE_FIND_VERSION_MAJOR}" STREQUAL "${CLBLAS_VERSION_MAJOR}")
+    set(PACKAGE_VERSION_COMPATIBLE TRUE)
+  else()
+    set(PACKAGE_VERSION_COMPATIBLE FALSE)
+  endif()
+
+  if( "${PACKAGE_FIND_VERSION}" STREQUAL "${PACKAGE_VERSION}")
+      set(PACKAGE_VERSION_EXACT TRUE)
+  endif()
+endif()
+
+
+# if the installed or the using project don't have CMAKE_SIZEOF_VOID_P set, ignore it:
+if("${CMAKE_SIZEOF_VOID_P}"  STREQUAL ""  OR "@CMAKE_SIZEOF_VOID_P@" STREQUAL "")
+   return()
+endif()
+
+# check that the installed version has the same 32/64bit-ness as the one which is currently searching:
+if(NOT "${CMAKE_SIZEOF_VOID_P}" STREQUAL "@CMAKE_SIZEOF_VOID_P@")
+  math(EXPR installedBits "@CMAKE_SIZEOF_VOID_P@ * 8")
+  set(PACKAGE_VERSION "${PACKAGE_VERSION} (${installedBits}bit)")
+  set(PACKAGE_VERSION_UNSUITABLE TRUE)
+endif()

src/client/CMakeLists.txt

@@ -52,7 +52,9 @@ include_directories(
 
 add_executable(client ${CLIENT_SRC} ${CLIENT_HEADER})
 target_link_libraries(client ${Boost_LIBRARIES} ${OPENCL_LIBRARIES} clBLAS)
-set_target_properties( client PROPERTIES RUNTIME_OUTPUT_DIRECTORY "${PROJECT_BINARY_DIR}/staging" )
+set_target_properties( client PROPERTIES
+  RUNTIME_OUTPUT_DIRECTORY "${PROJECT_BINARY_DIR}/staging"
+  OUTPUT_NAME clBLAS-client )
 
 add_executable(testPerfWrapper ${WRAPPER_SRC})
 target_link_libraries(testPerfWrapper ${Boost_LIBRARIES})

src/include/granulation.h

@@ -46,6 +46,8 @@ typedef struct PGranularity {
     unsigned int wfSize;
     /** Record number of work-groups spawned */
     unsigned int numWGSpawned[2];
+	/** max number of work group size */
+	unsigned int maxWorkGroupSize;
 } PGranularity;
 
 /**

src/library/CMakeLists.txt

@@ -279,7 +279,7 @@ endif()
 
 include( ExternalProject )
 ExternalProject_Add( tplgen
-    URL "${CMAKE_SOURCE_DIR}/library/tools/tplgen"
+    URL "${PROJECT_SOURCE_DIR}/library/tools/tplgen"
     INSTALL_COMMAND ""
 )
 
@@ -294,7 +294,7 @@ endif()
 
 add_custom_target( GENERATE_CLT
     COMMAND ${tplgenBinaryDir}/tplgen -o ${clBLAS_BINARY_DIR}/include ${SRC_CL_TEMPLATES}
-    WORKING_DIRECTORY ${CMAKE_SOURCE_DIR}/library/blas/gens/clTemplates
+    WORKING_DIRECTORY ${PROJECT_SOURCE_DIR}/library/blas/gens/clTemplates
 )
 
 add_dependencies( GENERATE_CLT tplgen )
@@ -307,7 +307,7 @@ if( CMAKE_COMPILER_IS_GNUCC )
              DESTINATION lib${SUFFIX_LIB}/pkgconfig )
 endif( )
 
-add_library(clBLAS SHARED ${CLBLAS_SOURCES} ${GLOBAL_HEADERS} ${SRC_BLAS_HEADERS} ${SRC_BLAS_GENS_HEADERS})
+add_library(clBLAS ${CLBLAS_SOURCES} ${GLOBAL_HEADERS} ${SRC_BLAS_HEADERS} ${SRC_BLAS_GENS_HEADERS})
 add_dependencies(clBLAS GENERATE_CLT)
 set_target_properties(clBLAS PROPERTIES VERSION ${clBLAS_VERSION})
 set_target_properties(clBLAS PROPERTIES SOVERSION ${clBLAS_SOVERSION})
@@ -316,8 +316,9 @@ target_link_libraries(clBLAS ${OPENCL_LIBRARIES} ${MATH_LIBRARY})
 
 # CPack configuration; include the executable into the package
 install( TARGETS clBLAS
-		RUNTIME DESTINATION bin${SUFFIX_BIN}
-		LIBRARY DESTINATION lib${SUFFIX_LIB}
+                EXPORT Library
+                RUNTIME DESTINATION bin${SUFFIX_BIN}
+                LIBRARY DESTINATION lib${SUFFIX_LIB}
 		ARCHIVE DESTINATION lib${SUFFIX_LIB}/import
 		)
 

src/library/blas/generic/common.c

@@ -527,7 +527,9 @@ setupBuildOpts(
     opts[0] = '\0';
 
 #if !defined NDEBUG
-    addBuildOpt(opts, BUILD_OPTS_MAXLEN, "-g");
+    // Nvidia runtime does not appear to support the -g flag, at least in their OpenCL v1.1 runtime
+    if( target.ident.vendor != VENDOR_NVIDIA )
+        addBuildOpt( opts, BUILD_OPTS_MAXLEN, "-g" );
 #endif  /* NDEBUG */
 
     if (target.ident.vendor == VENDOR_NVIDIA &&

src/library/blas/gens/blas_kgen.h

@@ -52,6 +52,7 @@
 #include "tile.h"
 #include "fetch.h"
 
+
 #define BLAS_KGEN_FORMAT 1
 
 #define genInternalLoopEnd(ctx) kgenEndBranch(ctx, NULL)
@@ -539,6 +540,18 @@ sprintfComplexMulUpdate(
     bool conjB,
     TileMulCore core);
 
+void
+sprintfComplexMulUpdate_syr2k_beta0(
+    Kstring *expr,
+    const Kstring *dst,
+    const Kstring *a,
+    const Kstring *b,
+    const Kstring *c,
+    bool isDouble,
+    bool conjA,
+    bool conjB,
+    TileMulCore core);
+
 /**
  * @brief Sprintf expression of fast scalar mad
  *
@@ -892,4 +905,6 @@ checkGenRestoreTailCoords(
 UpdateResultFlags
 tailStatusToUpresFlags(TailStatus status);
 
+
+
 #endif /* BLAS_KGEN_H_ */

src/library/blas/gens/fetch.c

@@ -359,7 +359,7 @@ sprintfNormalizedBaseCoord(
         kstrcpy(kstr, name);
     }
     else {
-        ksprintf(kstr, "(%s >> %d)", name, shift);
+        ksprintf(kstr, "(uint)(%s >> %d)", name, shift);
     }
 }
 
@@ -515,7 +515,7 @@ sprintfLeadingDimension(Kstring *ld, const FetchContext *fctx)
 
         shift = findHighestSetBit(fctx->physTile.vecLen);
         if (shift != 0) {
-            ksprintf(ld, "(%s >> %d)", varName, shift);
+            ksprintf(ld, "(uint)(%s >> %d)", varName, shift);
             done = true;
         }
     }
@@ -564,10 +564,10 @@ sprintfGboundK(Kstring *kstr, const FetchContext *fctx)
     }
     else {
         if (fctx->addrMode & FETCH_ADDR_TAILK_PADD) {
-            ksprintf(kstr, "((%s + %u) >> %d)", varK, vecLen - 1, shift);
+            ksprintf(kstr, "(uint)((%s + %u) >> %d)", varK, vecLen - 1, shift);
         }
         else {
-            ksprintf(kstr, "(%s >> %d)", varK, shift);
+            ksprintf(kstr, "(uint)(%s >> %d)", varK, shift);
         }
     }
 }

src/library/blas/gens/gemm.c

@@ -560,7 +560,7 @@ blockGen(
         kgenAddStmt(ctx, tmp);
     }
     else {
-        sprintf(globalIdB, "get_global_id(%d)", 1-i);
+        sprintf(globalIdB, "(uint)get_global_id(%d)", 1-i);
     }
 
     if (!(isColMajA || isColMajB)) {
@@ -758,7 +758,7 @@ subgGen(
     vecLenA = gset.tileA.vecLen;
 
     // channel offset based coordinate
-    ksprintf(&exprK, "( get_group_id(0)*%lu + k )", staggered/vecLenA*vecLenA);
+    ksprintf(&exprK, "( (uint)(get_group_id(0))*%lu + k )", staggered/vecLenA*vecLenA);
 
     // starting code generation--------------------------------------------------
     pCtx = createKgenContext(pBuf, buflen, true);
@@ -1104,6 +1104,8 @@ blockCheckCalcDecomp(
     int check)
 {
     bool ret = true;
+	bool ret_multiple = false;
+	int i;
 
     DUMMY_ARG_USAGE(subdimsNum);
 
@@ -1114,7 +1116,12 @@ blockCheckCalcDecomp(
         minSize = (dtype == TYPE_COMPLEX_DOUBLE) ? 1 : 2;
         ret = decompSanityCheck(subdims, minSize, maxSize, 24, dtype, true);
         ret = ret && (subdims[0].bwidth == subdims[1].bwidth);
-        ret = ret && (pgran->wgSize[0] * pgran->wgSize[1] == 64);
+		for(i = 0; i < ( (pgran->maxWorkGroupSize) / (pgran->wfSize) ); i++)
+		{
+			// returns true if wgSize[0] * wgSize[1] is multiples of the 64 but not bigger than maxWorkGroupSize
+			ret_multiple = ret_multiple || ( pgran->wgSize[0] * pgran->wgSize[1] == pgran->wfSize * (i + 1) );
+		}
+		ret = ret && ret_multiple;
     }
     else {
         calcPgranDedicated(pgran, subdims, 1, 3);