Windows gfx1201

CMakeLists.txt

@@ -62,11 +62,19 @@ else()
 endif()
 
 if (NOT DEFINED ENV{CXX} AND NOT CMAKE_CXX_COMPILER)
-  set(CMAKE_CXX_COMPILER "${rocm_bin}/amdclang++")
+  if(WIN32)
+    set(CMAKE_CXX_COMPILER "${rocm_bin}/clang++.exe")
+  else()
+    set(CMAKE_CXX_COMPILER "${rocm_bin}/amdclang++")
+  endif()
 endif()
 
 if (NOT DEFINED ENV{CC} AND NOT CMAKE_C_COMPILER)
-  set(CMAKE_C_COMPILER "${rocm_bin}/amdclang")
+  if(WIN32)
+    set(CMAKE_C_COMPILER "${rocm_bin}/clang")
+  else()
+    set(CMAKE_C_COMPILER "${rocm_bin}/amdclang")
+  endif()
 endif()
 
 # TODO: move FC and CXX and CC compiler vars above to new toolchain-linux.cmake (Fortran for clients)
@@ -116,7 +124,11 @@ option(Tensile_SEPARATE_ARCHITECTURES "Tensile to use GPU architecture specific
 option(Tensile_NO_LAZY_LIBRARY_LOADING "Diasble loading kernels on demand?" OFF)
 # For roctx
 include(CMakeDependentOption)
-cmake_dependent_option(HIPBLASLT_ENABLE_MARKER "Enable roctx marker in hipBLASLt" ON "BUILD_SHARED_LIBS" OFF)
+if(WIN32)
+  cmake_dependent_option(HIPBLASLT_ENABLE_MARKER "Disable roctx marker in hipBLASLt - roctracer does not support on Windows" OFF "BUILD_SHARED_LIBS" OFF)
+else()
+  cmake_dependent_option(HIPBLASLT_ENABLE_MARKER "Enable roctx marker in hipBLASLt" ON "BUILD_SHARED_LIBS" OFF)
+endif()
 
 if(BUILD_CODE_COVERAGE)
   add_compile_options(-fprofile-arcs -ftest-coverage)
@@ -195,7 +207,11 @@ else()
 
     set( Tensile_LOGIC "asm_full" CACHE STRING "Tensile to use which logic?")
     set( Tensile_CODE_OBJECT_VERSION "4" CACHE STRING "Tensile code_object_version")
-    set( Tensile_COMPILER "amdclang++" CACHE STRING "Tensile compiler")
+    if(WIN32)
+      set( Tensile_COMPILER "clang++.exe" CACHE STRING "Tensile compiler")
+    else()
+      set( Tensile_COMPILER "amdclang++" CACHE STRING "Tensile compiler")
+    endif()
     set( Tensile_LIBRARY_FORMAT "msgpack" CACHE STRING "Tensile library format")
     set( Tensile_CPU_THREADS "" CACHE STRING "Number of threads for Tensile parallel build")
 
@@ -255,9 +271,17 @@ else()
 endif()
 
 if( LEGACY_HIPBLAS_DIRECT )
-  find_package( hipblas REQUIRED CONFIG PATHS ${HIP_DIR} ${ROCM_PATH} /opt/rocm)
+  if (NOT WIN32)
+    find_package( hipblas REQUIRED CONFIG PATHS ${HIP_DIR} ${ROCM_PATH} /opt/rocm)
+  else()
+    find_package( hipblas REQUIRED CONFIG PATHS ${HIP_DIR})
+  endif()
 else()
-  find_package( hipblas-common REQUIRED CONFIG PATHS ${HIP_DIR} ${ROCM_PATH} /opt/rocm)
+    if (NOT WIN32)
+      find_package( hipblas-common REQUIRED CONFIG PATHS ${HIP_DIR} ${ROCM_PATH} /opt/rocm)
+    else()
+      find_package( hipblas-common REQUIRED CONFIG PATHS ${HIP_DIR})
+    endif()
 endif()
 
 if(HIPBLASLT_ENABLE_MARKER)
@@ -359,12 +383,27 @@ if(BUILD_DOCS)
   add_subdirectory(docs)
 endif()
 
+# The following code is setting variables to control the behavior of CPack to generate our
+if( WIN32 )
+    set( CPACK_SOURCE_GENERATOR "ZIP" )
+    set( CPACK_GENERATOR "ZIP" )
+endif( )
+
 # Package specific CPACK vars
 set( CPACK_RESOURCE_FILE_LICENSE "${CMAKE_CURRENT_SOURCE_DIR}/LICENSE.md" )
 set( CPACK_RPM_PACKAGE_LICENSE "MIT")
 
-if( NOT CPACK_PACKAGING_INSTALL_PREFIX )
-    set( CPACK_PACKAGING_INSTALL_PREFIX "${CMAKE_INSTALL_PREFIX}" )
+if (WIN32)
+  SET( CMAKE_INSTALL_PREFIX "C:/hipSDK" CACHE PATH "Install path" FORCE )
+  SET( INSTALL_PREFIX "C:/hipSDK" )
+  SET( CPACK_SET_DESTDIR FALSE )
+  SET( CPACK_PACKAGE_INSTALL_DIRECTORY "C:/hipSDK" )
+  SET( CPACK_PACKAGING_INSTALL_PREFIX "" )
+  set( CPACK_INCLUDE_TOPLEVEL_DIRECTORY OFF )
+else()
+  if( NOT CPACK_PACKAGING_INSTALL_PREFIX )
+      set( CPACK_PACKAGING_INSTALL_PREFIX "${CMAKE_INSTALL_PREFIX}" )
+  endif()
 endif()
 
 set( CPACK_RPM_EXCLUDE_FROM_AUTO_FILELIST_ADDITION "\${CPACK_PACKAGING_INSTALL_PREFIX}" "\${CPACK_PACKAGING_INSTALL_PREFIX}/include" "\${CPACK_PACKAGING_INSTALL_PREFIX}/lib" )

clients/CMakeLists.txt

@@ -51,7 +51,10 @@ find_package(OpenMP)
 
 if (TARGET OpenMP::OpenMP_CXX)
   set( COMMON_LINK_LIBS "OpenMP::OpenMP_CXX")
-  list( APPEND COMMON_LINK_LIBS "-L${HIP_CLANG_ROOT}/lib;-Wl,-rpath=${HIP_CLANG_ROOT}/lib")
+  list( APPEND COMMON_LINK_LIBS "-L\"${HIP_CLANG_ROOT}/lib\"")
+  if (NOT WIN32)
+    list( APPEND COMMON_LINK_LIBS "-Wl,-rpath=${HIP_CLANG_ROOT}/lib")
+  endif()
 endif()
 
 if (TARGET Threads::Threads)
@@ -87,12 +90,25 @@ endif( )
 if( BUILD_CLIENTS_BENCHMARKS OR BUILD_CLIENTS_TESTS)
 
   # Linking lapack library requires fortran flags
-  find_package( cblas REQUIRED CONFIG )
-  if(${BLIS_FOUND})
-    set( BLAS_LIBRARY ${BLIS_LIB} )
-    set( BLIS_CPP ../common/blis_interface.cpp )
-  else()
-    set( BLAS_LIBRARY "blas" )
+  if ( NOT WIN32 )
+    find_package( cblas REQUIRED CONFIG )
+    if(${BLIS_FOUND})
+      set( BLAS_LIBRARY ${BLIS_LIB} )
+      set( BLIS_CPP ../common/blis_interface.cpp )
+    else()
+      set( BLAS_LIBRARY "blas" )
+    endif()
+  else() #WIN32
+    set( BLAS_INCLUDE_DIR ${OPENBLAS_DIR}/include CACHE PATH "OpenBLAS library include path" )
+    find_library( BLAS_LIBRARY libopenblas
+                  PATHS ${OPENBLAS_DIR}/lib
+                  NO_DEFAULT_PATH
+                )
+    if (NOT BLAS_LIBRARY)
+      find_package( OPENBLAS CONFIG REQUIRED )
+      set( BLAS_LIBRARY OpenBLAS::OpenBLAS )
+      set( BLAS_INCLUDE_DIR "" )
+    endif()
   endif()
 
   # Find the package ROCmSMI

clients/benchmarks/CMakeLists.txt

@@ -88,6 +88,9 @@ if (NOT WIN32)
     list( APPEND COMMON_LINK_LIBS "-lflang -lflangrti") # for lapack
   endif()
 else()
+  find_package(lapack REQUIRED)
+  message("LAPACK: ${LAPACK_LIBRARIES}")
+  target_link_libraries(hipblaslt-bench PRIVATE ${LAPACK_LIBRARIES})
   list( APPEND COMMON_LINK_LIBS "libomp")
 endif()
 target_link_libraries( hipblaslt-bench PRIVATE ${COMMON_LINK_LIBS} )

clients/common/blis_interface.cpp

@@ -29,7 +29,7 @@
 
 void setup_blis()
 {
-#ifndef WIN32
+#ifndef _WIN32
     bli_init();
 #endif
 }

clients/common/hipblaslt_init_device.cpp

@@ -116,7 +116,7 @@ __device__ int8_t random_hpl(size_t idx)
 }
 
 template <typename T>
-void hipblaslt_init_device(ABC                      abc,
+void hipblaslt_init_device(ABC_dims                      abc,
                            hipblaslt_initialization init,
                            bool                     is_nan,
                            T*                       A,
@@ -140,11 +140,11 @@ void hipblaslt_init_device(ABC                      abc,
         switch(init)
         {
         case hipblaslt_initialization::rand_int:
-            if(abc == ABC::A || abc == ABC::C)
+            if(abc == ABC_dims::A || abc == ABC_dims::C)
                 fill_batch(A, M, N, lda, stride, batch_count, [](size_t idx) -> T {
                     return random_int<T>(idx);
                 });
-            else if(abc == ABC::B)
+            else if(abc == ABC_dims::B)
             {
                 stride = std::max(lda * N, stride);
                 fill_batch(A, M, N, lda, stride, batch_count, [stride, lda](size_t idx) -> T {
@@ -158,14 +158,14 @@ void hipblaslt_init_device(ABC                      abc,
             break;
         case hipblaslt_initialization::trig_float:
             stride = std::max(lda * N, stride);
-            if(abc == ABC::A || abc == ABC::C)
+            if(abc == ABC_dims::A || abc == ABC_dims::C)
                 fill_batch(A, M, N, lda, stride, batch_count, [M, N, stride, lda](size_t idx) -> T {
                     auto b = idx / stride;
                     auto j = (idx - b * stride) / lda;
                     auto i = (idx - b * stride) - j * lda;
                     return T(sin(double(i + j*M + b*M*N)));
                 });
-            else if(abc == ABC::B)
+            else if(abc == ABC_dims::B)
                 fill_batch(A, M, N, lda, stride, batch_count, [M, N, stride, lda](size_t idx) -> T {
                     auto b = idx / stride;
                     auto j = (idx - b * stride) / lda;
@@ -179,15 +179,15 @@ void hipblaslt_init_device(ABC                      abc,
             });
             break;
         case hipblaslt_initialization::special:
-            if(abc == ABC::A)
+            if(abc == ABC_dims::A)
                 fill_batch(A, M, N, lda, stride, batch_count, [](size_t idx) -> T {
                     return T(hipblasLtHalf(65280.0));
                 });
-            else if(abc == ABC::B)
+            else if(abc == ABC_dims::B)
                 fill_batch(A, M, N, lda, stride, batch_count, [](size_t idx) -> T {
                     return T(hipblasLtHalf(0.0000607967376708984375));
                 });
-            else if(abc == ABC::C)
+            else if(abc == ABC_dims::C)
                 fill_batch(A, M, N, lda, stride, batch_count, [](size_t idx) -> T {
                     return T(pseudo_random_device(idx) % 10 + 1.f);
                 });
@@ -213,7 +213,7 @@ void hipblaslt_init_device(ABC                      abc,
     }
 }
 
-void hipblaslt_init_device(ABC                      abc,
+void hipblaslt_init_device(ABC_dims                      abc,
                            hipblaslt_initialization init,
                            bool                     is_nan,
                            void*                    A,

clients/common/hipblaslt_parse_data.cpp

@@ -43,7 +43,7 @@ static std::string hipblaslt_parse_yaml(const std::string& yaml)
                + "hipblaslt_template.yaml -o " + tmp + " " + yaml;
     hipblaslt_cerr << cmd << std::endl;
 
-#ifdef WIN32
+#ifdef _WIN32
     int status = std::system(cmd.c_str());
     if(status == -1)
         exit(EXIT_FAILURE);

clients/common/utility.cpp

@@ -32,7 +32,12 @@
 #include <stdexcept>
 #include <stdlib.h>
 
+#ifdef _WIN32
+#include <windows.h>
+#include <libloaderapi.h>
+#else
 #include <fcntl.h>
+#endif
 
 #include "Tensile/Source/client/include/Utility.hpp"
 
@@ -50,6 +55,27 @@ namespace fs = std::experimental::filesystem;
 // Return path of this executable
 std::string hipblaslt_exepath()
 {
+#ifdef _WIN32
+    std::vector<TCHAR> result(MAX_PATH + 1);
+    // Ensure result is large enough to accommodate the path
+    DWORD length = 0;
+    for(;;)
+    {
+        length = GetModuleFileNameA(nullptr, result.data(), result.size());
+        if(length < result.size() - 1)
+        {
+            result.resize(length + 1);
+            break;
+        }
+        result.resize(result.size() * 2);
+    }
+
+    fs::path exepath(result.begin(), result.end());
+    exepath = exepath.remove_filename();
+    // Add trailing "/" to exepath if required
+    exepath += exepath.empty() ? "" : "/";
+    return exepath.string();
+#else
     std::string pathstr;
     char*       path = realpath("/proc/self/exe", 0);
     if(path)
@@ -63,12 +89,26 @@ std::string hipblaslt_exepath()
         free(path);
     }
     return pathstr;
+#endif
 }
 
 /* ============================================================================================ */
 // Temp directory rooted random path
 std::string hipblaslt_tempname()
 {
+#ifdef _WIN32
+    // Generate "/tmp/rocblas-XXXXXX" like file name
+    const std::string alphanum     = "0123456789ABCDEFGHIJKLMNOPQRSTUVWXYZabcdefghijklmnopqrstuv";
+    int               stringlength = alphanum.length() - 1;
+    std::string       uniquestr    = "hipblaslt-";
+
+    for(auto n : {0, 1, 2, 3, 4, 5})
+        uniquestr += alphanum.at(rand() % stringlength);
+
+    fs::path tmpname = fs::temp_directory_path() / uniquestr;
+
+    return tmpname.string();
+#else
     char tmp[] = "/tmp/hipblaslt-XXXXXX";
     int  fd    = mkostemp(tmp, O_CLOEXEC);
     if(fd == -1)
@@ -78,6 +118,7 @@ std::string hipblaslt_tempname()
     }
 
     return std::string(tmp);
+#endif
 }
 
 /* ============================================================================================ */
@@ -239,7 +280,11 @@ hipblaslt_local_handle::hipblaslt_local_handle(const Arguments& arg)
         if(sol_selec_env)
             m_sol_selec_saved_status = std::string(sol_selec_env);
         m_sol_selec_env_set = true;
+#ifdef _WIN32
+        _putenv_s("TENSILE_SOLUTION_SELECTION_METHOD", std::to_string(arg.tensile_solution_selection_method).c_str());
+#else
         setenv("TENSILE_SOLUTION_SELECTION_METHOD", std::to_string(arg.tensile_solution_selection_method).c_str(), true);
+#endif    
     }
     // memory guard control, with multi-threading should not change values across threads
     d_vector_set_pad_length(arg.pad);
@@ -249,7 +294,11 @@ hipblaslt_local_handle::~hipblaslt_local_handle()
 {
     if(m_sol_selec_env_set)
     {
-        setenv("TENSILE_SOLUTION_SELECTION_METHOD", m_sol_selec_saved_status.c_str(), true);
+#ifdef _WIN32
+        _putenv_s("TENSILE_SOLUTION_SELECTION_METHOD", m_sol_selec_saved_status.c_str());
+#else
+setenv("TENSILE_SOLUTION_SELECTION_METHOD", m_sol_selec_saved_status.c_str(), true);
+#endif  
     }
     hipblasLtDestroy(m_handle);
 }