Skip to content

Commit

Permalink
Adding conversion between LinearLayout and DPASLayout (#1684)
Browse files Browse the repository at this point in the history 
This PR fixes #1295.

Upstream Triton tries to abstract all tensor layouts with the
LinearLayout abstraction.
Solving #1295 requires using LinearLayout, so I've added codes to
convert DPASLayout to LinearLayout.
Currently, the conversion is only correct for DPAS OperandC layout.
  • Loading branch information
@hwnam831
hwnam831 authored Jul 31, 2024
1 parent 87005f4 commit cc57449
Show file tree
Hide file tree
Showing 12 changed files with 799 additions and 456 deletions.
177 changes: 0 additions & 177 deletions scripts/skiplist/a770/language.txt
View file

Large diffs are not rendered by default.

49 changes: 0 additions & 49 deletions scripts/skiplist/conda/language.txt
View file
Original file line number Diff line number Diff line change
@@ -1,52 +1,3 @@
# https://github.com/intel/intel-xpu-backend-for-triton/issues/1295
test/unit/language/test_core.py::test_convert2d[dst_layout0-interm_layout1-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout0-interm_layout1-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout0-interm_layout2-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout0-interm_layout2-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout0-interm_layout3-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout0-interm_layout3-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout1-interm_layout1-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout1-interm_layout1-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout1-interm_layout2-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout1-interm_layout2-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout1-interm_layout3-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout1-interm_layout3-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout2-interm_layout1-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout2-interm_layout1-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout2-interm_layout2-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout2-interm_layout2-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout2-interm_layout3-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout2-interm_layout3-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout3-interm_layout1-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout3-interm_layout1-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout3-interm_layout2-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout3-interm_layout2-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout3-interm_layout3-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout3-interm_layout3-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout4-interm_layout1-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout4-interm_layout1-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout4-interm_layout2-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout4-interm_layout2-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout4-interm_layout3-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout4-interm_layout3-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout5-interm_layout1-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout5-interm_layout1-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout5-interm_layout2-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout5-interm_layout2-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout5-interm_layout3-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout5-interm_layout3-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout6-interm_layout1-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout6-interm_layout1-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout6-interm_layout2-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout6-interm_layout2-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout6-interm_layout3-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout6-interm_layout3-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout7-interm_layout1-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout7-interm_layout1-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout7-interm_layout2-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout7-interm_layout2-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout7-interm_layout3-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout7-interm_layout3-src_layout8-float16-64-1]
# https://github.com/intel/intel-xpu-backend-for-triton/issues/1434
test/unit/language/test_core.py::test_precise_math[1-tl.math.sqrt_rn(x)-tl.math.sqrt(x.to(tl.float64)).to(tl.float32)]
test/unit/language/test_core.py::test_dot3d[1-1-32-32-32-32-32-float16-float16]
Expand Down
49 changes: 0 additions & 49 deletions scripts/skiplist/default/language.txt
View file
Original file line number Diff line number Diff line change
@@ -1,52 +1,3 @@
# https://github.com/intel/intel-xpu-backend-for-triton/issues/1295
test/unit/language/test_core.py::test_convert2d[dst_layout0-interm_layout1-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout0-interm_layout1-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout0-interm_layout2-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout0-interm_layout2-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout0-interm_layout3-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout0-interm_layout3-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout1-interm_layout1-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout1-interm_layout1-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout1-interm_layout2-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout1-interm_layout2-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout1-interm_layout3-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout1-interm_layout3-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout2-interm_layout1-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout2-interm_layout1-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout2-interm_layout2-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout2-interm_layout2-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout2-interm_layout3-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout2-interm_layout3-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout3-interm_layout1-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout3-interm_layout1-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout3-interm_layout2-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout3-interm_layout2-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout3-interm_layout3-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout3-interm_layout3-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout4-interm_layout1-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout4-interm_layout1-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout4-interm_layout2-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout4-interm_layout2-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout4-interm_layout3-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout4-interm_layout3-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout5-interm_layout1-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout5-interm_layout1-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout5-interm_layout2-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout5-interm_layout2-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout5-interm_layout3-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout5-interm_layout3-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout6-interm_layout1-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout6-interm_layout1-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout6-interm_layout2-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout6-interm_layout2-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout6-interm_layout3-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout6-interm_layout3-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout7-interm_layout1-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout7-interm_layout1-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout7-interm_layout2-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout7-interm_layout2-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout7-interm_layout3-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout7-interm_layout3-src_layout8-float16-64-1]
# https://github.com/intel/intel-xpu-backend-for-triton/issues/1434
test/unit/language/test_core.py::test_precise_math[1-tl.math.sqrt_rn(x)-tl.math.sqrt(x.to(tl.float64)).to(tl.float32)]
test/unit/language/test_core.py::test_dot3d[1-1-32-32-32-32-32-float16-float16]
Expand Down
49 changes: 0 additions & 49 deletions scripts/skiplist/lts/language.txt
View file
Original file line number Diff line number Diff line change
@@ -1,52 +1,3 @@
# https://github.com/intel/intel-xpu-backend-for-triton/issues/1295
test/unit/language/test_core.py::test_convert2d[dst_layout0-interm_layout1-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout0-interm_layout1-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout0-interm_layout2-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout0-interm_layout2-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout0-interm_layout3-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout0-interm_layout3-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout1-interm_layout1-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout1-interm_layout1-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout1-interm_layout2-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout1-interm_layout2-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout1-interm_layout3-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout1-interm_layout3-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout2-interm_layout1-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout2-interm_layout1-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout2-interm_layout2-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout2-interm_layout2-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout2-interm_layout3-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout2-interm_layout3-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout3-interm_layout1-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout3-interm_layout1-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout3-interm_layout2-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout3-interm_layout2-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout3-interm_layout3-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout3-interm_layout3-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout4-interm_layout1-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout4-interm_layout1-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout4-interm_layout2-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout4-interm_layout2-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout4-interm_layout3-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout4-interm_layout3-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout5-interm_layout1-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout5-interm_layout1-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout5-interm_layout2-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout5-interm_layout2-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout5-interm_layout3-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout5-interm_layout3-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout6-interm_layout1-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout6-interm_layout1-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout6-interm_layout2-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout6-interm_layout2-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout6-interm_layout3-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout6-interm_layout3-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout7-interm_layout1-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout7-interm_layout1-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout7-interm_layout2-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout7-interm_layout2-src_layout8-float16-64-1]
test/unit/language/test_core.py::test_convert2d[dst_layout7-interm_layout3-src_layout8-float16-1-64]
test/unit/language/test_core.py::test_convert2d[dst_layout7-interm_layout3-src_layout8-float16-64-1]
# LTS failures
test/unit/language/test_core.py::test_dot[1-64-128-128-4-True-True-none-tf32-int8-int8-1_0]
test/unit/language/test_core.py::test_dot[1-64-128-128-4-True-True-none-tf32-int8-int8-1_1]
Expand Down
19 changes: 19 additions & 0 deletions third_party/intel/include/Dialect/TritonIntelGPU/IR/LinearLayoutConversions.h
View file
Original file line number Diff line number Diff line change
@@ -0,0 +1,19 @@
// Conversions from TritonIntelGPU DpasEncodingAttr to LinearLayout.

#ifndef TRITON_DIALECT_TRITONINTELGPU_IR_LINEARLAYOUTCONVERSIONS_H
#define TRITON_DIALECT_TRITONINTELGPU_IR_LINEARLAYOUTCONVERSIONS_H

#include <optional>

#include "intel/include/Dialect/TritonIntelGPU/IR/Attributes.h"
#include "intel/include/Dialect/TritonIntelGPU/IR/Dialect.h"
#include "triton/Tools/LinearLayout.h"

namespace mlir::triton::gpu {

std::optional<LinearLayout> DPAStoLinearLayout(ArrayRef<int64_t> shape,
Attribute layout);

} // namespace mlir::triton::gpu

#endif // TRITON_DIALECT_TRITONINTELGPU_IR_LINEARLAYOUTCONVERSIONS_H
2 changes: 1 addition & 1 deletion third_party/intel/lib/Dialect/TritonIntelGPU/IR/CMakeLists.txt
View file
Original file line number Diff line number Diff line change
@@ -1,7 +1,7 @@
add_triton_library(TritonIntelGPUIR
Dialect.cpp
Ops.cpp

LinearLayoutConversions.cpp
DEPENDS
TritonIntelGPUTableGen
TritonIntelGPUAttrDefsIncGen
Expand Down
Loading

0 comments on commit cc57449

Please sign in to comment.