Skip to content

Route TLE GPU through TileIR conversion #75

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
kateyijian wants to merge 2 commits into
base: triton_v3.6.x
Choose a base branch
from
Open

Route TLE GPU through TileIR conversion #75

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
e86732c
Route TLE GPU through TileIR conversion
macqueen09 Aug 10, 2026
b6fe3db
Refine GPU CommonIR to TTGIR conversion
macqueen09 Aug 20, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
19 changes: 15 additions & 4 deletions CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -16,6 +16,13 @@ if(FLAGTREE_BACKEND)
add_definitions(-DFLAGTREE_BACKEND=\"${FLAGTREE_BACKEND}\")
endif()

# CommonIR is an explicit build-time feature. Keep the public contract as the
# C/C++ FLAGTREE_COMMON_IR macro rather than adding a CMake cache option.
set(FLAGTREE_COMMON_IR_ENABLED "$ENV{FLAGTREE_COMMON_IR}")
if(FLAGTREE_COMMON_IR_ENABLED)
add_compile_definitions(FLAGTREE_COMMON_IR)
endif()

set(FLAGTREE_TLE ON)
if(FLAGTREE_TLE)
add_definitions(-D__TLE__)
Expand Down Expand Up @@ -258,10 +265,14 @@ endif()
include_directories(${PROJECT_SOURCE_DIR}/third_party)
include_directories(${PROJECT_BINARY_DIR}/third_party) # Tablegen'd files

# TileIR dialect from FLIR is used by the TLE GPU -> Common IR POC path.
# Keep it optional at configure time so non-POC builds without FLIR are not
# forced to download the dependency.
if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/third_party/flir/CMakeLists.txt" AND NOT TARGET TileIRIR)
# TileIR is only part of builds that explicitly enable CommonIR.
if(FLAGTREE_COMMON_IR_ENABLED AND NOT TARGET TileIRIR)
if(FLAGTREE_BACKEND)
message(FATAL_ERROR "FLAGTREE_COMMON_IR on triton_v3.6.x requires the default NVIDIA backend")
endif()
if(NOT EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/third_party/flir/CMakeLists.txt")
message(FATAL_ERROR "FLAGTREE_COMMON_IR requires FLIR at third_party/flir")
endif()
add_definitions(-D__FLIR_TILEIR__)
include_directories(${CMAKE_CURRENT_SOURCE_DIR}/third_party/flir/include)
include_directories(${CMAKE_CURRENT_BINARY_DIR}/third_party/flir/include)
Expand Down
6 changes: 6 additions & 0 deletions bin/RegisterTritonDialects.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -17,6 +17,9 @@
#ifdef __FLIR_TILEIR__
#include "mlir-ext/Dialect/TileIR/IR/TileIRDialect.h"
#endif
#ifdef FLAGTREE_COMMON_IR
#include "nvidia/include/CommonIRToTTGIR/Passes.h"
#endif
#include "triton/Dialect/Gluon/Transforms/Passes.h"
#include "triton/Dialect/Triton/IR/Dialect.h"
#include "triton/Dialect/TritonGPU/IR/Dialect.h"
Expand Down Expand Up @@ -77,6 +80,9 @@ inline void registerTritonDialects(mlir::DialectRegistry &registry) {
mlir::triton::gluon::registerGluonPasses();
#ifdef __TLE__
mlir::triton::tle::registerPasses(); // flagtree tle
#endif
#ifdef FLAGTREE_COMMON_IR
mlir::triton::registerCommonIRToTTGIRPasses();
#endif
mlir::test::registerTestAliasPass();
mlir::test::registerTestAlignmentPass();
Expand Down
4 changes: 4 additions & 0 deletions python/triton/_common_ir.py
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,4 @@
from triton._C.libtriton import tle as tle_ir


ENABLED = tle_ir.is_common_ir_enabled()
Loading
Loading