Adding support for building for AMD on Windows (#1846)

* Adding support for building for AMD on Windows

* Simplify building AMD on Windows by auto-detecting compiler and GPU arch

* Standardize Windows detection to use platform.system()

* Made some of the suggested changes

* Reverted return to previous behavior

* Changed to use find_path for lib path

Author: rwfsmith

CMakeLists.txt

@@ -16,6 +16,32 @@
 #                  libbitsandbytes_rocm70.so even if the system has ROCm 7.2.
 cmake_minimum_required(VERSION 3.22.1)
 
+# On Windows with HIP backend, auto-detect compilers from ROCM_PATH before project()
+if(WIN32 AND COMPUTE_BACKEND STREQUAL "hip")
+    if(DEFINED ENV{ROCM_PATH})
+        set(ROCM_PATH $ENV{ROCM_PATH})
+    endif()
+    if(ROCM_PATH AND NOT DEFINED CMAKE_CXX_COMPILER)
+        set(CMAKE_CXX_COMPILER "${ROCM_PATH}/lib/llvm/bin/clang++.exe")
+    endif()
+    if(ROCM_PATH AND NOT DEFINED CMAKE_HIP_COMPILER)
+        set(CMAKE_HIP_COMPILER "${ROCM_PATH}/lib/llvm/bin/clang++.exe")
+    endif()
+    # On Windows, the HIP compiler needs explicit paths to find device libraries.
+    if(ROCM_PATH)
+        find_path(ROCM_DEVICE_LIB_PATH
+            NAMES oclc_abi_version_400.bc ocml.bc
+            PATHS "${ROCM_PATH}/amdgcn/bitcode"
+                  "${ROCM_PATH}/lib/llvm/amdgcn/bitcode"
+            NO_DEFAULT_PATH
+        )
+        set(CMAKE_HIP_FLAGS "--rocm-path=${ROCM_PATH}")
+        if(ROCM_DEVICE_LIB_PATH)
+            set(CMAKE_HIP_FLAGS "${CMAKE_HIP_FLAGS} --rocm-device-lib-path=${ROCM_DEVICE_LIB_PATH}")
+        endif()
+    endif()
+endif()
+
 project(bitsandbytes LANGUAGES CXX)
 
 # If run without specifying a build type, default to using the Release configuration:
@@ -204,17 +230,18 @@ if(BUILD_CUDA)
     string(APPEND BNB_OUTPUT_NAME "_cuda${CUDA_VERSION_SHORT}")
     add_compile_definitions(BUILD_CUDA)
 elseif(BUILD_HIP)
-    enable_language(HIP)
-    message(STATUS "HIP Compiler: ${CMAKE_HIP_COMPILER}")
+    # Set target architectures before enable_language(HIP), which would otherwise
+    # auto-detect a single GPU and override the defaults.
     if(DEFINED BNB_ROCM_ARCH)
       set(CMAKE_HIP_ARCHITECTURES ${BNB_ROCM_ARCH})
-    else()
-      if (NOT AMDGPU_TARGETS AND NOT CMAKE_HIP_ARCHITECTURES)
-        set(CMAKE_HIP_ARCHITECTURES "gfx90a;gfx942;gfx1100;gfx1101;gfx1150;gfx1151;gfx1200;gfx1201")
-      elseif (AMDGPU_TARGETS AND NOT CMAKE_HIP_ARCHITECTURES)
-        set(CMAKE_HIP_ARCHITECTURES ${AMDGPU_TARGETS})
-      endif()
+    elseif(AMDGPU_TARGETS AND NOT CMAKE_HIP_ARCHITECTURES)
+      set(CMAKE_HIP_ARCHITECTURES ${AMDGPU_TARGETS})
+    elseif(NOT CMAKE_HIP_ARCHITECTURES)
+      set(CMAKE_HIP_ARCHITECTURES "gfx90a;gfx942;gfx1100;gfx1101;gfx1150;gfx1151;gfx1200;gfx1201")
     endif()
+
+    enable_language(HIP)
+    message(STATUS "HIP Compiler: ${CMAKE_HIP_COMPILER}")
     message(STATUS "HIP Targets: ${CMAKE_HIP_ARCHITECTURES}")
 
     list(APPEND SRC_FILES ${HIP_FILES})
@@ -275,6 +302,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)
@@ -327,10 +356,11 @@ if(BUILD_CUDA)
     )
 endif()
 if(BUILD_HIP)
-    if(NOT DEFINED ENV{ROCM_PATH})
-      set(ROCM_PATH /opt/rocm)
-    else()
+    # Determine ROCM_PATH from environment variable, fallback to /opt/rocm on Linux
+    if(DEFINED ENV{ROCM_PATH})
       set(ROCM_PATH $ENV{ROCM_PATH})
+    else()
+      set(ROCM_PATH /opt/rocm)
     endif()
     list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH})
     macro(find_package_and_print_version PACKAGE_NAME)
@@ -342,14 +372,23 @@ 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, rocblas is not pulled in transitively by roc::hipblas
+    # and is needed because ops_hip.cuh uses rocblas_handle directly.
+    if(WIN32)
+        target_link_libraries(bitsandbytes PUBLIC 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 platform
 import re
 import subprocess
 from typing import Optional
@@ -83,8 +84,16 @@ 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 platform.system() == "Windows":
+                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)
             else:
@@ -107,8 +116,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 platform.system() == "Windows":
+                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_hip.cuh

@@ -11,7 +11,14 @@
 #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 <functional>