Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
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
2 changes: 1 addition & 1 deletion README.md
Original file line number Diff line number Diff line change
Expand Up @@ -6,7 +6,7 @@ Triton-XDNA provides an end-to-end compilation flow that lowers standard Triton

### How it works

Triton kernels are first lowered to compact Linalg compute graphs via [triton-shared](https://github.com/microsoft/triton-shared), then tiled and mapped onto parallel NPU cores using the MLIR Transform dialect, and finally compiled through [MLIR-AIR](https://github.com/Xilinx/mlir-air) and [MLIR-AIE](https://github.com/Xilinx/mlir-aie) to produce device binaries.
Triton kernels are first lowered to compact Linalg compute graphs via [triton-shared](https://github.com/facebookincubator/triton-shared), then tiled and mapped onto parallel NPU cores using the MLIR Transform dialect, and finally compiled through [MLIR-AIR](https://github.com/Xilinx/mlir-air) and [MLIR-AIE](https://github.com/Xilinx/mlir-aie) to produce device binaries.

```
Triton kernel (@triton.jit)
Expand Down
2 changes: 1 addition & 1 deletion third_party/triton_shared
Submodule triton_shared updated 71 files
+1 βˆ’1 CMakeLists.txt
+1 βˆ’1 LICENSE
+24 βˆ’2 README.md
+0 βˆ’1 include/triton-shared/Conversion/CMakeLists.txt
+0 βˆ’9 include/triton-shared/Conversion/TritonToLinalg/CMakeLists.txt
+0 βˆ’22 include/triton-shared/Conversion/TritonToLinalg/Passes.h
+0 βˆ’18 include/triton-shared/Conversion/TritonToLinalg/Passes.td
+0 βˆ’33 include/triton-shared/Conversion/TritonToLinalg/TritonToLinalg.h
+1 βˆ’1 include/triton-shared/Conversion/UnstructuredToMemref/CMakeLists.txt
+0 βˆ’1 lib/Conversion/CMakeLists.txt
+0 βˆ’27 lib/Conversion/TritonToLinalg/CMakeLists.txt
+0 βˆ’95 lib/Conversion/TritonToLinalg/TritonToLinalg.cpp
+0 βˆ’229 lib/Conversion/TritonToLinalg/TritonToLinalgPass.cpp
+1 βˆ’1 lib/Conversion/UnstructuredToMemref/CMakeLists.txt
+0 βˆ’69 test/Conversion/TritonToLinalg/addptr_2d_example.mlir
+0 βˆ’68 test/Conversion/TritonToLinalg/addptr_add_value.mlir
+0 βˆ’113 test/Conversion/TritonToLinalg/addptr_dim1.mlir
+0 βˆ’92 test/Conversion/TritonToLinalg/addptr_for_accumulation.mlir
+0 βˆ’73 test/Conversion/TritonToLinalg/addptr_for_expand_ptr.mlir
+0 βˆ’71 test/Conversion/TritonToLinalg/addptr_for_more_init_args.mlir
+0 βˆ’98 test/Conversion/TritonToLinalg/addptr_for_used_after_update.mlir
+0 βˆ’55 test/Conversion/TritonToLinalg/addptr_for_used_before_update.mlir
+0 βˆ’53 test/Conversion/TritonToLinalg/addptr_loopback.mlir
+0 βˆ’49 test/Conversion/TritonToLinalg/addptr_mul_const_const.mlir
+0 βˆ’51 test/Conversion/TritonToLinalg/addptr_mul_value_const.mlir
+0 βˆ’73 test/Conversion/TritonToLinalg/addptr_nested.mlir
+0 βˆ’43 test/Conversion/TritonToLinalg/addptr_reshape_broadcast.mlir
+0 βˆ’65 test/Conversion/TritonToLinalg/addptr_scalar_broadcast.mlir
+0 βˆ’70 test/Conversion/TritonToLinalg/addptr_scalar_for.mlir
+0 βˆ’92 test/Conversion/TritonToLinalg/addptr_scalar_for_2d.mlir
+0 βˆ’27 test/Conversion/TritonToLinalg/addptr_scalar_loopback.mlir
+0 βˆ’57 test/Conversion/TritonToLinalg/addptr_scalar_nested.mlir
+0 βˆ’45 test/Conversion/TritonToLinalg/addptr_scalar_splat.mlir
+0 βˆ’56 test/Conversion/TritonToLinalg/addptr_scalar_splat_2d.mlir
+0 βˆ’39 test/Conversion/TritonToLinalg/arith_not_ptr_arith.mlir
+0 βˆ’44 test/Conversion/TritonToLinalg/bitcast.mlir
+0 βˆ’90 test/Conversion/TritonToLinalg/block_ptr_advance.mlir
+0 βˆ’72 test/Conversion/TritonToLinalg/convert_1d_elemwise_arith_binary.mlir
+0 βˆ’49 test/Conversion/TritonToLinalg/convert_1d_elemwise_arith_ternary.mlir
+0 βˆ’88 test/Conversion/TritonToLinalg/convert_1d_elemwise_arith_unary.mlir
+0 βˆ’55 test/Conversion/TritonToLinalg/convert_2d_elemwise_arith_binary.mlir
+0 βˆ’55 test/Conversion/TritonToLinalg/convert_2d_elemwise_arith_ternary.mlir
+0 βˆ’94 test/Conversion/TritonToLinalg/convert_2d_elemwise_arith_unary.mlir
+0 βˆ’32 test/Conversion/TritonToLinalg/convert_addi_reduce.mlir
+0 βˆ’141 test/Conversion/TritonToLinalg/convert_argmin_argmax.mlir
+0 βˆ’215 test/Conversion/TritonToLinalg/convert_argmin_argmax_2d.mlir
+0 βˆ’809 test/Conversion/TritonToLinalg/convert_extern_elementwise.mlir
+0 βˆ’50 test/Conversion/TritonToLinalg/convert_minmax.mlir
+0 βˆ’68 test/Conversion/TritonToLinalg/convert_minmax_fp_reduce.mlir
+0 βˆ’126 test/Conversion/TritonToLinalg/convert_minmax_reduce.mlir
+0 βˆ’23 test/Conversion/TritonToLinalg/convert_splat_float.mlir
+0 βˆ’45 test/Conversion/TritonToLinalg/convert_tensor_reshape.mlir
+0 βˆ’68 test/Conversion/TritonToLinalg/cumsum.mlir
+0 βˆ’84 test/Conversion/TritonToLinalg/dot.mlir
+0 βˆ’45 test/Conversion/TritonToLinalg/get_num_programs.mlir
+0 βˆ’58 test/Conversion/TritonToLinalg/reducemax_32_256_bf16.mlir
+0 βˆ’51 test/Conversion/TritonToLinalg/reducesum_512_256_bf16_axis0.mlir
+0 βˆ’53 test/Conversion/TritonToLinalg/reducesum_512_256_bf16_axis1.mlir
+0 βˆ’51 test/Conversion/TritonToLinalg/reducesum_512_256_f32_axis0.mlir
+0 βˆ’53 test/Conversion/TritonToLinalg/reducesum_512_256_f32_axis1.mlir
+0 βˆ’60 test/Conversion/TritonToLinalg/reducesum_middle_dim.mlir
+0 βˆ’38 test/Conversion/TritonToLinalg/reducesum_scalar.mlir
+0 βˆ’50 test/Conversion/TritonToLinalg/triton_assert.mlir
+0 βˆ’35 test/Conversion/TritonToLinalg/unsupported_extern_elementwise.mlir
+0 βˆ’76 test/Conversion/TritonToLinalg/use_dot_opc.mlir
+0 βˆ’95 test/Conversion/TritonToLinalg/use_end_chain.mlir
+0 βˆ’64 test/Conversion/TritonToLinalg/use_mid_chain.mlir
+0 βˆ’133 test/Conversion/TritonToLinalg/wraparound_side_by_side.mlir
+0 βˆ’129 test/Conversion/TritonToLinalg/wraparound_stacked.mlir
+0 βˆ’57 test/Conversion/TritonToLinalg/wraparound_unsupported_add_offset.mlir
+0 βˆ’2 tools/triton-shared-opt/RegisterTritonSharedDialects.h
Loading
Loading