[filesystem] rename tdl -> triton

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)

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)

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

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

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{

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
 {

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
 {

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
 {

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
 {

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>

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

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
 {

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
 {

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>
 

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
 {

include/triton/driver/platform.h

@@ -26,7 +26,7 @@
 #include <vector>
 #include <string>
 
-#include "driver/handle.h"
+#include "triton/driver/handle.h"
 
 namespace tdl
 {

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
 {

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{

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{

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{

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

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>
 
 

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{

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 {
 

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 {
 

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{

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{

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"

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 {
 

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>
 

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 {
 

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>

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

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
 {

lib/driver/device.cpp

@@ -26,7 +26,7 @@
 #include <cstring>
 #include <memory>
 
-#include "driver/device.h"
+#include "triton/driver/device.h"
 
 namespace tdl
 {

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
 {

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
 {

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
 {

lib/driver/handle.cpp

@@ -22,7 +22,7 @@
 
 #include <cassert>
 #include <memory>
-#include "driver/handle.h"
+#include "triton/driver/handle.h"
 
 namespace tdl
 {

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
 {

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
 {

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>
 

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
 {

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 {

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{

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{

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{

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{

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{

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{

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;
+  }
+}
+
+
+}
+}