Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Miopen dialect opt step13 : apply index diff maps in threadwise_copy and threadwise_copy_v2. #226

Merged
merged 50 commits into from
Jun 5, 2021
Merged
Show file tree
Hide file tree
Changes from 49 commits
Commits
Show all changes
50 commits
Select commit Hold shift + click to select a range
494df01
Apply index diff map in miopen.threadwise_copy_v2 for loads.
whchung May 21, 2021
453ec6a
Remove unused codes.
whchung May 21, 2021
79a32b2
Apply index diff map in miopen.threadwise_copy_v2 for stores.
whchung May 21, 2021
6cd4fd9
Remove unused codes.
whchung May 21, 2021
7154b29
Factor out common logic.
whchung May 22, 2021
12cb501
Fix clang-format
whchung May 26, 2021
b64af6a
Start to use populayeLayerIndices to compute lower-level coordinates.
whchung May 29, 2021
7ffced2
Remove composedSource/DestTransform plus renaming some variables.
whchung May 29, 2021
cd16a7d
Fix unit tests.
whchung May 29, 2021
5d770e8
Consolidate default lengths of SmallVector instances.
whchung May 29, 2021
47ba8a9
Make loopIVsPerAccessOrder be an argument rather than a captured object.
whchung May 29, 2021
86aa4fe
Supply coordinate transformations metadata to index diff map lambda.
whchung May 29, 2021
92a3494
Split source / dest coordinate transformation specifications.
whchung May 29, 2021
b980382
Populate coord transform attributes for source vectors.
whchung May 29, 2021
8147ea0
Amend unit tests.
whchung May 29, 2021
aa03177
Start to use the metadata and remove inputType from the lambda.
whchung May 29, 2021
10ffc7e
Fix clang-format.
whchung May 29, 2021
d20f8aa
Revise computeIndexDiffMap interface.
whchung May 30, 2021
47a6a88
Move logic to where it's truly needed.
whchung May 30, 2021
d80636a
Start to progressively apply index diff maps.
whchung May 30, 2021
56a2d98
Extract common logic to a lambda.
whchung May 30, 2021
84292ff
Remove unused codes.
whchung May 30, 2021
034fe3f
Rename some variables.
whchung May 30, 2021
5e01f82
Reorder dim_access_order.
whchung May 23, 2021
51bb56d
Change default lengths of SmallVector instances.
whchung May 30, 2021
6a3721e
Switch between legacy and new approach.
whchung May 23, 2021
6f31f58
Carve out lambda from threadwise_copy_v2 to a function.
whchung May 30, 2021
9a31115
Carve out lambda from threadwise_copy_v2 to a function.
whchung May 30, 2021
539200d
Populate initial upper and lower indices for index diff map computation.
whchung May 30, 2021
4c9c4a4
Add logic to cope with incomplete metadata.
whchung May 30, 2021
8e97f92
Adopt index diff map logic in threadwise_copy. Disabled by default.
whchung May 30, 2021
6df96e8
Experimental commit to test 5->3 transfer.
whchung May 31, 2021
4f9fdb7
Change comments.
whchung May 31, 2021
c8bfeed
Supply UnMerge parameters for matrix C write out logic.
whchung May 31, 2021
464b0ac
Supply transform metadata for users of subview op.
whchung May 31, 2021
b2ead3c
Fix clang-format.
whchung May 31, 2021
b86a83f
Proper implementation of index diff maps logic. F_infinite algorithm.
whchung May 31, 2021
96eb946
Fix one unit test.
whchung Jun 1, 2021
734c407
XXX FIXME disable a test for conv2d_bwd_data.
whchung Jun 1, 2021
555dce3
Merge branch 'miopen-dialect' into miopen-dialect-opt-step13
whchung Jun 1, 2021
cecef1a
Fix clang-format.
whchung Jun 1, 2021
f6739e4
Merge branch 'miopen-dialect' into miopen-dialect-opt-step13
whchung Jun 1, 2021
99a5b42
Consider Slice in computeIndexDiffMap. Fix Embed parameters in bwd_data.
whchung Jun 1, 2021
cd94277
XXX. HACKS for bwd_data.
whchung Jun 1, 2021
9f2d243
Merge branch 'miopen-dialect' into miopen-dialect-opt-step13
whchung Jun 2, 2021
ea603b1
Revert "XXX FIXME disable a test for conv2d_bwd_data."
whchung Jun 2, 2021
6486c63
Embed affine maps within the metadata of transformations and use them.
whchung Jun 2, 2021
7347f4d
Use constantFold whenever possible.
whchung Jun 2, 2021
0558040
Merge branch 'miopen-dialect' into miopen-dialect-opt-step13
whchung Jun 3, 2021
a20b61d
Merge branch 'miopen-dialect' into miopen-dialect-opt-step13
whchung Jun 5, 2021
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
20 changes: 0 additions & 20 deletions mlir/include/mlir/Dialect/MIOpen/AffineMapHelper.h
Original file line number Diff line number Diff line change
Expand Up @@ -49,26 +49,6 @@ inline AffineMap composeTransforms(ArrayAttr affineMaps) {
return transform;
}

//===----------------------------------------------------------------------===//
// Check if an AffineMap has division or remainder inside.
//===----------------------------------------------------------------------===//
inline bool hasDivisionOrRemainder(AffineMap map) {
bool ret = false;
if (!map)
return false;
map.walkExprs([&ret](AffineExpr expr) {
if (expr.getKind() == AffineExprKind::Mod ||
expr.getKind() == AffineExprKind::FloorDiv ||
expr.getKind() == AffineExprKind::CeilDiv)
ret = true;
});

// XXX. hack. always return false for now for performance reason.
// May need more sophisticated checks to determine if we would truly go OOB.
// return ret;
return false;
}

//===----------------------------------------------------------------------===//
// Check if an AffineExpr has padding, which is represented as a minus
// expression with a constant operand.
Expand Down
Loading