diff --git a/CMakeLists.txt b/CMakeLists.txt index 23fa7270b9..bb5985b3de 100644 --- 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 "$<$:-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 deleted file mode 100644 index 0e89371e07..0000000000 --- a/examples/CMakeLists.txt +++ /dev/null @@ -1 +0,0 @@ -add_subdirectory(plugins) diff --git a/examples/plugins/CMakeLists.txt b/examples/plugins/CMakeLists.txt index cd87443650..fe352dc0bd 100644 --- 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 index 2e02718000..ff7c66a350 100644 --- 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 + "$<$:-undefined dynamic_lookup>" ) diff --git a/examples/plugins/README.md b/examples/plugins/README.md index c5615e9ae9..a5523935cd 100644 --- 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 index a8c174d7f4..4980d02ddd 100644 --- 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 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 index 55ba6c3695..5c42895345 100644 --- 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 index c56c17b3af..4caaeef5bf 100644 --- 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 index f06ad0c961..e45356d5c3 100755 --- 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 index ab821f3535..7cbadda25f 100644 --- 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 index 0bc060cea1..fc789a9102 100644 --- 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 index da21f002fa..41a4d118d7 100644 --- 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 index 915533fd2c..3d2b0ee50f 100644 --- 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 index d540d8f868..1dfc0c5cf5 100644 --- 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 index a891cba83f..90e4ae71ef 100644 --- 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