[Proton] Enable build on windows (#5533)

anmyachev · web-flow · commit dcb8ff06ad00 · 2025-11-26T11:26:41.000+01:00
Signed-off-by: Anatoly Myachev &lt;anatoly.myachev@intel.com&gt;
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -16,15 +16,8 @@ include(CTest)
 list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/cmake")
 
 # Options
-if(WIN32)
-  set(DEFAULT_BUILD_PROTON OFF)
-else()
-  set(DEFAULT_BUILD_PROTON ON)
-endif()
-
-# Define the option with the determined default value
-option(TRITON_BUILD_PROTON "Build the Triton Proton profiler" ${DEFAULT_BUILD_PROTON})
 option(TRITON_BUILD_PYTHON_MODULE "Build Python Triton bindings" OFF)
+option(TRITON_BUILD_PROTON "Build the Triton Proton profiler" ON)
 option(TRITON_BUILD_UT "Build C++ Triton Unit Tests" ON)
 option(TRITON_BUILD_WITH_CCACHE "Build with ccache (if available)" ON)
 set(TRITON_CODEGEN_BACKENDS "" CACHE STRING "Enable different codegen backends")
diff --git a/lib/Instrumentation/PrintLoadStoreMemSpaces.cpp b/lib/Instrumentation/PrintLoadStoreMemSpaces.cpp
@@ -101,6 +101,6 @@ extern "C" LLVM_ATTRIBUTE_WEAK PassPluginLibraryInfo llvmGetPassPluginInfo() {
   return getPassPluginInfo();
 }
 
-#ifdef WIN32
+#if defined(_WIN32)
 #pragma comment(linker, "/export:llvmGetPassPluginInfo")
 #endif
diff --git a/test/lib/Instrumentation/GPUHello.cpp b/test/lib/Instrumentation/GPUHello.cpp
@@ -75,6 +75,6 @@ llvmGetPassPluginInfo() {
   return getPassPluginInfo();
 }
 
-#ifdef WIN32
+#if defined(_WIN32)
 #pragma comment(linker, "/export:llvmGetPassPluginInfo")
 #endif
diff --git a/third_party/amd/backend/include/hsa/hsa_ext_amd.h b/third_party/amd/backend/include/hsa/hsa_ext_amd.h
@@ -61,6 +61,12 @@
 #define HSA_AMD_INTERFACE_VERSION_MAJOR 1
 #define HSA_AMD_INTERFACE_VERSION_MINOR 6
 
+#if defined(_WIN32)
+#define ALWAYS_INLINE __forceinline
+#else
+#define ALWAYS_INLINE __attribute__((always_inline)) inline
+#endif
+
 #ifdef __cplusplus
 extern "C" {
 #endif
@@ -73,8 +79,7 @@ extern "C" {
  * @brief Macro to use to determine that a  flag is set when querying flags within uint8_t[8]
  * types
  */
-static __inline__ __attribute__((always_inline)) bool hsa_flag_isset64(uint8_t* value,
-                                                                       uint32_t bit) {
+static ALWAYS_INLINE bool hsa_flag_isset64(uint8_t* value, uint32_t bit) {
   unsigned int index = bit / 8;
   unsigned int subBit = bit % 8;
   return ((uint8_t*)value)[index] & (1 << subBit);
diff --git a/third_party/amd/include/hipblas_instance.h b/third_party/amd/include/hipblas_instance.h
@@ -2,7 +2,7 @@
 #define TRITON_HIPBLAS_INSTANCE_H
 
 #include "hipblas_types.h"
-#ifdef WIN32
+#if defined(_WIN32)
 #define WIN32_LEAN_AND_MEAN
 #define NOMINMAX
 #include <windows.h>
@@ -84,7 +84,7 @@ class HipblasLtInstance {
   hipblasLtMatmulPreference_t preference = NULL;
 
   void loadHipBlasDylib() {
-#ifdef WIN32
+#if defined(_WIN32)
     assert(0 && "Not implemented on Windows");
 #else
     if (dylibHandle == nullptr) {
@@ -137,7 +137,7 @@ class HipblasLtInstance {
   }
 
   void unloadHipBlasDylib() {
-#ifdef WIN32
+#if defined(_WIN32)
     assert(0 && "Not implemented on Windows");
 #else
     dlclose(dylibHandle);
diff --git a/third_party/intel/backend/driver.c b/third_party/intel/backend/driver.c
@@ -169,7 +169,7 @@ struct BuildFlags {
 
 sycl::context get_default_context(const sycl::device &sycl_device) {
   const auto &platform = sycl_device.get_platform();
-#ifdef WIN32
+#if defined(_WIN32)
   sycl::context ctx;
   try {
 #if __SYCL_COMPILER_VERSION >= 20250604
diff --git a/third_party/intel/backend/proton/include/pti/pti_export.h b/third_party/intel/backend/proton/include/pti/pti_export.h
@@ -6,23 +6,46 @@
 #  define PTI_EXPORT
 #  define PTI_NO_EXPORT
 #else
-#  ifndef PTI_EXPORT
-#    ifdef pti_EXPORTS
-        /* We are building this library */
-#      define PTI_EXPORT __attribute__((visibility("default")))
-#    else
-        /* We are using this library */
-#      define PTI_EXPORT __attribute__((visibility("default")))
+
+#  ifdef WIN32
+     /* Windows (MSVC/MinGW) */
+#    ifndef PTI_EXPORT
+#      ifdef pti_EXPORTS
+         /* We are building this library */
+#        define PTI_EXPORT __declspec(dllexport)
+#      else
+         /* We are using this library */
+#        define PTI_EXPORT __declspec(dllimport)
+#      endif
 #    endif
-#  endif
 
-#  ifndef PTI_NO_EXPORT
-#    define PTI_NO_EXPORT __attribute__((visibility("hidden")))
-#  endif
-#endif
+#    ifndef PTI_NO_EXPORT
+#      define PTI_NO_EXPORT
+#    endif
+
+#  else
+
+     /* Linux / Unix — GCC/Clang visibility */
+#    ifndef PTI_EXPORT
+#      ifdef pti_EXPORTS
+         /* We are building this library */
+#        define PTI_EXPORT __attribute__((visibility("default")))
+#      else
+         /* We are using this library */
+#        define PTI_EXPORT __attribute__((visibility("default")))
+#      endif
+#    endif
+
+#    ifndef PTI_NO_EXPORT
+#      define PTI_NO_EXPORT __attribute__((visibility("hidden")))
+#    endif
+
+#  endif /* WIN32 */
+
+#endif /* PTI_STATIC_DEFINE */
 
 #ifndef PTI_DEPRECATED
-#  define PTI_DEPRECATED 
+#  define PTI_DEPRECATED
 #endif
 
 #ifndef PTI_DEPRECATED_EXPORT
diff --git a/third_party/nvidia/backend/driver.c b/third_party/nvidia/backend/driver.c
@@ -1,5 +1,5 @@
 #include "cuda.h"
-#ifdef WIN32
+#if defined(_WIN32)
 #define WIN32_LEAN_AND_MEAN
 #define NOMINMAX
 #include <windows.h>
@@ -180,7 +180,7 @@ typedef CUresult (*cuTensorMapEncodeTiled_t)(
     CUtensorMapSwizzle swizzle, CUtensorMapL2promotion l2Promotion,
     CUtensorMapFloatOOBfill oobFill);
 
-#ifdef WIN32
+#if defined(_WIN32)
 #define defineGetFunctionHandle(name, symbolName)                              \
   static symbolName##_t name() {                                               \
     /* Open the shared library */                                              \
@@ -196,6 +196,7 @@ typedef CUresult (*cuTensorMapEncodeTiled_t)(
     if (err) {                                                                 \
       PyErr_SetString(PyExc_RuntimeError,                                      \
                       "Failed to retrieve " #symbolName " from nvcuda.dll");   \
+      FreeLibrary(handle);                                                     \
       return NULL;                                                             \
     }                                                                          \
     return funcHandle;                                                         \
diff --git a/third_party/nvidia/include/cublas_instance.h b/third_party/nvidia/include/cublas_instance.h
@@ -2,7 +2,7 @@
 #define TRITON_CUBLAS_INSTANCE_H
 
 #include "cublas_types.h"
-#ifdef WIN32
+#if defined(_WIN32)
 #define WIN32_LEAN_AND_MEAN
 #define NOMINMAX
 #include <windows.h>
@@ -70,7 +70,7 @@ class CublasLtInstance {
   cublasLtMatmulPreference_t preference = NULL;
 
   void loadCublasDylib() {
-#ifdef WIN32
+#if defined(_WIN32)
     assert(0 && "Not implemented on Windows");
 #else
     if (dylibHandle == nullptr) {
@@ -121,7 +121,7 @@ class CublasLtInstance {
   }
 
   void unloadCublasDylib() {
-#ifdef WIN32
+#if defined(_WIN32)
     assert(0 && "Not implemented on Windows");
 #else
     dlclose(dylibHandle);
diff --git a/third_party/proton/CMakeLists.txt b/third_party/proton/CMakeLists.txt
@@ -83,6 +83,17 @@ add_library(proton SHARED ${_proton_obj_sources})
 
 target_link_libraries(proton PRIVATE Python3::Module)
 
+if(WIN32)
+  set_target_properties(proton PROPERTIES SUFFIX ".pyd")
+  set_target_properties(proton PROPERTIES PREFIX "lib")
+  if(DEFINED TRITON_PYD_PATH)
+    string(TOUPPER "${CMAKE_BUILD_TYPE}" UPPER_CMAKE_BUILD_TYPE)
+    set_target_properties(proton PROPERTIES
+      RUNTIME_OUTPUT_DIRECTORY_${UPPER_CMAKE_BUILD_TYPE}
+    "${TRITON_PYD_PATH}")
+  endif(DEFINED TRITON_PYD_PATH)
+endif()
+
 # Apply any macOS linker flags or extra link options
 if(PROTON_PYTHON_LDFLAGS)
   target_link_options(proton PRIVATE ${PROTON_PYTHON_LDFLAGS})
diff --git a/third_party/proton/csrc/include/Driver/Dispatch.h b/third_party/proton/csrc/include/Driver/Dispatch.h
@@ -1,7 +1,13 @@
 #ifndef PROTON_DRIVER_DISPATCH_H_
 #define PROTON_DRIVER_DISPATCH_H_
 
+#if defined(_WIN32)
+#define WIN32_LEAN_AND_MEAN
+#define NOMINMAX
+#include <windows.h>
+#else
 #include <dlfcn.h>
+#endif
 
 #include "Utility/Env.h"
 #include <stdexcept>
@@ -59,6 +65,31 @@ template <typename ExternLib> class Dispatch {
 public:
   Dispatch() = delete;
 
+#if defined(_WIN32)
+  static void init(const char *name, void **lib) {
+    if (*lib == nullptr) {
+      // If not found, try to load it from the default path
+      auto dir =
+          ExternLib::pathEnv == nullptr ? "" : getStrEnv(ExternLib::pathEnv);
+      if (!dir.empty()) {
+        auto fullPath = dir + "/" + name;
+        *lib = (void *)LoadLibrary(fullPath.c_str());
+      } else {
+        // Only if the default path is not set, we try to load it from the
+        // system.
+        // First reuse the existing handle
+        *lib = (void *)GetModuleHandle(name);
+        if (*lib == nullptr) {
+          // If not found, try to load it from LD_LIBRARY_PATH
+          *lib = (void *)LoadLibrary(name);
+        }
+      }
+    }
+    if (*lib == nullptr) {
+      throw std::runtime_error("Could not load `" + std::string(name) + "`");
+    }
+  }
+#else
   static void init(const char *name, void **lib) {
     if (*lib == nullptr) {
       // If not found, try to load it from the default path
@@ -82,6 +113,7 @@ template <typename ExternLib> class Dispatch {
       throw std::runtime_error("Could not load `" + std::string(name) + "`");
     }
   }
+#endif
 
   static void check(typename ExternLib::RetType ret, const char *functionName) {
     if (ret != ExternLib::success) {
@@ -96,7 +128,12 @@ template <typename ExternLib> class Dispatch {
   exec(FnT &handler, const char *functionName, Args... args) {
     init(ExternLib::name, &ExternLib::lib);
     if (handler == nullptr) {
+#if defined(_WIN32)
+      handler = reinterpret_cast<FnT>(
+          GetProcAddress((HMODULE)ExternLib::lib, functionName));
+#else
       handler = reinterpret_cast<FnT>(dlsym(ExternLib::lib, functionName));
+#endif
       if (handler == nullptr) {
         throw std::runtime_error("Failed to load " +
                                  std::string(ExternLib::name));
@@ -118,12 +155,30 @@ template <typename ExternLib> class Dispatch {
       }
     }
     if (ExternLib::lib != nullptr) {
+#if defined(_WIN32)
+      void *sym =
+          GetProcAddress((HMODULE)ExternLib::lib,
+                         ExternLib::symbolName); // pick any known symbol
+      HMODULE mod = nullptr;
+
+      if (GetModuleHandleEx(GET_MODULE_HANDLE_EX_FLAG_FROM_ADDRESS |
+                                GET_MODULE_HANDLE_EX_FLAG_UNCHANGED_REFCOUNT,
+                            (LPCSTR)sym, &mod)) {
+        static char path[MAX_PATH];
+        DWORD len = GetModuleFileName(mod, path, MAX_PATH);
+        if (len > 0)
+          return path;
+      }
+
+      return nullptr;
+#else
       void *sym = dlsym(ExternLib::lib,
                         ExternLib::symbolName); // pick any known symbol
       Dl_info info;
       if (dladdr(sym, &info)) {
         return info.dli_fname;
       }
+#endif
     }
     return "";
   }
diff --git a/third_party/proton/csrc/lib/Driver/GPU/HipApi.cpp b/third_party/proton/csrc/lib/Driver/GPU/HipApi.cpp
@@ -88,8 +88,13 @@ const char *getKernelNameRef(const hipFunction_t f) {
   static hipKernelNameRef_t func = nullptr;
   Dispatch<ExternLibHip>::init(ExternLibHip::name, &ExternLibHip::lib);
   if (func == nullptr)
+#if defined(_WIN32)
+    func = reinterpret_cast<hipKernelNameRef_t>(
+        GetProcAddress((HMODULE)ExternLibHip::lib, "hipKernelNameRef"));
+#else
     func = reinterpret_cast<hipKernelNameRef_t>(
         dlsym(ExternLibHip::lib, "hipKernelNameRef"));
+#endif
   return (func ? func(f) : NULL);
 }
 
@@ -99,8 +104,13 @@ const char *getKernelNameRefByPtr(const void *hostFunction,
   static hipKernelNameRefByPtr_t func = nullptr;
   Dispatch<ExternLibHip>::init(ExternLibHip::name, &ExternLibHip::lib);
   if (func == nullptr)
+#if defined(_WIN32)
+    func = reinterpret_cast<hipKernelNameRefByPtr_t>(
+        GetProcAddress((HMODULE)ExternLibHip::lib, "hipKernelNameRefByPtr"));
+#else
     func = reinterpret_cast<hipKernelNameRefByPtr_t>(
         dlsym(ExternLibHip::lib, "hipKernelNameRefByPtr"));
+#endif
   return (func ? func(hostFunction, stream) : NULL);
 }
 
diff --git a/third_party/proton/csrc/lib/Driver/GPU/HsaApi.cpp b/third_party/proton/csrc/lib/Driver/GPU/HsaApi.cpp
@@ -26,8 +26,13 @@ hsa_status_t iterateAgents(hsa_status_t (*callback)(hsa_agent_t agent,
   static hsa_iterate_agents_t func = nullptr;
   Dispatch<ExternLibHsa>::init(ExternLibHsa::name, &ExternLibHsa::lib);
   if (func == nullptr)
+#if defined(_WIN32)
+    func = reinterpret_cast<hsa_iterate_agents_t>(
+        GetProcAddress((HMODULE)ExternLibHsa::lib, "hsa_iterate_agents"));
+#else
     func = reinterpret_cast<hsa_iterate_agents_t>(
         dlsym(ExternLibHsa::lib, "hsa_iterate_agents"));
+#endif
   return (func ? func(callback, data) : HSA_STATUS_ERROR_FATAL);
 }
 
diff --git a/third_party/proton/csrc/lib/Driver/GPU/NvtxApi.cpp b/third_party/proton/csrc/lib/Driver/GPU/NvtxApi.cpp
@@ -22,11 +22,21 @@ void enable() {
   const std::string cuptiLibPath =
       Dispatch<cupti::ExternLibCupti>::getLibPath();
   if (!cuptiLibPath.empty()) {
+#if defined(_WIN32)
+    _putenv_s("NVTX_INJECTION64_PATH", cuptiLibPath.c_str());
+#else
     setenv("NVTX_INJECTION64_PATH", cuptiLibPath.c_str(), 1);
+#endif
   }
 }
 
-void disable() { unsetenv("NVTX_INJECTION64_PATH"); }
+void disable() {
+#if defined(_WIN32)
+  _putenv_s("NVTX_INJECTION64_PATH", "");
+#else
+  unsetenv("NVTX_INJECTION64_PATH");
+#endif
+}
 
 std::string getMessageFromRangePushA(const void *params) {
   if (const auto *p = static_cast<const RangePushAParams *>(params))
diff --git a/third_party/proton/csrc/lib/Driver/GPU/RoctracerApi.cpp b/third_party/proton/csrc/lib/Driver/GPU/RoctracerApi.cpp
diff --git a/third_party/proton/csrc/lib/Driver/GPU/XpuApi.cpp b/third_party/proton/csrc/lib/Driver/GPU/XpuApi.cpp
diff --git a/third_party/proton/csrc/lib/Profiler/Cupti/CuptiProfiler.cpp b/third_party/proton/csrc/lib/Profiler/Cupti/CuptiProfiler.cpp
diff --git a/third_party/proton/csrc/lib/Profiler/RocTracer/RoctracerProfiler.cpp b/third_party/proton/csrc/lib/Profiler/RocTracer/RoctracerProfiler.cpp
diff --git a/third_party/proton/csrc/lib/Profiler/Xpupti/XpuptiProfiler.cpp b/third_party/proton/csrc/lib/Profiler/Xpupti/XpuptiProfiler.cpp
diff --git a/utils/SPIRVRunner/SPIRVRunner.cpp b/utils/SPIRVRunner/SPIRVRunner.cpp

Original file line number	Diff line number	Diff line change
`@@ -101,6 +101,6 @@ extern "C" LLVM_ATTRIBUTE_WEAK PassPluginLibraryInfo llvmGetPassPluginInfo() {`
`101`	`101`	`return getPassPluginInfo();`
`102`	`102`	`}`
`103`	`103`
`104`		`-#ifdef WIN32`
	`104`	`+#if defined(_WIN32)`
`105`	`105`	`#pragma comment(linker, "/export:llvmGetPassPluginInfo")`
`106`	`106`	`#endif`
Original file line number	Diff line number	Diff line change
`@@ -75,6 +75,6 @@ llvmGetPassPluginInfo() {`
`75`	`75`	`return getPassPluginInfo();`
`76`	`76`	`}`
`77`	`77`
`78`		`-#ifdef WIN32`
	`78`	`+#if defined(_WIN32)`
`79`	`79`	`#pragma comment(linker, "/export:llvmGetPassPluginInfo")`
`80`	`80`	`#endif`