diff --git a/CMakeLists.txt b/CMakeLists.txt
index 814206cfc..3326b3ff6 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -1,12 +1,12 @@
 cmake_minimum_required(VERSION 2.8)
-project(TDL)
+project(triton)
 include(CTest)
 
 # FLEX/YACC
 find_package(BISON)
 find_package(FLEX)
-BISON_TARGET(Parser ${CMAKE_CURRENT_SOURCE_DIR}/include/ast/parser.y ${CMAKE_CURRENT_BINARY_DIR}/parser.cpp)
-FLEX_TARGET(Lexer ${CMAKE_CURRENT_SOURCE_DIR}/include/ast/scanner.l ${CMAKE_CURRENT_BINARY_DIR}/scanner.cpp)
+BISON_TARGET(Parser ${CMAKE_CURRENT_SOURCE_DIR}/include/triton/ast/parser.y ${CMAKE_CURRENT_BINARY_DIR}/parser.cpp)
+FLEX_TARGET(Lexer ${CMAKE_CURRENT_SOURCE_DIR}/include/triton/ast/scanner.l ${CMAKE_CURRENT_BINARY_DIR}/scanner.cpp)
 get_filename_component(BISON_Parser_INCLUDE_DIRECTORIES ${BISON_Parser_OUTPUT_HEADER} DIRECTORY)
 include_directories(${BISON_Parser_INCLUDE_DIRECTORIES})
 
@@ -31,10 +31,10 @@ add_custom_target( ALL SOURCES ${ALL_SRC} )
 include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include)
 set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${LLVM_CXXFLAGS} -std=c++11")
 
-# TDL
-file(GLOB_RECURSE LIBTDL_SRC lib/*.cpp)
-add_library(tdl SHARED ${LIBTDL_SRC} ${BISON_Parser_OUTPUTS} ${FLEX_Lexer_OUTPUTS})
-target_link_libraries(tdl LLVM)
+# Triton
+file(GLOB_RECURSE LIBTRITON_SRC lib/*.cpp)
+add_library(triton SHARED ${LIBTRITON_SRC} ${BISON_Parser_OUTPUTS} ${FLEX_Lexer_OUTPUTS})
+target_link_libraries(triton LLVM)
 
 # Examples
 add_subdirectory(examples)
diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt
index 8419125c2..4a235d45a 100644
--- a/examples/CMakeLists.txt
+++ b/examples/CMakeLists.txt
@@ -2,5 +2,5 @@ foreach(PROG matrix)
      add_executable(${PROG} ${PROG}.cpp)
      set_target_properties(${PROG} PROPERTIES OUTPUT_NAME ${PROG})
      include_directories(/usr/local/cuda/include/)
-     target_link_libraries(${PROG} tdl cuda)
+     target_link_libraries(${PROG} triton cuda)
 endforeach(PROG)
diff --git a/examples/matrix.cpp b/examples/matrix.cpp
index 9c425c6f3..e0fd34646 100644
--- a/examples/matrix.cpp
+++ b/examples/matrix.cpp
@@ -2,19 +2,19 @@
 #include <cstdio>
 #include "cuda.h"
 #include "llvm/IR/Verifier.h"
-#include "ast/ast.h"
-#include "ir/context.h"
-#include "ir/module.h"
-#include "ir/print.h"
-#include "ir/context_impl.h"
-#include "codegen/selection.h"
-#include "codegen/tune.h"
-#include "codegen/shared_copy.h"
-#include "codegen/allocation.h"
-#include "codegen/liveness.h"
-#include "codegen/vectorize.h"
-#include "codegen/buffer_info.h"
-#include "codegen/barriers.h"
+#include "triton/ast/ast.h"
+#include "triton/ir/context.h"
+#include "triton/ir/module.h"
+#include "triton/ir/print.h"
+#include "triton/ir/context_impl.h"
+#include "triton/codegen/selection.h"
+#include "triton/codegen/tune.h"
+#include "triton/codegen/shared_copy.h"
+#include "triton/codegen/allocation.h"
+#include "triton/codegen/liveness.h"
+#include "triton/codegen/vectorize.h"
+#include "triton/codegen/buffer_info.h"
+#include "triton/codegen/barriers.h"
 #include "llvm/IR/IRPrintingPasses.h"
 #include "llvm/IR/Module.h"
 #include "llvm/IR/LLVMContext.h"
@@ -40,33 +40,35 @@ const char src[] =
 "\
 const tunable int32 TM;\
 const tunable int32 TN;\
-void test(fp32 *a, fp32 *b, fp32 *c, int32 M, int32 N, int32 K, int32 bound){\
+const tunable int32 TK;\
+\
+void matmul(fp32 *a, fp32 *b, fp32 *c, int32 M, int32 N, int32 K, int32 bound){\
   int32 rxa[TM] = get_global_range[TM](0);\
   int32 ryb[TN] = get_global_range[TN](1);\
-  int32 rka[8] = 0 ... 8;\
-  int32 rkb[8] = 0 ... 8;\
+  int32 rka[TK] = 0 ... TK;\
+  int32 rkb[TK] = 0 ... TK;\
   int32 rxc[TM] = get_global_range[TM](0);\
   int32 ryc[TN] = get_global_range[TN](1);\
   fp32 C[TM, TN] = 0;\
   int32 k;\
-  fp32* pa[TM, 8] = a + rxa[:, newaxis] + rka[newaxis, :]*M;\
-  fp32* pb[TN, 8] = b + ryb[:, newaxis] + rkb[newaxis, :]*K;\
+  fp32* pa[TM, TK] = a + rxa[:, newaxis] + rka[newaxis, :]*M;\
+  fp32* pb[TN, TK] = b + ryb[:, newaxis] + rkb[newaxis, :]*K;\
   fp32* pc[TM, TN] = c + rxc[:, newaxis] + ryc[newaxis, :]*M;\
-  fp32 a[TM, 8] = *pa;\
-  fp32 b[TN, 8] = *pb;\
+  fp32 a[TM, TK] = *pa;\
+  fp32 b[TN, TK] = *pb;\
   int1 checkc0[TM] = rxc < M;\
   int1 checkc1[TN] = ryc < N;\
   int1 checkc[TM, TN] = checkc0[:, newaxis] && checkc1[newaxis, :];\
-  for(k = K; k > 0; k = k - 8){\
-    int1 checka[TM, 8] = (k > 8);\
-    int1 checkb[TN, 8] = (k > 8);\
+  for(k = K; k > 0; k = k - TK){\
+    int1 checka[TM, TK] = (k > 8);\
+    int1 checkb[TN, TK] = (k > 8);\
     int1 checka0[TM];\
-    int1 checka1[8];\
+    int1 checka1[TK];\
     int1 checkb0[TN];\
-    int1 checkb1[8];\
+    int1 checkb1[TK];\
     C = dot(a, b, C);\
-    pa = pa + 8*M;\
-    pb = pb + 8*K;\
+    pa = pa + TK*M;\
+    pb = pb + TK*K;\
     @checka a = *pa;\
     @checkb b = *pb;\
     if(k > 8)\
@@ -171,6 +173,24 @@ void simple_gemm(std::vector<T> &c, const std::vector<T> &a, const std::vector<T
   }
 }
 
+void loop_nest(std::vector<size_t> const & ranges, std::function<void(std::vector<size_t> const &)> const & f){
+  size_t D = ranges.size();
+  std::vector<size_t> values(D, 0);
+  // Start with innermost loop
+  size_t i = D - 1;
+  while(true){
+    //Execute function
+    f(values);
+    //Increment counters
+    while(values[i]++ == ranges[i] - 1){
+      if(i == 0)
+        return;
+      values[i--] = 0;
+    }
+    i = D - 1;
+  }
+}
+
 int main() {
   // create AST from Triton-C source
   YY_BUFFER_STATE buffer = yy_scan_string(src);
@@ -183,11 +203,9 @@ int main() {
   tdl::ir::module module("matrix", context);
   program->codegen(&module);
   llvm::LLVMContext llvm_context;
-  llvm::Module llvm_module("test", llvm_context);
+  llvm::Module llvm_module("matmul", llvm_context);
+
 
-  context.p_impl->mp_constants_[0]->set_value(16);
-  context.p_impl->mp_constants_[1]->set_value(16);
-//  context.p_impl->mp_constants_[2]->set_value(8);
 
   // create passes
   tdl::codegen::buffer_info_pass buffer_info;
@@ -202,6 +220,8 @@ int main() {
   // tuning parameters
   tune.run(module);
   std::vector<unsigned> params = {
+    // shapes
+    16, 16, 8,
     // a0
     2, 8, 1,
     // b0
@@ -215,10 +235,15 @@ int main() {
     // b1
     1, 8, 1
   };
-  std::map<tdl::ir::value*, std::vector<std::string>> errors;
+  // meta-parameters
   unsigned i = 0;
+  context.p_impl->mp_constants_[0]->set_value(params[0]);
+  context.p_impl->mp_constants_[1]->set_value(params[1]);
+  context.p_impl->mp_constants_[2]->set_value(params[2]);
   for(unsigned *x: tune.get_params(module))
-    *x = params[i++];
+    *x = params[3 + i++];
+  // constraints
+  std::map<tdl::ir::value*, std::vector<std::string>> errors;
   tune.check_constraints(module, errors);
   std::cout << "errors: " << errors.size() << std::endl;
   for(auto &x: errors){
@@ -255,7 +280,7 @@ int main() {
   CUfunction cu_kernel;
   CUstream cu_stream;
   int major, minor;
-  compile_machine_code(cu_device, cu_context, cu_module, cu_kernel, cu_stream, major, minor, src, "test");
+  compile_machine_code(cu_device, cu_context, cu_module, cu_kernel, cu_stream, major, minor, src, "matmul");
 
   // execute machine code
   // Allocate buffers
@@ -284,8 +309,8 @@ int main() {
   void *args[] = { &d_a, &d_b, &d_c, &M, &N, &K, &bound};
   int num_regs;
   cuFuncGetAttribute(&num_regs, CU_FUNC_ATTRIBUTE_NUM_REGS, cu_kernel);
-  unsigned TM = 16;
-  unsigned TN = 16;
+  unsigned TM = context.p_impl->mp_constants_[0]->get_value();
+  unsigned TN = context.p_impl->mp_constants_[1]->get_value();
   unsigned nthreads = 32;
   checkCudaErrors(cuLaunchKernel(cu_kernel, (M + TM - 1)/TM, (N + TN - 1)/TN, 1, nthreads, 1, 1, 0, cu_stream, args, NULL));
   checkCudaErrors(cuStreamSynchronize(cu_stream));
diff --git a/include/ast/ast.h b/include/triton/ast/ast.h
similarity index 100%
rename from include/ast/ast.h
rename to include/triton/ast/ast.h
diff --git a/include/ast/parser.y b/include/triton/ast/parser.y
similarity index 98%
rename from include/ast/parser.y
rename to include/triton/ast/parser.y
index 960ae25a5..3fdcbd60b 100644
--- a/include/ast/parser.y
+++ b/include/triton/ast/parser.y
@@ -6,7 +6,7 @@ class node;
 }
 using namespace tdl::ast;
 #define YYSTYPE node*
-#include "../include/ast/ast.h"
+#include "../include/triton/ast/ast.h"
 
 extern char* yytext;
 void yyerror(const char *s);
@@ -117,7 +117,7 @@ builtin
 primary_expression
   : identifier         { $$ = new named_expression($1); }
   | constant           { $$ = $1; }
-  | constant ELLIPSIS constant { $$ = new constant_range($1, $3); }
+  | primary_expression ELLIPSIS primary_expression { $$ = new constant_range($1, $3); }
   | builtin            { $$ = $1; }
   | STRING_LITERAL     { $$ = new string_literal(yytext); }
   | '(' expression ')' { $$ = $2; }
diff --git a/include/ast/scanner.l b/include/triton/ast/scanner.l
similarity index 100%
rename from include/ast/scanner.l
rename to include/triton/ast/scanner.l
diff --git a/include/codegen/allocation.h b/include/triton/codegen/allocation.h
similarity index 100%
rename from include/codegen/allocation.h
rename to include/triton/codegen/allocation.h
diff --git a/include/codegen/barriers.h b/include/triton/codegen/barriers.h
similarity index 100%
rename from include/codegen/barriers.h
rename to include/triton/codegen/barriers.h
diff --git a/include/codegen/buffer_info.h b/include/triton/codegen/buffer_info.h
similarity index 100%
rename from include/codegen/buffer_info.h
rename to include/triton/codegen/buffer_info.h
diff --git a/include/codegen/layout.h b/include/triton/codegen/layout.h
similarity index 100%
rename from include/codegen/layout.h
rename to include/triton/codegen/layout.h
diff --git a/include/codegen/liveness.h b/include/triton/codegen/liveness.h
similarity index 100%
rename from include/codegen/liveness.h
rename to include/triton/codegen/liveness.h
diff --git a/include/codegen/selection.h b/include/triton/codegen/selection.h
similarity index 96%
rename from include/codegen/selection.h
rename to include/triton/codegen/selection.h
index ec733ed57..3f515ee4d 100644
--- a/include/codegen/selection.h
+++ b/include/triton/codegen/selection.h
@@ -3,11 +3,11 @@
 
 #include "llvm/IR/Module.h"
 #include "llvm/IR/IRBuilder.h"
-#include "ir/context.h"
-#include "ir/module.h"
-#include "ir/function.h"
-#include "ir/type.h"
-#include "codegen/buffer_info.h"
+#include "triton/ir/context.h"
+#include "triton/ir/module.h"
+#include "triton/ir/function.h"
+#include "triton/ir/type.h"
+#include "triton/codegen/buffer_info.h"
 
 
 namespace llvm{
diff --git a/include/codegen/shared_copy.h b/include/triton/codegen/shared_copy.h
similarity index 100%
rename from include/codegen/shared_copy.h
rename to include/triton/codegen/shared_copy.h
diff --git a/include/codegen/tune.h b/include/triton/codegen/tune.h
similarity index 100%
rename from include/codegen/tune.h
rename to include/triton/codegen/tune.h
diff --git a/include/codegen/vectorize.h b/include/triton/codegen/vectorize.h
similarity index 100%
rename from include/codegen/vectorize.h
rename to include/triton/codegen/vectorize.h
diff --git a/include/driver/backend.h b/include/triton/driver/backend.h
similarity index 100%
rename from include/driver/backend.h
rename to include/triton/driver/backend.h
diff --git a/include/driver/buffer.h b/include/triton/driver/buffer.h
similarity index 95%
rename from include/driver/buffer.h
rename to include/triton/driver/buffer.h
index 475cf2273..1d4130cd0 100755
--- a/include/driver/buffer.h
+++ b/include/triton/driver/buffer.h
@@ -23,8 +23,8 @@
 #ifndef TDL_INCLUDE_DRIVER_BUFFER_H
 #define TDL_INCLUDE_DRIVER_BUFFER_H
 
-#include "driver/handle.h"
-#include "driver/context.h"
+#include "triton/driver/handle.h"
+#include "triton/driver/context.h"
 
 namespace tdl
 {
diff --git a/include/driver/context.h b/include/triton/driver/context.h
similarity index 96%
rename from include/driver/context.h
rename to include/triton/driver/context.h
index bd98faded..339a25c72 100755
--- a/include/driver/context.h
+++ b/include/triton/driver/context.h
@@ -23,8 +23,8 @@
 #ifndef TDL_INCLUDE_DRIVER_CONTEXT_H
 #define TDL_INCLUDE_DRIVER_CONTEXT_H
 
-#include "driver/device.h"
-#include "driver/handle.h"
+#include "triton/driver/device.h"
+#include "triton/driver/handle.h"
 
 namespace tdl
 {
diff --git a/include/driver/cublas.h b/include/triton/driver/cublas.h
similarity index 97%
rename from include/driver/cublas.h
rename to include/triton/driver/cublas.h
index 9e1688a97..175b7f089 100755
--- a/include/driver/cublas.h
+++ b/include/triton/driver/cublas.h
@@ -24,13 +24,13 @@
 #define TDL_INCLUDE_DRIVER_CUBLAS_H
 
 #include "isaac/templates/common.hpp"
-#include "driver/dispatch.h"
-#include "driver/buffer.h"
-#include "driver/stream.h"
-#include "driver/backend.h"
-#include "driver/error.h"
-#include "tools/bench.hpp"
-#include "tools/collections.hpp"
+#include "triton/driver/dispatch.h"
+#include "triton/driver/buffer.h"
+#include "triton/driver/stream.h"
+#include "triton/driver/backend.h"
+#include "triton/driver/error.h"
+#include "triton/tools/bench.hpp"
+#include "triton/tools/collections.hpp"
 
 namespace tdl
 {
diff --git a/include/driver/device.h b/include/triton/driver/device.h
similarity index 97%
rename from include/driver/device.h
rename to include/triton/driver/device.h
index cffaf64b2..3be7ca04f 100755
--- a/include/driver/device.h
+++ b/include/triton/driver/device.h
@@ -23,8 +23,8 @@
 #ifndef TDL_INCLUDE_DRIVER_DEVICE_H
 #define TDL_INCLUDE_DRIVER_DEVICE_H
 
-#include "driver/platform.h"
-#include "driver/handle.h"
+#include "triton/driver/platform.h"
+#include "triton/driver/handle.h"
 
 namespace tdl
 {
diff --git a/include/driver/dispatch.h b/include/triton/driver/dispatch.h
similarity index 98%
rename from include/driver/dispatch.h
rename to include/triton/driver/dispatch.h
index 910fdc001..42ce6729f 100755
--- a/include/driver/dispatch.h
+++ b/include/triton/driver/dispatch.h
@@ -27,11 +27,11 @@
 #include <dlfcn.h>
 
 //CUDA Backend
-#include "external/CUDA/cuda.h"
-#include "external/CUDA/nvrtc.h"
-#include "external/CUDA/cublas_v2.h"
-#include "external/CUDA/cudnn.h"
-#include "external/CUDA/nvml.h"
+#include "triton/external/CUDA/cuda.h"
+#include "triton/external/CUDA/nvrtc.h"
+#include "triton/external/CUDA/cublas_v2.h"
+#include "triton/external/CUDA/cudnn.h"
+#include "triton/external/CUDA/nvml.h"
 
 //Exceptions
 #include <iostream>
diff --git a/include/driver/error.h b/include/triton/driver/error.h
similarity index 99%
rename from include/driver/error.h
rename to include/triton/driver/error.h
index d1589aad5..b837dea92 100755
--- a/include/driver/error.h
+++ b/include/triton/driver/error.h
@@ -24,7 +24,7 @@
 #define TDL_INCLUDE_DRIVER_ERROR_H
 
 #include <exception>
-#include "driver/dispatch.h"
+#include "triton/driver/dispatch.h"
 
 
 namespace tdl
diff --git a/include/driver/event.h b/include/triton/driver/event.h
similarity index 97%
rename from include/driver/event.h
rename to include/triton/driver/event.h
index 23f2c557f..79fbbb56f 100755
--- a/include/driver/event.h
+++ b/include/triton/driver/event.h
@@ -23,7 +23,7 @@
 #ifndef TDL_INCLUDE_DRIVER_EVENT_H
 #define TDL_INCLUDE_DRIVER_EVENT_H
 
-#include "driver/handle.h"
+#include "triton/driver/handle.h"
 
 namespace tdl
 {
diff --git a/include/driver/handle.h b/include/triton/driver/handle.h
similarity index 98%
rename from include/driver/handle.h
rename to include/triton/driver/handle.h
index eb7c90705..cb8463584 100755
--- a/include/driver/handle.h
+++ b/include/triton/driver/handle.h
@@ -27,7 +27,7 @@
 #include <iostream>
 #include <functional>
 #include <type_traits>
-#include "driver/dispatch.h"
+#include "triton/driver/dispatch.h"
 
 namespace tdl
 {
diff --git a/include/driver/kernel.h b/include/triton/driver/kernel.h
similarity index 96%
rename from include/driver/kernel.h
rename to include/triton/driver/kernel.h
index 60d4dc108..1fbf7935a 100755
--- a/include/driver/kernel.h
+++ b/include/triton/driver/kernel.h
@@ -23,8 +23,8 @@
 #ifndef TDL_INCLUDE_DRIVER_KERNEL_H
 #define TDL_INCLUDE_DRIVER_KERNEL_H
 
-#include "driver/module.h"
-#include "driver/handle.h"
+#include "triton/driver/module.h"
+#include "triton/driver/handle.h"
 
 #include <memory>
 
diff --git a/include/driver/module.h b/include/triton/driver/module.h
similarity index 94%
rename from include/driver/module.h
rename to include/triton/driver/module.h
index 2a1093233..913e90853 100755
--- a/include/driver/module.h
+++ b/include/triton/driver/module.h
@@ -24,9 +24,9 @@
 #define TDL_INCLUDE_DRIVER_MODULE_H
 
 #include <map>
-#include "driver/handle.h"
-#include "driver/context.h"
-#include "driver/buffer.h"
+#include "triton/driver/handle.h"
+#include "triton/driver/context.h"
+#include "triton/driver/buffer.h"
 
 namespace tdl
 {
diff --git a/include/driver/platform.h b/include/triton/driver/platform.h
similarity index 97%
rename from include/driver/platform.h
rename to include/triton/driver/platform.h
index add506e82..d39c48e72 100755
--- a/include/driver/platform.h
+++ b/include/triton/driver/platform.h
@@ -26,7 +26,7 @@
 #include <vector>
 #include <string>
 
-#include "driver/handle.h"
+#include "triton/driver/handle.h"
 
 namespace tdl
 {
diff --git a/include/driver/stream.h b/include/triton/driver/stream.h
similarity index 94%
rename from include/driver/stream.h
rename to include/triton/driver/stream.h
index 5ff59356c..8e5783892 100755
--- a/include/driver/stream.h
+++ b/include/triton/driver/stream.h
@@ -24,10 +24,10 @@
 #define TDL_INCLUDE_DRIVER_STREAM_H
 
 #include <map>
-#include "driver/context.h"
-#include "driver/device.h"
-#include "driver/handle.h"
-#include "driver/buffer.h"
+#include "triton/driver/context.h"
+#include "triton/driver/device.h"
+#include "triton/driver/handle.h"
+#include "triton/driver/buffer.h"
 
 namespace tdl
 {
diff --git a/include/external/CUDA/builtin_types.h b/include/triton/external/CUDA/builtin_types.h
similarity index 100%
rename from include/external/CUDA/builtin_types.h
rename to include/triton/external/CUDA/builtin_types.h
diff --git a/include/external/CUDA/channel_descriptor.h b/include/triton/external/CUDA/channel_descriptor.h
similarity index 100%
rename from include/external/CUDA/channel_descriptor.h
rename to include/triton/external/CUDA/channel_descriptor.h
diff --git a/include/external/CUDA/crt/host_config.h b/include/triton/external/CUDA/crt/host_config.h
similarity index 100%
rename from include/external/CUDA/crt/host_config.h
rename to include/triton/external/CUDA/crt/host_config.h
diff --git a/include/external/CUDA/crt/host_defines.h b/include/triton/external/CUDA/crt/host_defines.h
similarity index 100%
rename from include/external/CUDA/crt/host_defines.h
rename to include/triton/external/CUDA/crt/host_defines.h
diff --git a/include/external/CUDA/cuComplex.h b/include/triton/external/CUDA/cuComplex.h
similarity index 100%
rename from include/external/CUDA/cuComplex.h
rename to include/triton/external/CUDA/cuComplex.h
diff --git a/include/external/CUDA/cublas.h b/include/triton/external/CUDA/cublas.h
similarity index 100%
rename from include/external/CUDA/cublas.h
rename to include/triton/external/CUDA/cublas.h
diff --git a/include/external/CUDA/cublas_api.h b/include/triton/external/CUDA/cublas_api.h
similarity index 100%
rename from include/external/CUDA/cublas_api.h
rename to include/triton/external/CUDA/cublas_api.h
diff --git a/include/external/CUDA/cublas_v2.h b/include/triton/external/CUDA/cublas_v2.h
similarity index 100%
rename from include/external/CUDA/cublas_v2.h
rename to include/triton/external/CUDA/cublas_v2.h
diff --git a/include/external/CUDA/cuda.h b/include/triton/external/CUDA/cuda.h
similarity index 100%
rename from include/external/CUDA/cuda.h
rename to include/triton/external/CUDA/cuda.h
diff --git a/include/external/CUDA/cuda_device_runtime_api.h b/include/triton/external/CUDA/cuda_device_runtime_api.h
similarity index 100%
rename from include/external/CUDA/cuda_device_runtime_api.h
rename to include/triton/external/CUDA/cuda_device_runtime_api.h
diff --git a/include/external/CUDA/cuda_fp16.h b/include/triton/external/CUDA/cuda_fp16.h
similarity index 100%
rename from include/external/CUDA/cuda_fp16.h
rename to include/triton/external/CUDA/cuda_fp16.h
diff --git a/include/external/CUDA/cuda_fp16.hpp b/include/triton/external/CUDA/cuda_fp16.hpp
similarity index 100%
rename from include/external/CUDA/cuda_fp16.hpp
rename to include/triton/external/CUDA/cuda_fp16.hpp
diff --git a/include/external/CUDA/cuda_runtime.h b/include/triton/external/CUDA/cuda_runtime.h
similarity index 100%
rename from include/external/CUDA/cuda_runtime.h
rename to include/triton/external/CUDA/cuda_runtime.h
diff --git a/include/external/CUDA/cuda_runtime_api.h b/include/triton/external/CUDA/cuda_runtime_api.h
similarity index 100%
rename from include/external/CUDA/cuda_runtime_api.h
rename to include/triton/external/CUDA/cuda_runtime_api.h
diff --git a/include/external/CUDA/cudnn.h b/include/triton/external/CUDA/cudnn.h
similarity index 100%
rename from include/external/CUDA/cudnn.h
rename to include/triton/external/CUDA/cudnn.h
diff --git a/include/external/CUDA/cusparse.h b/include/triton/external/CUDA/cusparse.h
similarity index 100%
rename from include/external/CUDA/cusparse.h
rename to include/triton/external/CUDA/cusparse.h
diff --git a/include/external/CUDA/device_types.h b/include/triton/external/CUDA/device_types.h
similarity index 100%
rename from include/external/CUDA/device_types.h
rename to include/triton/external/CUDA/device_types.h
diff --git a/include/external/CUDA/driver_functions.h b/include/triton/external/CUDA/driver_functions.h
similarity index 100%
rename from include/external/CUDA/driver_functions.h
rename to include/triton/external/CUDA/driver_functions.h
diff --git a/include/external/CUDA/driver_types.h b/include/triton/external/CUDA/driver_types.h
similarity index 100%
rename from include/external/CUDA/driver_types.h
rename to include/triton/external/CUDA/driver_types.h
diff --git a/include/external/CUDA/host_config.h b/include/triton/external/CUDA/host_config.h
similarity index 100%
rename from include/external/CUDA/host_config.h
rename to include/triton/external/CUDA/host_config.h
diff --git a/include/external/CUDA/host_defines.h b/include/triton/external/CUDA/host_defines.h
similarity index 100%
rename from include/external/CUDA/host_defines.h
rename to include/triton/external/CUDA/host_defines.h
diff --git a/include/external/CUDA/library_types.h b/include/triton/external/CUDA/library_types.h
similarity index 100%
rename from include/external/CUDA/library_types.h
rename to include/triton/external/CUDA/library_types.h
diff --git a/include/external/CUDA/nvml.h b/include/triton/external/CUDA/nvml.h
similarity index 100%
rename from include/external/CUDA/nvml.h
rename to include/triton/external/CUDA/nvml.h
diff --git a/include/external/CUDA/nvrtc.h b/include/triton/external/CUDA/nvrtc.h
similarity index 100%
rename from include/external/CUDA/nvrtc.h
rename to include/triton/external/CUDA/nvrtc.h
diff --git a/include/external/CUDA/surface_types.h b/include/triton/external/CUDA/surface_types.h
similarity index 100%
rename from include/external/CUDA/surface_types.h
rename to include/triton/external/CUDA/surface_types.h
diff --git a/include/external/CUDA/texture_types.h b/include/triton/external/CUDA/texture_types.h
similarity index 100%
rename from include/external/CUDA/texture_types.h
rename to include/triton/external/CUDA/texture_types.h
diff --git a/include/external/CUDA/vector_functions.h b/include/triton/external/CUDA/vector_functions.h
similarity index 100%
rename from include/external/CUDA/vector_functions.h
rename to include/triton/external/CUDA/vector_functions.h
diff --git a/include/external/CUDA/vector_functions.hpp b/include/triton/external/CUDA/vector_functions.hpp
similarity index 100%
rename from include/external/CUDA/vector_functions.hpp
rename to include/triton/external/CUDA/vector_functions.hpp
diff --git a/include/external/CUDA/vector_types.h b/include/triton/external/CUDA/vector_types.h
similarity index 100%
rename from include/external/CUDA/vector_types.h
rename to include/triton/external/CUDA/vector_types.h
diff --git a/include/ir/basic_block.h b/include/triton/ir/basic_block.h
similarity index 100%
rename from include/ir/basic_block.h
rename to include/triton/ir/basic_block.h
diff --git a/include/ir/builder.h b/include/triton/ir/builder.h
similarity index 100%
rename from include/ir/builder.h
rename to include/triton/ir/builder.h
diff --git a/include/ir/constant.h b/include/triton/ir/constant.h
similarity index 100%
rename from include/ir/constant.h
rename to include/triton/ir/constant.h
diff --git a/include/ir/context.h b/include/triton/ir/context.h
similarity index 90%
rename from include/ir/context.h
rename to include/triton/ir/context.h
index c7382a0cb..d3018aa6f 100644
--- a/include/ir/context.h
+++ b/include/triton/ir/context.h
@@ -2,7 +2,7 @@
 #define TDL_INCLUDE_IR_CONTEXT_H
 
 #include <memory>
-#include "ir/type.h"
+#include "triton/ir/type.h"
 
 namespace tdl{
 namespace ir{
diff --git a/include/ir/context_impl.h b/include/triton/ir/context_impl.h
similarity index 97%
rename from include/ir/context_impl.h
rename to include/triton/ir/context_impl.h
index b9017b39c..623f58e40 100644
--- a/include/ir/context_impl.h
+++ b/include/triton/ir/context_impl.h
@@ -3,7 +3,7 @@
 
 #include <memory>
 #include <map>
-#include "ir/type.h"
+#include "triton/ir/type.h"
 
 namespace tdl{
 namespace ir{
diff --git a/include/ir/function.h b/include/triton/ir/function.h
similarity index 100%
rename from include/ir/function.h
rename to include/triton/ir/function.h
diff --git a/include/ir/instructions.h b/include/triton/ir/instructions.h
similarity index 99%
rename from include/ir/instructions.h
rename to include/triton/ir/instructions.h
index ae752f78e..93dc0b78e 100644
--- a/include/ir/instructions.h
+++ b/include/triton/ir/instructions.h
@@ -3,7 +3,7 @@
 
 #include <vector>
 #include "value.h"
-#include "ir/type.h"
+#include "triton/ir/type.h"
 #include "llvm/IR/Instructions.h"
 
 namespace tdl{
diff --git a/include/ir/module.h b/include/triton/ir/module.h
similarity index 100%
rename from include/ir/module.h
rename to include/triton/ir/module.h
diff --git a/include/triton/ir/print.h b/include/triton/ir/print.h
new file mode 100644
index 000000000..a31929282
--- /dev/null
+++ b/include/triton/ir/print.h
@@ -0,0 +1,17 @@
+#ifndef TDL_INCLUDE_IR_PRINT_H
+#define TDL_INCLUDE_IR_PRINT_H
+
+
+#include "builder.h"
+
+namespace tdl{
+namespace ir{
+
+class module;
+
+void print(module &mod, std::ostream& os);
+
+}
+}
+
+#endif
diff --git a/include/ir/type.h b/include/triton/ir/type.h
similarity index 100%
rename from include/ir/type.h
rename to include/triton/ir/type.h
diff --git a/include/ir/value.h b/include/triton/ir/value.h
similarity index 100%
rename from include/ir/value.h
rename to include/triton/ir/value.h
diff --git a/include/tools/sys/getenv.hpp b/include/triton/tools/sys/getenv.hpp
similarity index 100%
rename from include/tools/sys/getenv.hpp
rename to include/triton/tools/sys/getenv.hpp
diff --git a/include/tools/sys/mkdir.hpp b/include/triton/tools/sys/mkdir.hpp
similarity index 100%
rename from include/tools/sys/mkdir.hpp
rename to include/triton/tools/sys/mkdir.hpp
diff --git a/lib/ast/lowering.cpp b/lib/ast/lowering.cpp
index 714d8b16f..4d1fa5048 100644
--- a/lib/ast/lowering.cpp
+++ b/lib/ast/lowering.cpp
@@ -1,12 +1,12 @@
 #include <functional>
 #include <algorithm>
-#include "ast/ast.h"
-#include "ir/constant.h"
-#include "ir/function.h"
-#include "ir/module.h"
-#include "ir/basic_block.h"
-#include "ir/builder.h"
-#include "ir/type.h"
+#include "triton/ast/ast.h"
+#include "triton/ir/constant.h"
+#include "triton/ir/function.h"
+#include "triton/ir/module.h"
+#include "triton/ir/basic_block.h"
+#include "triton/ir/builder.h"
+#include "triton/ir/type.h"
 #include <iostream>
 
 
diff --git a/lib/codegen/allocation.cpp b/lib/codegen/allocation.cpp
index 696b46cb9..d396ec790 100644
--- a/lib/codegen/allocation.cpp
+++ b/lib/codegen/allocation.cpp
@@ -1,12 +1,12 @@
-#include "codegen/allocation.h"
-#include "codegen/liveness.h"
-#include "codegen/layout.h"
-#include "codegen/buffer_info.h"
-#include "ir/basic_block.h"
-#include "ir/type.h"
-#include "ir/value.h"
-#include "ir/function.h"
-#include "ir/instructions.h"
+#include "triton/codegen/allocation.h"
+#include "triton/codegen/liveness.h"
+#include "triton/codegen/layout.h"
+#include "triton/codegen/buffer_info.h"
+#include "triton/ir/basic_block.h"
+#include "triton/ir/type.h"
+#include "triton/ir/value.h"
+#include "triton/ir/function.h"
+#include "triton/ir/instructions.h"
 
 namespace tdl{
 namespace codegen{
diff --git a/lib/codegen/barriers.cpp b/lib/codegen/barriers.cpp
index df017931b..c40c08186 100644
--- a/lib/codegen/barriers.cpp
+++ b/lib/codegen/barriers.cpp
@@ -1,11 +1,11 @@
 #include <algorithm>
-#include "codegen/barriers.h"
-#include "codegen/allocation.h"
-#include "codegen/buffer_info.h"
-#include "ir/module.h"
-#include "ir/function.h"
-#include "ir/basic_block.h"
-#include "ir/instructions.h"
+#include "triton/codegen/barriers.h"
+#include "triton/codegen/allocation.h"
+#include "triton/codegen/buffer_info.h"
+#include "triton/ir/module.h"
+#include "triton/ir/function.h"
+#include "triton/ir/basic_block.h"
+#include "triton/ir/instructions.h"
 
 namespace tdl {
 
diff --git a/lib/codegen/buffer_info.cpp b/lib/codegen/buffer_info.cpp
index 4d20fa6e9..c8fe08df6 100644
--- a/lib/codegen/buffer_info.cpp
+++ b/lib/codegen/buffer_info.cpp
@@ -1,9 +1,9 @@
-#include "codegen/buffer_info.h"
-#include "ir/module.h"
-#include "ir/function.h"
-#include "ir/basic_block.h"
-#include "ir/instructions.h"
-#include "ir/type.h"
+#include "triton/codegen/buffer_info.h"
+#include "triton/ir/module.h"
+#include "triton/ir/function.h"
+#include "triton/ir/basic_block.h"
+#include "triton/ir/instructions.h"
+#include "triton/ir/type.h"
 
 namespace tdl {
 
diff --git a/lib/codegen/layout.cpp b/lib/codegen/layout.cpp
index 58f81227a..040954caa 100644
--- a/lib/codegen/layout.cpp
+++ b/lib/codegen/layout.cpp
@@ -1,8 +1,8 @@
-#include "codegen/layout.h"
-#include "ir/function.h"
-#include "ir/module.h"
-#include "ir/basic_block.h"
-#include "ir/instructions.h"
+#include "triton/codegen/layout.h"
+#include "triton/ir/function.h"
+#include "triton/ir/module.h"
+#include "triton/ir/basic_block.h"
+#include "triton/ir/instructions.h"
 
 namespace tdl{
 namespace codegen{
diff --git a/lib/codegen/liveness.cpp b/lib/codegen/liveness.cpp
index 05b803f8f..f5e5a79c6 100644
--- a/lib/codegen/liveness.cpp
+++ b/lib/codegen/liveness.cpp
@@ -1,10 +1,10 @@
-#include "codegen/liveness.h"
-#include "codegen/buffer_info.h"
-#include "ir/basic_block.h"
-#include "ir/function.h"
-#include "ir/module.h"
-#include "ir/instructions.h"
-#include "ir/value.h"
+#include "triton/codegen/liveness.h"
+#include "triton/codegen/buffer_info.h"
+#include "triton/ir/basic_block.h"
+#include "triton/ir/function.h"
+#include "triton/ir/module.h"
+#include "triton/ir/instructions.h"
+#include "triton/ir/value.h"
 
 namespace tdl{
 namespace codegen{
diff --git a/lib/codegen/selection.cpp b/lib/codegen/selection.cpp
index a028134fa..4cb2cf827 100644
--- a/lib/codegen/selection.cpp
+++ b/lib/codegen/selection.cpp
@@ -1,13 +1,13 @@
-#include "codegen/selection.h"
-#include "codegen/tune.h"
-#include "codegen/allocation.h"
+#include "triton/codegen/selection.h"
+#include "triton/codegen/tune.h"
+#include "triton/codegen/allocation.h"
 #include "llvm/IR/InstrTypes.h"
 #include "llvm/IR/Module.h"
 #include "llvm/IR/IRBuilder.h"
-#include "ir/context.h"
-#include "ir/module.h"
-#include "ir/function.h"
-#include "ir/type.h"
+#include "triton/ir/context.h"
+#include "triton/ir/module.h"
+#include "triton/ir/function.h"
+#include "triton/ir/type.h"
 #include "llvm/Transforms/Scalar/EarlyCSE.h"
 #include "llvm/Analysis/LoopInfo.h"
 #include "llvm/Transforms/Utils/BasicBlockUtils.h"
diff --git a/lib/codegen/shared_copy.cpp b/lib/codegen/shared_copy.cpp
index 60c31199f..b7276bbfe 100644
--- a/lib/codegen/shared_copy.cpp
+++ b/lib/codegen/shared_copy.cpp
@@ -1,10 +1,10 @@
 #include <algorithm>
-#include "codegen/shared_copy.h"
-#include "codegen/buffer_info.h"
-#include "ir/module.h"
-#include "ir/function.h"
-#include "ir/basic_block.h"
-#include "ir/instructions.h"
+#include "triton/codegen/shared_copy.h"
+#include "triton/codegen/buffer_info.h"
+#include "triton/ir/module.h"
+#include "triton/ir/function.h"
+#include "triton/ir/basic_block.h"
+#include "triton/ir/instructions.h"
 
 namespace tdl {
 
diff --git a/lib/codegen/tune.cpp b/lib/codegen/tune.cpp
index 3dc5c4e87..302676697 100644
--- a/lib/codegen/tune.cpp
+++ b/lib/codegen/tune.cpp
@@ -1,9 +1,9 @@
-#include "codegen/tune.h"
-#include "ir/instructions.h"
-#include "ir/type.h"
-#include "ir/module.h"
-#include "ir/function.h"
-#include "ir/context_impl.h"
+#include "triton/codegen/tune.h"
+#include "triton/ir/instructions.h"
+#include "triton/ir/type.h"
+#include "triton/ir/module.h"
+#include "triton/ir/function.h"
+#include "triton/ir/context_impl.h"
 
 #include <cstdlib>
 
diff --git a/lib/codegen/vectorize.cpp b/lib/codegen/vectorize.cpp
index 41a1afd10..4c45928a2 100644
--- a/lib/codegen/vectorize.cpp
+++ b/lib/codegen/vectorize.cpp
@@ -1,9 +1,9 @@
-#include "codegen/vectorize.h"
-#include "codegen/tune.h"
-#include "ir/module.h"
-#include "ir/function.h"
-#include "ir/basic_block.h"
-#include "ir/instructions.h"
+#include "triton/codegen/vectorize.h"
+#include "triton/codegen/tune.h"
+#include "triton/ir/module.h"
+#include "triton/ir/function.h"
+#include "triton/ir/basic_block.h"
+#include "triton/ir/instructions.h"
 
 namespace tdl {
 
diff --git a/lib/driver/backend.cpp b/lib/driver/backend.cpp
index bddb419df..cdf72027d 100755
--- a/lib/driver/backend.cpp
+++ b/lib/driver/backend.cpp
@@ -20,12 +20,12 @@
 * SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
 */
 
-#include "driver/dispatch.h"
-#include "driver/backend.h"
-#include "driver/buffer.h"
-#include "driver/context.h"
-#include "driver/stream.h"
-#include "driver/kernel.h"
+#include "triton/driver/dispatch.h"
+#include "triton/driver/backend.h"
+#include "triton/driver/buffer.h"
+#include "triton/driver/context.h"
+#include "triton/driver/stream.h"
+#include "triton/driver/kernel.h"
 
 #include <assert.h>
 #include <stdexcept>
diff --git a/lib/driver/buffer.cpp b/lib/driver/buffer.cpp
index aa770a05d..069e8abe1 100755
--- a/lib/driver/buffer.cpp
+++ b/lib/driver/buffer.cpp
@@ -21,10 +21,10 @@
 */
 
 #include <iostream>
-#include "driver/stream.h"
-#include "driver/buffer.h"
-#include "driver/context.h"
-#include "driver/dispatch.h"
+#include "triton/driver/stream.h"
+#include "triton/driver/buffer.h"
+#include "triton/driver/context.h"
+#include "triton/driver/dispatch.h"
 
 
 namespace tdl
diff --git a/lib/driver/context.cpp b/lib/driver/context.cpp
index 9da2c6978..8e365db81 100755
--- a/lib/driver/context.cpp
+++ b/lib/driver/context.cpp
@@ -23,11 +23,11 @@
 #include <iostream>
 #include <cassert>
 
-#include "driver/context.h"
-#include "driver/module.h"
+#include "triton/driver/context.h"
+#include "triton/driver/module.h"
 
-#include "tools/sys/getenv.hpp"
-#include "tools/sys/mkdir.hpp"
+#include "triton/tools/sys/getenv.hpp"
+#include "triton/tools/sys/mkdir.hpp"
 
 namespace tdl
 {
diff --git a/lib/driver/device.cpp b/lib/driver/device.cpp
index 13f10f6a0..22f640d7b 100755
--- a/lib/driver/device.cpp
+++ b/lib/driver/device.cpp
@@ -26,7 +26,7 @@
 #include <cstring>
 #include <memory>
 
-#include "driver/device.h"
+#include "triton/driver/device.h"
 
 namespace tdl
 {
diff --git a/lib/driver/dispatch.cpp b/lib/driver/dispatch.cpp
index 2d0cd5232..d7d19727d 100755
--- a/lib/driver/dispatch.cpp
+++ b/lib/driver/dispatch.cpp
@@ -21,8 +21,8 @@
 */
 
 #include <map>
-#include "driver/dispatch.h"
-#include "driver/context.h"
+#include "triton/driver/dispatch.h"
+#include "triton/driver/context.h"
 
 namespace tdl
 {
diff --git a/lib/driver/error.cpp b/lib/driver/error.cpp
index 7e7dc9d75..9759a1323 100755
--- a/lib/driver/error.cpp
+++ b/lib/driver/error.cpp
@@ -20,7 +20,7 @@
 * SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
 */
 
-#include "driver/error.h"
+#include "triton/driver/error.h"
 
 namespace tdl
 {
diff --git a/lib/driver/event.cpp b/lib/driver/event.cpp
index dc554d808..b8841dc0f 100755
--- a/lib/driver/event.cpp
+++ b/lib/driver/event.cpp
@@ -20,7 +20,7 @@
 * SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
 */
 
-#include "driver/event.h"
+#include "triton/driver/event.h"
 
 namespace tdl
 {
diff --git a/lib/driver/handle.cpp b/lib/driver/handle.cpp
index a01a099bd..cd7ee4195 100755
--- a/lib/driver/handle.cpp
+++ b/lib/driver/handle.cpp
@@ -22,7 +22,7 @@
 
 #include <cassert>
 #include <memory>
-#include "driver/handle.h"
+#include "triton/driver/handle.h"
 
 namespace tdl
 {
diff --git a/lib/driver/kernel.cpp b/lib/driver/kernel.cpp
index 6e536b767..180e46cc7 100755
--- a/lib/driver/kernel.cpp
+++ b/lib/driver/kernel.cpp
@@ -23,8 +23,8 @@
 #include <iostream>
 #include <cstring>
 
-#include "driver/kernel.h"
-#include "driver/buffer.h"
+#include "triton/driver/kernel.h"
+#include "triton/driver/buffer.h"
 
 namespace tdl
 {
diff --git a/lib/driver/module.cpp b/lib/driver/module.cpp
index c61482cbc..d897489fe 100755
--- a/lib/driver/module.cpp
+++ b/lib/driver/module.cpp
@@ -23,11 +23,11 @@
 #include <iostream>
 #include <fstream>
 
-#include "driver/module.h"
-#include "driver/context.h"
-#include "driver/error.h"
+#include "triton/driver/module.h"
+#include "triton/driver/context.h"
+#include "triton/driver/error.h"
 
-#include "tools/sys/getenv.hpp"
+#include "triton/tools/sys/getenv.hpp"
 
 namespace tdl
 {
diff --git a/lib/driver/platform.cpp b/lib/driver/platform.cpp
index 2fa9933d3..ada13de41 100755
--- a/lib/driver/platform.cpp
+++ b/lib/driver/platform.cpp
@@ -21,8 +21,8 @@
 */
 
 
-#include "driver/platform.h"
-#include "driver/device.h"
+#include "triton/driver/platform.h"
+#include "triton/driver/device.h"
 
 #include <string>
 
diff --git a/lib/driver/stream.cpp b/lib/driver/stream.cpp
index 0296eba40..39996d473 100755
--- a/lib/driver/stream.cpp
+++ b/lib/driver/stream.cpp
@@ -24,13 +24,13 @@
 #include <cassert>
 #include <array>
 
-#include "driver/backend.h"
-#include "driver/stream.h"
-#include "driver/context.h"
-#include "driver/device.h"
-#include "driver/event.h"
-#include "driver/kernel.h"
-#include "driver/buffer.h"
+#include "triton/driver/backend.h"
+#include "triton/driver/stream.h"
+#include "triton/driver/context.h"
+#include "triton/driver/device.h"
+#include "triton/driver/event.h"
+#include "triton/driver/kernel.h"
+#include "triton/driver/buffer.h"
 
 namespace tdl
 {
diff --git a/lib/ir/basic_block.cpp b/lib/ir/basic_block.cpp
index c7d8493e8..15fa3188c 100644
--- a/lib/ir/basic_block.cpp
+++ b/lib/ir/basic_block.cpp
@@ -1,7 +1,7 @@
-#include "ir/basic_block.h"
-#include "ir/instructions.h"
-#include "ir/type.h"
-#include "ir/function.h"
+#include "triton/ir/basic_block.h"
+#include "triton/ir/instructions.h"
+#include "triton/ir/type.h"
+#include "triton/ir/function.h"
 
 namespace tdl {
 namespace ir {
diff --git a/lib/ir/builder.cpp b/lib/ir/builder.cpp
index c2b338236..467a7ef71 100644
--- a/lib/ir/builder.cpp
+++ b/lib/ir/builder.cpp
@@ -1,9 +1,9 @@
 #include <string>
-#include "ir/basic_block.h"
-#include "ir/builder.h"
-#include "ir/constant.h"
-#include "ir/instructions.h"
-#include "ir/type.h"
+#include "triton/ir/basic_block.h"
+#include "triton/ir/builder.h"
+#include "triton/ir/constant.h"
+#include "triton/ir/instructions.h"
+#include "triton/ir/type.h"
 #include "llvm/IR/Instruction.h"
 
 namespace tdl{
diff --git a/lib/ir/constant.cpp b/lib/ir/constant.cpp
index 87f669e4d..5a6cfbc7a 100644
--- a/lib/ir/constant.cpp
+++ b/lib/ir/constant.cpp
@@ -1,8 +1,8 @@
 #include <cassert>
-#include "ir/constant.h"
-#include "ir/type.h"
-#include "ir/context.h"
-#include "ir/context_impl.h"
+#include "triton/ir/constant.h"
+#include "triton/ir/type.h"
+#include "triton/ir/context.h"
+#include "triton/ir/context_impl.h"
 
 namespace tdl{
 namespace ir{
diff --git a/lib/ir/context.cpp b/lib/ir/context.cpp
index a3fd665e1..33a852bda 100644
--- a/lib/ir/context.cpp
+++ b/lib/ir/context.cpp
@@ -1,6 +1,6 @@
-#include "ir/context_impl.h"
-#include "ir/context.h"
-#include "ir/type.h"
+#include "triton/ir/context_impl.h"
+#include "triton/ir/context.h"
+#include "triton/ir/type.h"
 
 namespace tdl{
 namespace ir{
diff --git a/lib/ir/function.cpp b/lib/ir/function.cpp
index 6d3329f21..75a465e8f 100644
--- a/lib/ir/function.cpp
+++ b/lib/ir/function.cpp
@@ -1,6 +1,6 @@
-#include "ir/function.h"
-#include "ir/type.h"
-#include "ir/module.h"
+#include "triton/ir/function.h"
+#include "triton/ir/type.h"
+#include "triton/ir/module.h"
 
 namespace tdl{
 namespace ir{
diff --git a/lib/ir/instructions.cpp b/lib/ir/instructions.cpp
index 38adcc377..93ba4602a 100644
--- a/lib/ir/instructions.cpp
+++ b/lib/ir/instructions.cpp
@@ -1,8 +1,8 @@
-#include "ir/context.h"
-#include "ir/basic_block.h"
-#include "ir/instructions.h"
-#include "ir/constant.h"
-#include "ir/type.h"
+#include "triton/ir/context.h"
+#include "triton/ir/basic_block.h"
+#include "triton/ir/instructions.h"
+#include "triton/ir/constant.h"
+#include "triton/ir/type.h"
 
 namespace tdl{
 namespace ir{
diff --git a/lib/ir/module.cpp b/lib/ir/module.cpp
index 66a4784cc..28825fe4b 100644
--- a/lib/ir/module.cpp
+++ b/lib/ir/module.cpp
@@ -1,8 +1,8 @@
-#include "ir/basic_block.h"
-#include "ir/module.h"
-#include "ir/type.h"
-#include "ir/constant.h"
-#include "ir/function.h"
+#include "triton/ir/basic_block.h"
+#include "triton/ir/module.h"
+#include "triton/ir/type.h"
+#include "triton/ir/constant.h"
+#include "triton/ir/function.h"
 
 namespace tdl{
 namespace ir{
diff --git a/lib/ir/print.cpp b/lib/ir/print.cpp
new file mode 100644
index 000000000..34829aec0
--- /dev/null
+++ b/lib/ir/print.cpp
@@ -0,0 +1,59 @@
+#include "triton/ir/basic_block.h"
+#include "triton/ir/module.h"
+#include "triton/ir/type.h"
+#include "triton/ir/constant.h"
+#include "triton/ir/function.h"
+#include "triton/ir/instructions.h"
+#include "triton/ir/print.h"
+
+namespace tdl{
+namespace ir{
+
+std::string get_name(ir::value *v, unsigned i) {
+  if(v->get_name().empty()){
+    std::string name = "%" + std::to_string(i);
+    v->set_name(name);
+  }
+  return v->get_name();
+}
+
+
+void print(module &mod, std::ostream& os) {
+  unsigned cnt = 0;
+  for(ir::function *fn: mod.get_function_list()){
+    os << "{" << std::endl;
+    for(ir::basic_block *block: fn->blocks()){
+      auto const &predecessors = block->get_predecessors();
+      os << block->get_name() << ":";
+      if(!predecessors.empty()){
+        os << "                 ";
+        os << "; preds = ";
+        auto const &predecessors = block->get_predecessors();
+        for(ir::basic_block *pred: predecessors)
+          os << pred->get_name() << (pred!=predecessors.back()?", ":"");
+      }
+      os << std::endl;
+      for(ir::instruction *inst: block->get_inst_list()){
+        os << "  ";
+        if(ir::value *pred = inst->get_mask_pred())
+          os << "@" << get_name(pred, cnt++) << " ";
+        if(!inst->get_type()->is_void_ty())
+          os << get_name(inst, cnt++) << " = ";
+        os << inst->repr();
+        ir::instruction::ops_t ops = inst->ops();
+        size_t num_ops = inst->get_num_operands();
+        if(num_ops > 0)
+          os << " ";
+        for(unsigned i = 0; i < num_ops; i++)
+          os << get_name(ops[i], cnt++) << (i < num_ops - 1?", ":"");
+        os << ";" << std::endl;
+      }
+      os << std::endl;
+    }
+    os << "}" << std::endl;
+  }
+}
+
+
+}
+}
diff --git a/lib/ir/type.cpp b/lib/ir/type.cpp
index 5aebd94a5..10a8d582e 100644
--- a/lib/ir/type.cpp
+++ b/lib/ir/type.cpp
@@ -1,9 +1,9 @@
 #include <cassert>
-#include "ir/type.h"
-#include "ir/context.h"
-#include "ir/context_impl.h"
-#include "ir/value.h"
-#include "ir/constant.h"
+#include "triton/ir/type.h"
+#include "triton/ir/context.h"
+#include "triton/ir/context_impl.h"
+#include "triton/ir/value.h"
+#include "triton/ir/constant.h"
 
 namespace tdl{
 namespace ir{
diff --git a/lib/ir/value.cpp b/lib/ir/value.cpp
index 2d43f8e9d..3827af220 100644
--- a/lib/ir/value.cpp
+++ b/lib/ir/value.cpp
@@ -1,5 +1,5 @@
-#include "ir/value.h"
-#include "ir/instructions.h"
+#include "triton/ir/value.h"
+#include "triton/ir/instructions.h"
 #include <iostream>
 #include <cassert>