build_libtriton_from_llvm.diff

diff --git a/CMakeLists.txt b/CMakeLists.txt
index de6ed2393..db776c0fc 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -47,6 +47,8 @@ if (TRITON_PARALLEL_LINK_JOBS)
     set(CMAKE_JOB_POOL_LINK link_job_pool)
 endif()
 
+string(REPLACE "-Wl,-z,defs", "" CMAKE_MODULE_LINKER_FLAGS ${CMAKE_MODULE_LINKER_FLAGS})
+string(REPLACE "-Wl,-z,defs", "" CMAKE_SHARED_LINKER_FLAGS ${CMAKE_SHARED_LINKER_FLAGS})
 
 # Ensure Python3 vars are set correctly
 # used conditionally in this file and by lit tests
@@ -143,7 +145,7 @@ endfunction()
 
 # Disable warnings that show up in external code (gtest;pybind11)
 if(NOT MSVC)
-  set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Werror -Wno-covered-switch-default -fvisibility=hidden")
+  set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wno-covered-switch-default -fvisibility=hidden")
 else()
   set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /wd4244 /wd4624 /wd4715 /wd4530")
 endif()
@@ -236,6 +238,7 @@ if(TRITON_BUILD_PYTHON_MODULE)
     MLIRIndexToLLVM
     MLIRGPUToROCDLTransforms
     MLIRUBToLLVM
+    MLIRTosaDialect
 
     # LLVM
     LLVMPasses
@@ -285,6 +288,7 @@ if(TRITON_BUILD_PYTHON_MODULE)
                   ${PYTHON_SRC_PATH}/passes.cc
                   ${PYTHON_SRC_PATH}/interpreter.cc
                   ${PYTHON_SRC_PATH}/llvm.cc)
+  set_target_properties(triton PROPERTIES LIBRARY_OUTPUT_DIRECTORY ${CMAKE_CURRENT_LIST_DIR}/python/triton/_C)
 
   # Link triton with its dependencies
   target_link_libraries(triton PRIVATE ${TRITON_LIBRARIES})
@@ -323,7 +327,7 @@ endif()
 find_package(Threads REQUIRED)
 
 add_subdirectory(third_party/f2reduce)
-add_subdirectory(bin)
+#add_subdirectory(bin)
 add_subdirectory(test)
 
 if(TRITON_BUILD_UT)
@@ -335,3 +339,30 @@ if(TRITON_BUILD_UT)
     USES_TERMINAL
   )
 endif()
+
+set(_uninstalled
+        NVGPUIR
+        NVGPUToLLVM
+        ProtonIR
+        TritonAMDGPUDialectToLLVM
+        TritonAMDGPUIR
+        TritonAMDGPUToLLVM
+        TritonAMDGPUTransforms
+        TritonAMDUtils
+        TritonAnalysis
+        TritonGPUIR
+        TritonGPUToLLVM
+        TritonGPUTransforms
+        TritonIR
+        TritonLLVMIR
+        TritonNVIDIAGPUToLLVM
+        TritonNvidiaGPUIR
+        TritonNvidiaGPUTransforms
+        TritonProtonToLLVM
+        TritonToTritonGPU
+        TritonTools
+        TritonTransforms
+        f2reduce
+        TritonAMDAnalysis
+)
+install(TARGETS ${_uninstalled} EXPORT MLIRTargets)
diff --git a/python/src/llvm.cc b/python/src/llvm.cc
index c86bf671a..3ab0c1b14 100644
--- a/python/src/llvm.cc
+++ b/python/src/llvm.cc
@@ -315,8 +315,8 @@ void init_triton_llvm(py::module &&m) {
         PassInstrumentationCallbacks passInstrCb;
         StandardInstrumentations standardInstr(mod->getContext(),
                                                /*DebugLogging*/ true);
+        auto optMap = llvm::cl::getRegisteredOptions();
         if (mlir::triton::tools::getBoolEnv("LLVM_IR_ENABLE_DUMP")) {
-          auto optMap = llvm::cl::getRegisteredOptions();
           auto optIt = optMap.find("print-after-all");
           if (optIt != optMap.end()) {
             auto optPtr = static_cast<llvm::cl::opt<bool> *>(optIt->second);
@@ -325,6 +325,9 @@ void init_triton_llvm(py::module &&m) {
           standardInstr.registerCallbacks(passInstrCb, &mam);
           instrCbPtr = &passInstrCb;
         }
+        auto optIt = optMap.find("global-isel");
+        auto optPtr = static_cast<llvm::cl::opt<bool> *>(optIt->second);
+        *optPtr = true;
 
         PipelineTuningOptions tuningOptions;
         tuningOptions.LoopUnrolling = true;
diff --git a/python/src/passes.cc b/python/src/passes.cc
index 04b76d17d..57bb11b6e 100644
--- a/python/src/passes.cc
+++ b/python/src/passes.cc
@@ -77,7 +77,7 @@ void init_triton_passes_ttgpuir(py::module &&m) {
 
 void init_triton_passes_convert(py::module &&m) {
   using namespace mlir;
-  ADD_PASS_WRAPPER_0("add_scf_to_cf", createConvertSCFToCFPass);
+  ADD_PASS_WRAPPER_0("add_scf_to_cf", createSCFToControlFlowPass);
   ADD_PASS_WRAPPER_0("add_cf_to_llvmir", createConvertControlFlowToLLVMPass);
   ADD_PASS_WRAPPER_0("add_index_to_llvmir", createConvertIndexToLLVMPass);
   ADD_PASS_WRAPPER_0("add_arith_to_llvmir", createArithToLLVMConversionPass);
diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt
index edef53315..e009599fb 100644
--- a/test/CMakeLists.txt
+++ b/test/CMakeLists.txt
@@ -12,9 +12,9 @@ configure_lit_site_cfg(
 )
 
 set(TRITON_TEST_DEPENDS
-  triton-opt
-  triton-tensor-layout
-  triton-llvm-opt
+#  triton-opt
+#  triton-tensor-layout
+#  triton-llvm-opt
 )
 
 set(FILECHECK_PATH "${LLVM_LIBRARY_DIR}/../bin/FileCheck")
diff --git a/third_party/amd/lib/TritonAMDGPUToLLVM/TargetUtils.cpp b/third_party/amd/lib/TritonAMDGPUToLLVM/TargetUtils.cpp
index 5bce008f9..211e5f95c 100644
--- a/third_party/amd/lib/TritonAMDGPUToLLVM/TargetUtils.cpp
+++ b/third_party/amd/lib/TritonAMDGPUToLLVM/TargetUtils.cpp
@@ -14,8 +14,8 @@ ISAFamily deduceISAFamily(llvm::StringRef arch) {
   case llvm::AMDGPU::GK_GFX950:
     return ISAFamily::CDNA4;
   case llvm::AMDGPU::GK_GFX942:
-  case llvm::AMDGPU::GK_GFX941:
-  case llvm::AMDGPU::GK_GFX940:
+  // case llvm::AMDGPU::GK_GFX941:
+  // case llvm::AMDGPU::GK_GFX940:
     return ISAFamily::CDNA3;
   case llvm::AMDGPU::GK_GFX90A:
     return ISAFamily::CDNA2;
diff --git a/third_party/amd/lib/TritonAMDGPUTransforms/CMakeLists.txt b/third_party/amd/lib/TritonAMDGPUTransforms/CMakeLists.txt
index 4ca78d3ea..0031934b4 100644
--- a/third_party/amd/lib/TritonAMDGPUTransforms/CMakeLists.txt
+++ b/third_party/amd/lib/TritonAMDGPUTransforms/CMakeLists.txt
@@ -17,5 +17,5 @@ add_triton_library(TritonAMDGPUTransforms
   TritonAMDAnalysis
 )
 
-target_include_directories(TritonAMDGPUTransforms PUBLIC ${CMAKE_CURRENT_SOURCE_DIR}/../../include)
-target_include_directories(TritonAMDGPUTransforms PUBLIC ${CMAKE_CURRENT_BINARY_DIR}/../../include)
+target_include_directories(TritonAMDGPUTransforms PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../../include)
+target_include_directories(TritonAMDGPUTransforms PRIVATE ${CMAKE_CURRENT_BINARY_DIR}/../../include)