8d163e5045
It adds a `shuffle_16x16` strategy LowerVectorTranspose and renames `shuffle` to `shuffle_1d`. The idea is similar to 8x8 cases in x86Vector::avx2. The general algorithm is:
```
interleave 32-bit lanes using
8x _mm512_unpacklo_epi32
8x _mm512_unpackhi_epi32
interleave 64-bit lanes using
8x _mm512_unpacklo_epi64
8x _mm512_unpackhi_epi64
permute 128-bit lanes using
16x _mm512_shuffle_i32x4
permute 256-bit lanes using again
16x _mm512_shuffle_i32x4
```
After the first stage, they got transposed to
```
0 16 1 17 4 20 5 21 8 24 9 25 12 28 13 29
2 18 3 19 6 22 7 23 10 26 11 27 14 30 15 31
32 48 33 49 ...
34 50 35 51 ...
64 80 65 81 ...
...
```
After the second stage, they got transposed to
```
0 16 32 48 ...
1 17 33 49 ...
2 18 34 49 ...
3 19 35 51 ...
64 80 96 112 ...
65 81 97 114 ...
66 82 98 113 ...
67 83 99 115 ...
...
```
After the thrid stage, they got transposed to
```
0 16 32 48 8 24 40 56 64 80 96 112 ...
1 17 33 49 ...
2 18 34 50 ...
3 19 35 51 ...
4 20 36 52 ...
5 21 37 53 ...
6 22 38 54 ...
7 23 39 55 ...
128 144 160 176 ...
129 145 161 177 ...
...
```
After the last stage, they got transposed to
```
0 16 32 48 64 80 96 112 ... 240
1 17 33 49 66 81 97 113 ... 241
2 18 34 50 67 82 98 114 ... 242
...
15 31 47 63 79 96 111 127 ... 255
```
Reviewed By: dcaballe
Differential Revision: https://reviews.llvm.org/D148685
63 lines
2.4 KiB
MLIR
63 lines
2.4 KiB
MLIR
// RUN: mlir-opt %s --test-transform-dialect-interpreter --split-input-file | FileCheck %s
|
|
|
|
// CHECK-LABEL: func @matmul_tensors
|
|
func.func @matmul_tensors(
|
|
%arg0: tensor<8x16xf32>, %arg1: tensor<16x32xf32>, %arg2: tensor<8x32xf32>)
|
|
-> tensor<8x32xf32> {
|
|
// CHECK-NOT: linalg
|
|
// CHECK: vector.extract {{.*}} : vector<8x4xf32>
|
|
// CHECK: vector.store {{.*}} : memref<8x32xf32>, vector<4xf32>
|
|
%0 = linalg.matmul ins(%arg0, %arg1: tensor<8x16xf32>, tensor<16x32xf32>)
|
|
outs(%arg2: tensor<8x32xf32>)
|
|
-> tensor<8x32xf32>
|
|
return %0 : tensor<8x32xf32>
|
|
}
|
|
|
|
transform.sequence failures(propagate) {
|
|
^bb1(%module_op: !pdl.operation):
|
|
%0 = transform.structured.match ops{["linalg.matmul"]} in %module_op : (!pdl.operation) -> !pdl.operation
|
|
%1, %loops:3 = transform.structured.tile %0 [8, 4, 2]
|
|
: (!pdl.operation) -> (!pdl.operation, !pdl.operation, !pdl.operation, !pdl.operation)
|
|
%2 = get_closest_isolated_parent %1 : (!pdl.operation) -> !pdl.operation
|
|
transform.structured.vectorize %2
|
|
%b = transform.bufferization.one_shot_bufferize
|
|
layout{IdentityLayoutMap} %module_op
|
|
{bufferize_function_boundaries = true, allow_return_allocs = true}
|
|
: (!pdl.operation) -> !pdl.operation
|
|
|
|
%f = transform.structured.match ops{["func.func"]} in %b
|
|
: (!pdl.operation) -> !pdl.operation
|
|
|
|
// TODO: group these lower-level controls into various properly named vector
|
|
// lowering TD macros.
|
|
%func = transform.vector.lower_contraction %f
|
|
lowering_strategy = "outerproduct"
|
|
: (!pdl.operation) -> !pdl.operation
|
|
|
|
%func_2 = transform.vector.apply_transfer_permutation_patterns %func
|
|
: (!pdl.operation) -> !pdl.operation
|
|
|
|
%func_3 = transform.vector.lower_multi_reduction %func_2
|
|
lowering_strategy = "innerparallel"
|
|
: (!pdl.operation) -> !pdl.operation
|
|
|
|
%func_4 = transform.vector.split_transfer_full_partial %func_3
|
|
split_transfer_strategy = "linalg-copy"
|
|
: (!pdl.operation) -> !pdl.operation
|
|
|
|
%func_5 = transform.vector.transfer_to_scf %func_4
|
|
max_transfer_rank = 1 full_unroll = true
|
|
: (!pdl.operation) -> !pdl.operation
|
|
|
|
%func_6 = transform.vector.lower_transfer %func_5
|
|
max_transfer_rank = 1
|
|
: (!pdl.operation) -> !pdl.operation
|
|
|
|
%func_7 = transform.vector.lower_shape_cast %func_6
|
|
: (!pdl.operation) -> !pdl.operation
|
|
|
|
%func_8 = transform.vector.lower_transpose %func_7
|
|
lowering_strategy = "shuffle_1d"
|
|
: (!pdl.operation) -> !pdl.operation
|
|
}
|