[triton-ext] Export triton+LLVM/MLIR symbols from triton-opt for use by plugin

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
 

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)

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)

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);
 

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

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;
 

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() {

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() {

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) {}