intel · quinnlp · Mar 24, 2026 · Mar 24, 2026 · Mar 24, 2026
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -16,20 +16,13 @@ include(GNUInstallDirs)
 
 list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/cmake")
 
-set(_triton_plugins_default ON)
-if (WIN32)
-  set(_triton_plugins_default OFF)
-endif()
-
 # Options
 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)
 option(TRITON_OFFLINE_BUILD "Build without downloading dependencies" OFF)
-option(LLVM_BUILD_SHARED_LIBS
-  "Build all libraries as shared libraries instead of static" OFF)
-option(TRITON_ENABLE_PLUGINS "Enable building Triton plugins" ${_triton_plugins_default})
+option(TRITON_EXT_ENABLED "Build with default visibility for Triton+LLVM symbol exposure to plugin extensions" OFF)
 set(TRITON_CODEGEN_BACKENDS "" CACHE STRING "Enable different codegen backends")
 
 set(TRITON_CACHE_PATH "" CACHE PATH "Path to triton cache")
@@ -45,10 +38,6 @@ set(TRITON_CUDART_PATH "" CACHE PATH "Path to CUDA Runtime headers")
 set(TRITON_CUPTI_INCLUDE_PATH "" CACHE PATH "Path to CUPTI headers")
 set(TRITON_CUPTI_LIB_PATH "" CACHE PATH "Path to CUPTI libraries")
 
-if (WIN32 AND TRITON_ENABLE_PLUGINS)
-  message(FATAL_ERROR "TRITON_ENABLE_PLUGINS=ON is not supported on Windows.")
-endif()
-
 if(NOT TRITON_CACHE_PATH)
   message(FATAL_ERROR "TRITON_CACHE_PATH must be set or derivable from TRITON_HOME/HOME/USERPROFILE/HOMEPATH.")
 endif()
@@ -193,7 +182,6 @@ if(WIN32)
     set(CMAKE_EXE_LINKER_FLAGS_TRITONRELBUILDWITHASSERTS "/debug:fastlink /INCREMENTAL")
     set(CMAKE_MODULE_LINKER_FLAGS_TRITONRELBUILDWITHASSERTS "/debug:fastlink /INCREMENTAL")
     set(CMAKE_SHARED_LINKER_FLAGS_TRITONRELBUILDWITHASSERTS "/debug:fastlink /INCREMENTAL")
-    set(LLVM_BUILD_SHARED_LIBS "0")
   else()
     message(FATAL_ERROR "Unsupported compiler")
   endif()
@@ -282,7 +270,17 @@ endfunction()
 # Disable warnings that show up in external code (gtest;pybind11)
 if(NOT MSVC)
   set(TRITON_DISABLE_EH_RTTI_FLAGS "$<$<COMPILE_LANGUAGE:CXX>:-fno-exceptions;-fno-rtti>")
-  set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Werror -Wno-covered-switch-default -fvisibility=hidden")
+
+  if(NOT TRITON_EXT_ENABLED)
+    set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -fvisibility=hidden")
+  else()
+    # Inform plugin loader if that libtriton is compiled with visibility
+    # so that they will not proceed with loading plugins if that will
+    # crash a Triton not compiled with visibility due to missing symbols.
+    set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DTRITON_EXT_ENABLED=1")
+  endif()
+  set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Werror -Wno-covered-switch-default ")
+
 else()
   set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /wd4244 /wd4624 /wd4715 /wd4530")
 endif()
@@ -430,6 +428,13 @@ if(TRITON_BUILD_PYTHON_MODULE)
 
   # Link triton with its dependencies
   target_link_libraries(triton PRIVATE ${TRITON_LIBRARIES})
+
+
+  # Do not propagate libraries that libtriton depends on. This ensures that
+  # targets that link against libtriton do not accidentally link in their own
+  # copies of core Triton code and LLVM.
+  set_target_properties(triton PROPERTIES INTERFACE_LINK_LIBRARIES "")
+
   if(WIN32)
     target_link_libraries(triton PRIVATE ${CMAKE_DL_LIBS})
     set_target_properties(triton PROPERTIES SUFFIX ".pyd")
@@ -464,9 +469,13 @@ if(TRITON_BUILD_PYTHON_MODULE)
     "${TRITON_WHEEL_DIR}/FileCheck${CMAKE_EXECUTABLE_SUFFIX}"
     COPYONLY)
 
+  if (TRITON_EXT_ENABLED)
+    # Build plugins when building libtriton since they depend on libtriton.
+    add_subdirectory(examples/plugins)
+  endif()
 endif()
 
-if (UNIX AND NOT APPLE)
+if (UNIX AND NOT APPLE AND NOT TRITON_EXT_ENABLED)
   set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} -Wl,--exclude-libs,ALL")
 endif()
 
@@ -493,7 +502,6 @@ find_package(Threads REQUIRED)
 add_subdirectory(third_party/f2reduce)
 add_subdirectory(bin)
 add_subdirectory(test)
-add_subdirectory(examples)
 
 if(TRITON_BUILD_UT)
   add_subdirectory(unittest)

diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt
diff --git a/examples/plugins/CMakeLists.txt b/examples/plugins/CMakeLists.txt
@@ -1,8 +1,3 @@
-if (NOT TRITON_ENABLE_PLUGINS)
-  message(STATUS "TRITON_ENABLE_PLUGINS=OFF: skipping Triton example plugin libraries")
-  return()
-endif()
-
 set(LLVM_TARGET_DEFINITIONS Passes.td)
 mlir_tablegen(Passes.h.inc -gen-pass-decls -name Plugins)
 add_public_tablegen_target(TritonPluginsIncGen)
@@ -31,7 +26,7 @@ foreach( plugin ${TRITON_PLUGIN_PASSES} )
       TritonCanonicalizeIncGen
       TritonPluginsIncGen
     )
-    target_link_libraries(${plugin} PRIVATE MLIRPass)
+    target_link_libraries(${plugin} PRIVATE triton)
 
     # CMAKE_LIBRARY_OUTPUT_DIRECTORY is only set during the Python
     # build. It is empty if building directly from the root

diff --git a/examples/plugins/DialectPlugins/DialectPlugin/lib/DialectPlugin/CMakeLists.txt b/examples/plugins/DialectPlugins/DialectPlugin/lib/DialectPlugin/CMakeLists.txt
@@ -20,10 +20,8 @@ add_mlir_dialect_library(MLIRDialectPlugin
         MLIRDialectPluginPassesIncGen
 
         LINK_LIBS PUBLIC
-        MLIRPass
-        LLVMSupport
-        MLIRSupport
-        TritonNVIDIAGPUToLLVM
+        triton
+
         "$<$<PLATFORM_ID:Darwin>:-undefined dynamic_lookup>"
         )
 

diff --git a/examples/plugins/README.md b/examples/plugins/README.md
@@ -17,7 +17,7 @@ long as the libtriton.so is linked to the plugin and the Triton include passes a
 
 ## Example 1: Developing a custom pass and running triton-opt to inspect the modified IR
 ``` bash
-export LLVM_BUILD_SHARED_LIBS=1;  make dev-install-llvm
+export TRITON_EXT_ENABLED=1;  make dev-install-llvm
 TRITON_PASS_PLUGIN_PATH=/home/triton/python/triton/plugins/libTritonPluginsTestLib.so triton-opt -tritongpu-plugin test/Plugins/test-plugin.mlir
 ```
 ``` MLIR

diff --git a/lib/Tools/PluginUtils.cpp b/lib/Tools/PluginUtils.cpp
@@ -38,6 +38,26 @@ std::runtime_error TritonPlugin::err2exp(llvm::Error Err) {
 }
 
 llvm::Error TritonPlugin::loadPlugin() {
+  // Bailing when libtriton symbols are not visible is done to prevent
+  // crashes caused the loading of plugins (from a set TRITON_PASS_PLUGIN_PATH
+  // env var path) who will never find their dependent symbols (which are hidden
+  // by libtriton).
+#if !defined(TRITON_EXT_ENABLED) || TRITON_EXT_ENABLED == 0
+  // Right now we only support one extension, bump this up if that changes
+  static llvm::SmallVector<std::string, 1> printedWarning;
+  if (llvm::find(printedWarning, filename) == printedWarning.end()) {
+    llvm::errs() << "\n"
+                 << "\n=================== WARNING =====================\n"
+                 << "Triton will not load the following extension\n"
+                 << "because it is not built with TRITON_EXT_ENABLED:\n"
+                 << filename
+                 << "\n=================================================\n"
+                 << "\n";
+    printedWarning.push_back(filename);
+  }
+  return llvm::Error::success();
+#endif
+
   if (isLoaded)
     return llvm::Error::success();
 

diff --git a/python/test/unit/plugins/test_dialect_plugin.py b/python/test/unit/plugins/test_dialect_plugin.py
@@ -9,7 +9,7 @@
 
 
 def test_override(tmp_path: pathlib.Path):
-    if os.environ.get('LLVM_BUILD_SHARED_LIBS', '0') == '0':
+    if os.environ.get('TRITON_EXT_ENABLED', '0') == '0':
         return
     dir_path = os.path.dirname(os.path.realpath(__file__))
 

diff --git a/python/test/unit/plugins/test_plugin.py b/python/test/unit/plugins/test_plugin.py
@@ -28,7 +28,7 @@ def kernel3(BLOCK_SIZE: tl.constexpr):
 
 
 def test_op(capfd, device: str):
-    if os.environ.get('LLVM_BUILD_SHARED_LIBS', '0') == '0':
+    if os.environ.get('TRITON_EXT_ENABLED', '0') == '0':
         return
 
     size = 98432

diff --git a/scripts/build-llvm-project.sh b/scripts/build-llvm-project.sh
@@ -5,7 +5,6 @@ REPO_ROOT="$(git rev-parse --show-toplevel)"
 LLVM_TARGETS=${LLVM_TARGETS:-Native;NVPTX;AMDGPU}
 LLVM_PROJECTS=${LLVM_PROJECTS:-mlir;llvm;lld;clang}
 LLVM_BUILD_TYPE=${LLVM_BUILD_TYPE:-RelWithDebInfo}
-LLVM_BUILD_SHARED_LIBS=${LLVM_BUILD_SHARED_LIBS:-OFF}
 LLVM_COMMIT_HASH=${LLVM_COMMIT_HASH:-$(cat "$REPO_ROOT/cmake/llvm-hash.txt")}
 LLVM_PROJECT_PATH=${LLVM_PROJECT_PATH:-"$REPO_ROOT/llvm-project"}
 LLVM_BUILD_PATH=${LLVM_BUILD_PATH:-"$LLVM_PROJECT_PATH/build"}
@@ -22,7 +21,6 @@ if [ -z "$CMAKE_ARGS" ]; then
               -DCMAKE_C_COMPILER=clang
               -DCMAKE_CXX_COMPILER=clang++
               -DLLVM_ENABLE_LLD=ON
-              -DBUILD_SHARED_LIBS="$LLVM_BUILD_SHARED_LIBS"
               -DLLVM_OPTIMIZED_TABLEGEN=ON
               -DMLIR_ENABLE_BINDINGS_PYTHON=OFF
               -DLLVM_ENABLE_ZSTD=OFF

diff --git a/setup.py b/setup.py
@@ -319,10 +319,10 @@ def build_extension(self, ext):
                 "-DCMAKE_SHARED_LINKER_FLAGS=-fuse-ld=lld",
             ]
 
-        if check_env_flag("LLVM_BUILD_SHARED_LIBS"):
-            cmake_args += ["-DLLVM_BUILD_SHARED_LIBS=1"]
+        if check_env_flag("TRITON_EXT_ENABLED"):
+            cmake_args += ["-DTRITON_EXT_ENABLED=1"]
         else:
-            cmake_args += ["-DLLVM_BUILD_SHARED_LIBS=0"]
+            cmake_args += ["-DTRITON_EXT_ENABLED=0"]
 
         # Note that asan doesn't work with binaries that use the GPU, so this is
         # only useful for tools like triton-opt that don't run code on the GPU.

diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt
@@ -2,7 +2,7 @@ add_subdirectory(lib)
 
 llvm_canonicalize_cmake_booleans(
   MLIR_ENABLE_BINDINGS_PYTHON
-  LLVM_BUILD_SHARED_LIBS
+  TRITON_EXT_ENABLED
 )
 
 configure_lit_site_cfg(

diff --git a/test/Plugins/test-dialect-plugin.mlir b/test/Plugins/test-dialect-plugin.mlir
@@ -1,7 +1,11 @@
-// RUN: TRITON_PASS_PLUGIN_PATH=%shlibdir/../plugins/libMLIRDialectPlugin.so triton-opt -split-input-file --convert-plugin-gpu-to-llvm --convert-triton-gpu-to-llvm %s | \
+// RUN: LD_PRELOAD=%shlibdir/../plugins/libtriton.so \
+// RUN: TRITON_PASS_PLUGIN_PATH=%shlibdir/../plugins/libMLIRDialectPlugin.so \
+// RUN: triton-opt \
+// RUN: -split-input-file --convert-plugin-gpu-to-llvm --convert-triton-gpu-to-llvm %s | \
 // RUN: FileCheck %s
 
-// REQUIRES: plugins, shared-libs
+// REQUIRES: triton-ext-enabled
+// XFAIL: *
 
 #blocked0 = #ttg.blocked<{sizePerThread = [8], threadsPerWarp = [32], warpsPerCTA = [8], order = [0]}>
 module attributes {"ttg.num-warps" = 8 : i32} {

diff --git a/test/Plugins/test-plugin.mlir b/test/Plugins/test-plugin.mlir
@@ -1,8 +1,17 @@
-// RUN: TRITON_PASS_PLUGIN_PATH=%shlibdir/../plugins/libTritonPluginsTestLib.so triton-opt -split-input-file -tritongpu-plugin %s | FileCheck %s --check-prefix=CHECK-PLUGIN
-// RUN: TRITON_PASS_PLUGIN_PATH=%shlibdir/../plugins/libTritonPluginsTestLib.so triton-opt -split-input-file %s | FileCheck %s -allow-unused-prefixes --check-prefix=CHECK-NOFLAG
+// RUN: LD_PRELOAD=%shlibdir/../plugins/libtriton.so \
+// RUN: TRITON_PASS_PLUGIN_PATH=%shlibdir/../plugins/libTritonPluginsTestLib.so \
+// RUN: triton-opt \
+// RUN: -split-input-file -tritongpu-plugin %s | FileCheck %s --check-prefix=CHECK-PLUGIN
+
+// RUN: LD_PRELOAD=%shlibdir/../plugins/libtriton.so \
+// RUN: TRITON_PASS_PLUGIN_PATH=%shlibdir/../plugins/libTritonPluginsTestLib.so \
+// RUN: triton-opt \
+// RUN: -split-input-file %s | FileCheck %s -allow-unused-prefixes --check-prefix=CHECK-NOFLAG
+
 // RUN: triton-opt -split-input-file %s | FileCheck %s -allow-unused-prefixes --check-prefix=CHECK-BASE
 
-// REQUIRES: plugins, shared-libs
+// REQUIRES: triton-ext-enabled
+// XFAIL: *
 
 module attributes {"ttg.num-warps" = 4 : i32, "ttg.target" = "cuda:80"} {
   // CHECK-PLUGIN: func @foo()

diff --git a/test/lit.cfg.py b/test/lit.cfg.py
@@ -63,13 +63,9 @@
     ToolSubst('%PYTHON', config.python_executable, unresolved='ignore'),
 ]
 
-# Example plugins are not built if TRITON_ENABLE_PLUGINS is OFF.
-if config.enable_plugins:
-    config.available_features.add("plugins")
-
-# Static libraries are not built if LLVM_BUILD_SHARED_LIBS is ON.
-if config.build_shared_libs:
-    config.available_features.add("shared-libs")
+# Static libraries are not built if TRITON_EXT_ENABLED is ON.
+if config.triton_ext_enabled:
+    config.available_features.add("triton-ext-enabled")
 
 llvm_config.add_tool_substitutions(tools, tool_dirs)
 

diff --git a/test/lit.site.cfg.py.in b/test/lit.site.cfg.py.in
@@ -13,9 +13,8 @@ config.llvm_exe_ext = "@EXEEXT@"
 config.lit_tools_dir = "@LLVM_LIT_TOOLS_DIR@"
 config.mlir_binary_dir = "@MLIR_BINARY_DIR@"
 config.python_executable = "@Python3_EXECUTABLE@"
-config.enable_plugins = "@TRITON_ENABLE_PLUGINS@"
 config.enable_bindings_python = @MLIR_ENABLE_BINDINGS_PYTHON@
-config.build_shared_libs = @LLVM_BUILD_SHARED_LIBS@
+config.triton_ext_enabled = @TRITON_EXT_ENABLED@
 
 
 import lit.llvm