diff --git a/CMakeLists.txt b/CMakeLists.txt index aae75779c421..6bef97bf7bf7 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -326,10 +326,8 @@ if(TRITON_BUILD_PYTHON_MODULE) get_property(triton_libs GLOBAL PROPERTY TRITON_LIBS) get_property(triton_plugins GLOBAL PROPERTY TRITON_PLUGINS) - set(TRITON_LIBRARIES + set(TRITON_COMPILER_LIBRARIES ${triton_libs} - ${triton_plugins} - # mlir MLIRAMDGPUDialect MLIRNVVMDialect @@ -360,7 +358,11 @@ if(TRITON_BUILD_PYTHON_MODULE) # LLVMNVPTXAsmPrinter LLVMAMDGPUCodeGen LLVMAMDGPUAsmParser + ) + set(TRITON_LIBRARIES + ${triton_plugins} + ${TRITON_COMPILER_LIBRARIES} Python3::Module pybind11::headers diff --git a/bin/CMakeLists.txt b/bin/CMakeLists.txt index 1d3f647c2c29..a2dc053a9851 100644 --- a/bin/CMakeLists.txt +++ b/bin/CMakeLists.txt @@ -1,10 +1,12 @@ get_property(triton_libs GLOBAL PROPERTY TRITON_LIBS) +add_subdirectory(FakeLibTriton) + add_executable(triton-opt triton-opt.cpp) target_compile_options(triton-opt PRIVATE ${TRITON_DISABLE_EH_RTTI_FLAGS}) target_link_libraries(triton-opt PRIVATE - ${triton_libs} + ${TRITON_COMPILER_LIBRARIES} # tests TritonTestAnalysis TritonTestDialect @@ -16,9 +18,11 @@ target_link_libraries(triton-opt PRIVATE MLIRRegisterAllDialects MLIRRegisterAllPasses MLIRTransforms + TritonNVIDIAGPUToLLVM ) mlir_check_all_link_libraries(triton-opt) +export_executable_symbols_for_plugins(triton-opt) add_executable(triton-reduce triton-reduce.cpp) mlir_check_all_link_libraries(triton-reduce) diff --git a/bin/FakeLibTriton/CMakeLists.txt b/bin/FakeLibTriton/CMakeLists.txt new file mode 100644 index 000000000000..268e6146d87f --- /dev/null +++ b/bin/FakeLibTriton/CMakeLists.txt @@ -0,0 +1,11 @@ +add_library(faketriton SHARED empty.c) + +set_target_properties(faketriton PROPERTIES + OUTPUT_NAME "triton" +) + +if(DEFINED CMAKE_LIBRARY_OUTPUT_DIRECTORY) + set_target_properties(faketriton PROPERTIES + LIBRARY_OUTPUT_DIRECTORY + "${CMAKE_BINARY_DIR}/bin/faketriton") +endif(DEFINED CMAKE_LIBRARY_OUTPUT_DIRECTORY) diff --git a/bin/FakeLibTriton/empty.c b/bin/FakeLibTriton/empty.c new file mode 100644 index 000000000000..e69de29bb2d1 diff --git a/bin/triton-opt.cpp b/bin/triton-opt.cpp index 2d2570771ab4..4e1fd7f305c1 100644 --- a/bin/triton-opt.cpp +++ b/bin/triton-opt.cpp @@ -1,8 +1,23 @@ #include "./RegisterTritonDialects.h" #include "mlir/Tools/mlir-opt/MlirOptMain.h" +#include "llvm/Support/DynamicLibrary.h" +#include "llvm/Support/FileSystem.h" +#include "llvm/Support/Path.h" int main(int argc, char **argv) { + if (std::string filename = + mlir::triton::tools::getStrEnv("TRITON_PASS_PLUGIN_PATH"); + !filename.empty()) { + auto executablePath = + llvm::sys::fs::getMainExecutable(argv[0], (void *)&main); + std::string fakeTriton = + llvm::sys::path::parent_path(executablePath).str() + + "/faketriton/libtriton.so"; + std::string error; + llvm::sys::DynamicLibrary::getPermanentLibrary(fakeTriton.c_str(), &error); + } + mlir::DialectRegistry registry; registerTritonDialects(registry); diff --git a/examples/plugins/CMakeLists.txt b/examples/plugins/CMakeLists.txt index 6b4a14952b7b..e89bd44b41eb 100644 --- a/examples/plugins/CMakeLists.txt +++ b/examples/plugins/CMakeLists.txt @@ -24,7 +24,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/include/triton/Conversion/TritonGPUToLLVM/TypeConverter.h b/include/triton/Conversion/TritonGPUToLLVM/TypeConverter.h index 1adbbee4e309..6e8d0ec8c441 100644 --- a/include/triton/Conversion/TritonGPUToLLVM/TypeConverter.h +++ b/include/triton/Conversion/TritonGPUToLLVM/TypeConverter.h @@ -10,7 +10,8 @@ using namespace mlir; using namespace mlir::triton; -class TritonGPUToLLVMTypeConverter : public LLVMTypeConverter { +class __attribute__((visibility("default"))) TritonGPUToLLVMTypeConverter + : public LLVMTypeConverter { public: using TypeConverter::convertType; diff --git a/test/Plugins/test-dialect-plugin.mlir b/test/Plugins/test-dialect-plugin.mlir index 5751e94b5ccc..d384857da730 100644 --- a/test/Plugins/test-dialect-plugin.mlir +++ b/test/Plugins/test-dialect-plugin.mlir @@ -1,8 +1,6 @@ // 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: FileCheck %s -// REQUIRES: shared-libs - #blocked0 = #ttg.blocked<{sizePerThread = [8], threadsPerWarp = [32], warpsPerCTA = [8], order = [0]}> module attributes {"ttg.num-warps" = 8 : i32} { tt.func @convert_plugin() { diff --git a/test/Plugins/test-plugin.mlir b/test/Plugins/test-plugin.mlir index f16ef0788240..7539a858baf1 100644 --- a/test/Plugins/test-plugin.mlir +++ b/test/Plugins/test-plugin.mlir @@ -2,8 +2,6 @@ // 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: triton-opt -split-input-file %s | FileCheck %s -allow-unused-prefixes --check-prefix=CHECK-BASE -// REQUIRES: shared-libs - module attributes {"ttg.num-warps" = 4 : i32, "ttg.target" = "cuda:80"} { // CHECK-PLUGIN: func @foo() tt.func @bar() { diff --git a/third_party/nvidia/lib/TritonNVIDIAGPUToLLVM/TargetInfo.h b/third_party/nvidia/lib/TritonNVIDIAGPUToLLVM/TargetInfo.h index 9908244b9a47..a3ce015cac98 100644 --- a/third_party/nvidia/lib/TritonNVIDIAGPUToLLVM/TargetInfo.h +++ b/third_party/nvidia/lib/TritonNVIDIAGPUToLLVM/TargetInfo.h @@ -5,7 +5,8 @@ namespace mlir::triton::NVIDIA { -class TargetInfo : public mlir::triton::TargetInfoBase { +class __attribute__((visibility("default"))) TargetInfo + : public mlir::triton::TargetInfoBase { public: TargetInfo(int computeCapability, int ptxVersion) : computeCapability(computeCapability), ptxVersion(ptxVersion) {}