Adding support for building for AMD on Windows

Author: rwfsmith

CMakeLists.txt

@@ -263,6 +263,8 @@ endif()
 if(WIN32)
     # Export all symbols
     set(CMAKE_WINDOWS_EXPORT_ALL_SYMBOLS ON)
+    # Prevent Windows SDK min/max macros from conflicting with std::min/std::max
+    add_compile_definitions(NOMINMAX)
 endif()
 
 if(MSVC)
@@ -330,14 +332,22 @@ if(BUILD_HIP)
     find_package_and_print_version(hipsparse REQUIRED)
 
     ## hacky way of excluding hip::amdhip64 (with it linked many tests unexpectedly fail e.g. adam8bit because of inaccuracies)
-    set_target_properties(hip::host PROPERTIES INTERFACE_LINK_LIBRARIES "")
-    set_target_properties(hip-lang::host PROPERTIES INTERFACE_LINK_LIBRARIES "")
-    set(CMAKE_HIP_IMPLICIT_LINK_LIBRARIES "")
+    ## On Windows, we need to link amdhip64 explicitly
+    if(NOT WIN32)
+        set_target_properties(hip::host PROPERTIES INTERFACE_LINK_LIBRARIES "")
+        set_target_properties(hip-lang::host PROPERTIES INTERFACE_LINK_LIBRARIES "")
+        set(CMAKE_HIP_IMPLICIT_LINK_LIBRARIES "")
+    endif()
 
     target_include_directories(bitsandbytes PRIVATE ${CMAKE_SOURCE_DIR} ${CMAKE_SOURCE_DIR}/include ${ROCM_PATH}/include /include)
     target_link_directories(bitsandbytes PRIVATE ${ROCM_PATH}/lib /lib)
     target_link_libraries(bitsandbytes PUBLIC roc::hipblas hip::hiprand roc::hipsparse)
 
+    # On Windows, link the HIP runtime and rocblas directly
+    if(WIN32)
+        target_link_libraries(bitsandbytes PUBLIC amdhip64 rocblas)
+    endif()
+
     target_compile_definitions(bitsandbytes PUBLIC BNB_USE_HIP)
     set_source_files_properties(${HIP_FILES} PROPERTIES LANGUAGE HIP)
     set_target_properties(bitsandbytes PROPERTIES LINKER_LANGUAGE CXX)

bitsandbytes/cuda_specs.py

@@ -1,6 +1,7 @@
 import dataclasses
 from functools import lru_cache
 import logging
+import os
 import re
 import subprocess
 from typing import Optional
@@ -83,10 +84,21 @@ def get_rocm_gpu_arch() -> str:
     logger = logging.getLogger(__name__)
     try:
         if torch.version.hip:
-            result = subprocess.run(["rocminfo"], capture_output=True, text=True)
-            match = re.search(r"Name:\s+gfx([a-zA-Z\d]+)", result.stdout)
+            # On Windows, use hipinfo.exe; on Linux, use rocminfo
+            if os.name == "nt":
+                cmd = ["hipinfo.exe"]
+                arch_pattern = r"gcnArchName:\s+(gfx[a-zA-Z\d]+)"
+            else:
+                cmd = ["rocminfo"]
+                arch_pattern = r"Name:\s+gfx([a-zA-Z\d]+)"
+
+            result = subprocess.run(cmd, capture_output=True, text=True)
+            match = re.search(arch_pattern, result.stdout)
             if match:
-                return "gfx" + match.group(1)
+                if os.name == "nt":
+                    return match.group(1)
+                else:
+                    return "gfx" + match.group(1)
             else:
                 return "unknown"
         else:
@@ -107,8 +119,17 @@ def get_rocm_warpsize() -> int:
     logger = logging.getLogger(__name__)
     try:
         if torch.version.hip:
-            result = subprocess.run(["rocminfo"], capture_output=True, text=True)
-            match = re.search(r"Wavefront Size:\s+([0-9]{2})\(0x[0-9]{2}\)", result.stdout)
+            # On Windows, use hipinfo.exe; on Linux, use rocminfo
+            if os.name == "nt":
+                cmd = ["hipinfo.exe"]
+                # hipinfo.exe output format: "warpSize: 32" or "warpSize: 64"
+                warp_pattern = r"warpSize:\s+(\d+)"
+            else:
+                cmd = ["rocminfo"]
+                warp_pattern = r"Wavefront Size:\s+([0-9]{2})\(0x[0-9]{2}\)"
+
+            result = subprocess.run(cmd, capture_output=True, text=True)
+            match = re.search(warp_pattern, result.stdout)
             if match:
                 return int(match.group(1))
             else:

csrc/ops.cuh

@@ -10,6 +10,13 @@
 #include <cstdint>
 #include <iostream>
 #include <stdio.h>
+#ifdef _WIN32
+#include <io.h>
+#include <process.h>
+#include <windows.h>
+#else
+#include <unistd.h>
+#endif
 
 #include <common.h>
 #include <cublasLt.h>

csrc/ops_hip.cuh

@@ -11,7 +11,17 @@
 #include <cstdint>
 #include <iostream>
 #include <stdio.h>
+
+#ifdef _WIN32
+#ifndef NOMINMAX
+#define NOMINMAX
+#endif
+#include <io.h>
+#include <process.h>
+#include <windows.h>
+#else
 #include <unistd.h>
+#endif
 
 #include <common.h>
 #include <functional>